From c213e5430c80fa9f63b33e836a4cbd8a0a98e844 Mon Sep 17 00:00:00 2001
From: Michael Kuron <mkuron@icp.uni-stuttgart.de>
Date: Mon, 10 Dec 2018 09:58:52 +0100
Subject: [PATCH] Remove default arguments from CUDA kernels

fixes build on AMD after recent update
---
 src/core/actor/Mmm1dgpuForce_cuda.cu | 16 ++++++++--------
 1 file changed, 8 insertions(+), 8 deletions(-)

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));