grid: Fix bug in GPU integrate kernel for large basis sets

cp2k · May 16, 2023 · c340a2d · c340a2d
1 parent 98d2d7a
commit c340a2d
Showing 1 changed file with 3 additions and 1 deletion.
diff --git a/src/grid/gpu/grid_gpu_integrate.cu b/src/grid/gpu/grid_gpu_integrate.cu
@@ -255,7 +255,9 @@ __device__ static void grid_to_cxyz(const kernel_params *params,
     // Add register values to coefficients stored in shared memory.
 #pragma unroll // avoid dynamic indexing of registers
     for (int i = 0; i < GRID_N_CXYZ_REGISTERS; i++) {
-      atomicAddDouble(&cxyz[i + offset], cxyz_regs[i]);
+      if (i + offset < ncoset(task->lp)) {
+        atomicAddDouble(&cxyz[i + offset], cxyz_regs[i]);
+      }
     }
   }
   __syncthreads(); // because of concurrent writes to cxyz