Minor CUDA Changes (#2947)

* Changes name of NVRTC program

* Adds launch bounds for findInteractingBlocks

* Replaces launch bound parameter with named constant
Co-authored-by: David Clark <daclark@nvidia.com>

platforms/cuda/src/kernels/findInteractingBlocks.cu

@@ -176,7 +176,7 @@ __device__ int saveSinglePairs(int x, int* atoms, int* flags, int length, unsign
  * [in] rebuildNeighbourList   - whether or not to execute this kernel
  *
  */
-extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodicBoxSize, real4 periodicBoxVecX, real4 periodicBoxVecY, real4 periodicBoxVecZ,
+extern "C" __global__ __launch_bounds__(GROUP_SIZE,1) void findBlocksWithInteractions(real4 periodicBoxSize, real4 invPeriodicBoxSize, real4 periodicBoxVecX, real4 periodicBoxVecY, real4 periodicBoxVecZ,
         unsigned int* __restrict__ interactionCount, int* __restrict__ interactingTiles, unsigned int* __restrict__ interactingAtoms,
         int2* __restrict__ singlePairs, const real4* __restrict__ posq, unsigned int maxTiles, unsigned int maxSinglePairs,
         unsigned int startBlockIndex, unsigned int numBlocks, real2* __restrict__ sortedBlocks, const real4* __restrict__ sortedBlockCenter,

plugins/cudacompiler/src/CudaCompilerKernels.cpp

@@ -64,7 +64,7 @@ string CudaRuntimeCompilerKernel::createModule(const string& source, const strin
     // Compile the program to PTX.
     
     nvrtcProgram program;
-    CHECK_RESULT(nvrtcCreateProgram(&program, source.c_str(), "", 0, NULL, NULL), "Error creating program");
+    CHECK_RESULT(nvrtcCreateProgram(&program, source.c_str(), NULL, 0, NULL, NULL), "Error creating program");
     try {
         nvrtcResult result = nvrtcCompileProgram(program, options.size(), &options[0]);
         if (result != NVRTC_SUCCESS) {