Merge pull request #2412 from mkuron/hip

Remove default arguments from CUDA kernels
espressomd · Dec 10, 2018 · 418db6b · 418db6b
2 parents 926ba87 + c213e54
commit 418db6b
Showing 1 changed file with 8 additions and 8 deletions.
diff --git 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));