fix cuda 11.8 compile problem

FindDefinition · Dec 9, 2024 · 3463b5f · 3463b5f
1 parent c5e40f9
commit 3463b5f
Show file tree

Hide file tree

Showing 4 changed files with 14 additions and 11 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -1,4 +1,7 @@
 # Changelog
+## [0.7.7] - 2024-12-09
+- fix cuda 11 compile problem
+
 ## [0.7.6] - 2024-12-09
 - fix missing windows cpu prebuilt
 

diff --git a/include/tensorview/core/arrayops/mathbase.h b/include/tensorview/core/arrayops/mathbase.h
@@ -786,13 +786,13 @@ template <> struct MathScalarOp<__half> {
 #endif
   }
 
-  TV_HOST_DEVICE_INLINE static __half max(__half x, __half y) { return __hmax(x, y); }
+  TV_DEVICE_INLINE static __half max(__half x, __half y) { return __hmax(x, y); }
 
-  TV_HOST_DEVICE_INLINE static __half min(__half x, __half y) { return __hmin(x, y); }
+  TV_DEVICE_INLINE static __half min(__half x, __half y) { return __hmin(x, y); }
 
-  TV_HOST_DEVICE_INLINE static __half clamp(__half v, __half lo, __half hi) { return min(hi, max(lo, v)); }
+  TV_DEVICE_INLINE static __half clamp(__half v, __half lo, __half hi) { return min(hi, max(lo, v)); }
 
-  TV_DEVICE_INLINE static __half mix(__half x, __half y, __half t) { return fma(t, y, fma(-t, x, x)); }
+  TV_DEVICE_INLINE static __half mix(__half x, __half y, __half t) { return fma(t, y, fma(neg(t), x, x)); }
 
 };
 
@@ -1107,15 +1107,15 @@ template <> struct MathScalarOp<__nv_bfloat16> {
 #endif
   }
 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800)
-  TV_HOST_DEVICE_INLINE static __nv_bfloat16 max(__nv_bfloat16 x, __nv_bfloat16 y) { return __hmax(x, y); }
+  TV_DEVICE_INLINE static __nv_bfloat16 max(__nv_bfloat16 x, __nv_bfloat16 y) { return __hmax(x, y); }
 
-  TV_HOST_DEVICE_INLINE static __nv_bfloat16 min(__nv_bfloat16 x, __nv_bfloat16 y) { return __hmin(x, y); }
+  TV_DEVICE_INLINE static __nv_bfloat16 min(__nv_bfloat16 x, __nv_bfloat16 y) { return __hmin(x, y); }
 
-  TV_HOST_DEVICE_INLINE static __nv_bfloat16 clamp(__nv_bfloat16 v, __nv_bfloat16 lo, __nv_bfloat16 hi) { return min(hi, max(lo, v)); }
+  TV_DEVICE_INLINE static __nv_bfloat16 clamp(__nv_bfloat16 v, __nv_bfloat16 lo, __nv_bfloat16 hi) { return min(hi, max(lo, v)); }
 
-  TV_HOST_DEVICE_INLINE static __nv_bfloat16 fma(__nv_bfloat16 x, __nv_bfloat16 y, __nv_bfloat16 z) { return __hfma(x, y, z); }
+  TV_DEVICE_INLINE static __nv_bfloat16 fma(__nv_bfloat16 x, __nv_bfloat16 y, __nv_bfloat16 z) { return __hfma(x, y, z); }
 
-  TV_HOST_DEVICE_INLINE static __nv_bfloat16 mix(__nv_bfloat16 x, __nv_bfloat16 y, __nv_bfloat16 t) { return fma(t, y, fma(-t, x, x)); }
+  TV_DEVICE_INLINE static __nv_bfloat16 mix(__nv_bfloat16 x, __nv_bfloat16 y, __nv_bfloat16 t) { return fma(t, y, fma(neg(t), x, x)); }
 
 #endif
 

diff --git a/tools/linux_test_build.sh b/tools/linux_test_build.sh
@@ -1 +1 @@
-sudo docker run --rm -it -e CUMM_CUDA_VERSION="126" -e PLAT=manylinux_2_28_x86_64 -v `pwd`:/io -v $HOME:/myhome manylinux2014-cuda:cu126-devel bash -c "source /etc/bashrc && /io/tools/build-wheels-dev.sh"
+sudo docker run --rm -it -e CUMM_CUDA_VERSION="118" -e PLAT=manylinux2014_x86_64 -v `pwd`:/io -v $HOME:/myhome scrin/manylinux2014-cuda:cu118-devel-1.0.0 bash -c "source /etc/bashrc && /io/tools/build-wheels-dev.sh"
diff --git a/version.txt b/version.txt
@@ -1 +1 @@
-0.7.6
+0.7.7
Original file line number	Diff line number	Diff line change
		@@ -1 +1 @@
		sudo docker run --rm -it -e CUMM_CUDA_VERSION="126" -e PLAT=manylinux_2_28_x86_64 -v `pwd`:/io -v $HOME:/myhome manylinux2014-cuda:cu126-devel bash -c "source /etc/bashrc && /io/tools/build-wheels-dev.sh"
		sudo docker run --rm -it -e CUMM_CUDA_VERSION="118" -e PLAT=manylinux2014_x86_64 -v `pwd`:/io -v $HOME:/myhome scrin/manylinux2014-cuda:cu118-devel-1.0.0 bash -c "source /etc/bashrc && /io/tools/build-wheels-dev.sh"