Fix cuda compiling bug

Critsium-xy · Critsium-xy · commit 0e0d7fb80051 · 2024-11-08T18:49:44.000+08:00
diff --git a/source/module_hsolver/kernels/cuda/math_kernel_op.cu b/source/module_hsolver/kernels/cuda/math_kernel_op.cu
@@ -12,7 +12,7 @@
 namespace hsolver
 {
 const int warp_size = 32;
-//const unsigned int full_mask = 0xffffffff;
+// const unsigned int full_mask = 0xffffffff;
 const int thread_per_block = 256;
 }
 
@@ -65,11 +65,11 @@ void destoryBLAShandle(){
     }
 }
 
-template <typename FPTYPE>
-__forceinline__ __device__ void warp_reduce(FPTYPE& val) {
-    for (int offset = 16; offset > 0; offset >>= 1)
-        val += __shfl_down_sync(full_mask, val, offset);
-}
+// template <typename FPTYPE>
+// __forceinline__ __device__ void warp_reduce(FPTYPE& val) {
+//     for (int offset = 16; offset > 0; offset >>= 1)
+//         val += __shfl_down_sync(full_mask, val, offset);
+// }
 
 template <typename Real>
 __global__ void line_minimize_with_block(

Original file line number	Diff line number	Diff line change
`@@ -12,7 +12,7 @@`
`12`	`12`	`namespace hsolver`
`13`	`13`	`{`
`14`	`14`	`const int warp_size = 32;`
`15`		`-//const unsigned int full_mask = 0xffffffff;`
	`15`	`+// const unsigned int full_mask = 0xffffffff;`
`16`	`16`	`const int thread_per_block = 256;`
`17`	`17`	`}`
`18`	`18`
`@@ -65,11 +65,11 @@ void destoryBLAShandle(){`
`65`	`65`	`}`
`66`	`66`	`}`
`67`	`67`
`68`		`-template <typename FPTYPE>`
`69`		`-__forceinline__ __device__ void warp_reduce(FPTYPE& val) {`
`70`		`- for (int offset = 16; offset > 0; offset >>= 1)`
`71`		`- val += __shfl_down_sync(full_mask, val, offset);`
`72`		`-}`
	`68`	`+// template <typename FPTYPE>`
	`69`	`+// __forceinline__ __device__ void warp_reduce(FPTYPE& val) {`
	`70`	`+// for (int offset = 16; offset > 0; offset >>= 1)`
	`71`	`+// val += __shfl_down_sync(full_mask, val, offset);`
	`72`	`+// }`
`73`	`73`
`74`	`74`	`template <typename Real>`
`75`	`75`	`__global__ void line_minimize_with_block(`