From c340a2d6083f34b9434cc4abbb236c02785e4ce0 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Ole=20Sch=C3=BCtt?= <ole.schuett@cp2k.org>
Date: Mon, 15 May 2023 21:57:33 +0200
Subject: [PATCH] grid: Fix bug in GPU integrate kernel for large basis sets

---
 src/grid/gpu/grid_gpu_integrate.cu | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/src/grid/gpu/grid_gpu_integrate.cu b/src/grid/gpu/grid_gpu_integrate.cu
index ba49a0efb7..4a38c43479 100644
--- 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