comp-chem/cp2k
1
0
Fork 0
mirror of https://github.com/cp2k/cp2k.git synced 2026-07-29 06:35:28 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

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

Browse source
This commit is contained in:
Ole Schütt 2023-05-15 21:57:33 +02:00 • committed by Ole Schütt
parent 98d2d7a0dd
commit c340a2d608
1 changed files with 3 additions and 1 deletions
Download patch file Download diff file Expand all files Collapse all files

4
src/grid/gpu/grid_gpu_integrate.cu
View file

@ -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