Fix mask in shfl sync.. (github issue #2339)

catboost · May 3, 2023 · 4e49a45 · 4e49a45
1 parent b6ebc3a
commit 4e49a45
Show file tree

Hide file tree

Showing 2 changed files with 3 additions and 3 deletions.
diff --git a/catboost/cuda/methods/kernel/linear_solver.cu b/catboost/cuda/methods/kernel/linear_solver.cu
@@ -132,7 +132,7 @@ namespace NKernel {
                 }
 
                 float sum = ShuffleReduce(x, tmp, min(reduceSize, 32));
-                sum = __shfl_sync(0xFFFFFF, sum, 0, logicalWarpSize);
+                sum = __shfl_sync(0xFFFFFFFF, sum, 0, logicalWarpSize);
 
 
                 const float ljj = Ljj[0];
@@ -158,7 +158,7 @@ namespace NKernel {
                 }
 
                 float sum = ShuffleReduce(x, tmp, min(reduceSize, 32));
-                sum = __shfl_sync(0xFFFFFF, sum, 0, logicalWarpSize);
+                sum = __shfl_sync(0xFFFFFFFF, sum, 0, logicalWarpSize);
 
                 __syncwarp();
 

diff --git a/catboost/cuda/targets/kernel/dcg.cu b/catboost/cuda/targets/kernel/dcg.cu
@@ -369,7 +369,7 @@ __global__ void RemoveGroupMeanImpl(
     }
 
     T mean = ShuffleReduce<T>(localThreadIdx, localMean, LogicalWarpSize);
-    mean = __shfl_sync(0xFFFFFF, mean, 0, LogicalWarpSize);
+    mean = __shfl_sync(0xFFFFFFFF, mean, 0, LogicalWarpSize);
 
     for (ui32 i = localThreadIdx; i < groupSize; i += LogicalWarpSize) {
         normalized[i] = __ldg(values + i) - mean;