zero.cu

__global__ void
#if (__CUDA_ARCH__ >= 300)
__launch_bounds__(SM_3X_GENERAL_THREADS_PER_BLOCK, 1)
#else
__launch_bounds__(SM_2X_GENERAL_THREADS_PER_BLOCK, 1)
#endif
kZeroSpecificForces_kernel()
{
    __threadfence();
    __syncthreads();
    unsigned int myatoms [] = {4, 6, 8, 14};         // zero these atoms
    for (int i = 0; i < 4; i++)
    {
        int p1 = cSim.pImageAtomLookup[myatoms[i]];
        cSim.pForceAccumulator[p1]                                      = FORCESCALE * 0;
        cSim.pForceAccumulator[p1 + cSim.stride]                        = FORCESCALE * 0;
        cSim.pForceAccumulator[p1 + cSim.stride2]                       = FORCESCALE * 0;
    }
}

void kZeroSpecificForces(gpuContext gpu)
{
    kZeroSpecificForces_kernel<<<gpu->blocks, gpu->generalThreadsPerBlock>>>();
    LAUNCHERROR("kZeroSpecificForces");
}