diff --git a/src/core/actor/Mmm1dgpuForce_cuda.cu b/src/core/actor/Mmm1dgpuForce_cuda.cu index 9023acc20a1..d8e399cbc61 100644 --- a/src/core/actor/Mmm1dgpuForce_cuda.cu +++ b/src/core/actor/Mmm1dgpuForce_cuda.cu @@ -319,7 +319,7 @@ void Mmm1dgpuForce::set_params(mmm1dgpu_real _boxz, __global__ void forcesKernel(const mmm1dgpu_real *__restrict__ r, const mmm1dgpu_real *__restrict__ q, mmm1dgpu_real *__restrict__ force, int N, - int pairs, int tStart = 0, int tStop = -1) { + int pairs, int tStart, int tStop) { if (tStop < 0) tStop = N * N; @@ -406,7 +406,7 @@ __global__ void forcesKernel(const mmm1dgpu_real *__restrict__ r, __global__ void energiesKernel(const mmm1dgpu_real *__restrict__ r, const mmm1dgpu_real *__restrict__ q, mmm1dgpu_real *__restrict__ energy, int N, - int pairs, int tStart = 0, int tStop = -1) { + int pairs, int tStart, int tStop) { if (tStop < 0) tStop = N * N; @@ -474,7 +474,7 @@ __global__ void energiesKernel(const mmm1dgpu_real *__restrict__ r, } __global__ void vectorReductionKernel(mmm1dgpu_real *src, mmm1dgpu_real *dst, - int N, int tStart = 0, int tStop = -1) { + int N, int tStart, int tStop) { if (tStop < 0) tStop = N * N; @@ -511,12 +511,12 @@ void Mmm1dgpuForce::computeForces(SystemInterface &s) { { int blocksRed = s.npart_gpu() / numThreads + 1; KERNELCALL(forcesKernel, numBlocks(s), numThreads, s.rGpuBegin(), - s.qGpuBegin(), dev_forcePairs, s.npart_gpu(), pairs) + s.qGpuBegin(), dev_forcePairs, s.npart_gpu(), pairs, 0, -1) KERNELCALL(vectorReductionKernel, blocksRed, numThreads, dev_forcePairs, - s.fGpuBegin(), s.npart_gpu()) + s.fGpuBegin(), s.npart_gpu(), 0, -1) } else { KERNELCALL(forcesKernel, numBlocks(s), numThreads, s.rGpuBegin(), - s.qGpuBegin(), s.fGpuBegin(), s.npart_gpu(), pairs) + s.qGpuBegin(), s.fGpuBegin(), s.npart_gpu(), pairs, 0, -1) } } @@ -547,7 +547,7 @@ void Mmm1dgpuForce::computeEnergy(SystemInterface &s) { KERNELCALL_shared(energiesKernel, numBlocks(s), numThreads, shared, s.rGpuBegin(), s.qGpuBegin(), dev_energyBlocks, - s.npart_gpu(), 0); + s.npart_gpu(), 0, 0, -1); KERNELCALL_shared(sumKernel, 1, numThreads, shared, dev_energyBlocks, numBlocks(s)); KERNELCALL(scaleAndAddKernel, 1, 1, &(((CUDA_energy *)s.eGpu())->coulomb), @@ -567,7 +567,7 @@ float Mmm1dgpuForce::force_benchmark(SystemInterface &s) { cuda_safe_mem(cudaEventCreate(&eventStop)); cuda_safe_mem(cudaEventRecord(eventStart, stream[0])); KERNELCALL(forcesKernel, numBlocks(s), numThreads, s.rGpuBegin(), - s.qGpuBegin(), dev_f_benchmark, s.npart_gpu(), 0) + s.qGpuBegin(), dev_f_benchmark, s.npart_gpu(), 0, 0, -1) cuda_safe_mem(cudaEventRecord(eventStop, stream[0])); cuda_safe_mem(cudaEventSynchronize(eventStop)); cuda_safe_mem(cudaEventElapsedTime(&elapsedTime, eventStart, eventStop));