Merge pull request #2745 from rbberger/gpu_updates

Updates to the GPU package

cmake/Modules/GenerateOpenCLHeader.cmake

@@ -0,0 +1,3 @@
+# utility script to call WriteOpenCLHeader function
+include(${SOURCE_DIR}/Modules/OpenCLUtils.cmake)
+WriteOpenCLHeader(${VARNAME} ${HEADER_FILE} ${SOURCE_FILES})

cmake/Modules/OpenCLUtils.cmake

@@ -1,10 +1,8 @@
-function(GenerateOpenCLHeader varname outfile files)
-    message("Creating ${outfile}...")
+function(WriteOpenCLHeader varname outfile files)
     file(WRITE ${outfile} "const char * ${varname} = \n")
-    math(EXPR ARG_END   "${ARGC}-1")
+    separate_arguments(files)
 
-    foreach(IDX RANGE 2 ${ARG_END})
-        list(GET ARGV ${IDX} filename)
+    foreach(filename ${files})
         file(READ ${filename} content)
         string(REGEX REPLACE "\\s*//[^\n]*\n" "\n" content "${content}")
         string(REGEX REPLACE "\\\\" "\\\\\\\\" content "${content}")
@@ -15,4 +13,16 @@ function(GenerateOpenCLHeader varname outfile files)
     endforeach()
 
     file(APPEND ${outfile} ";\n")
+endfunction(WriteOpenCLHeader)
+
+function(GenerateOpenCLHeader varname outfile files)
+    list(REMOVE_AT ARGV 0 1)
+    add_custom_command(OUTPUT ${outfile}
+      COMMAND ${CMAKE_COMMAND} -D SOURCE_DIR=${CMAKE_CURRENT_SOURCE_DIR}
+                               -D VARNAME=${varname}
+                               -D HEADER_FILE=${outfile}
+                               -D SOURCE_FILES="${ARGV}"
+                               -P ${CMAKE_CURRENT_SOURCE_DIR}/Modules/GenerateOpenCLHeader.cmake
+      DEPENDS ${ARGV}
+      COMMENT "Generating ${outfile}...")
 endfunction(GenerateOpenCLHeader)

lib/gpu/geryon/nvd_device.h

@@ -24,10 +24,6 @@
 #ifndef NVD_DEVICE
 #define NVD_DEVICE
 
-// workaround after GPU package Feb2021 update
-// todo: make new neighbor code work with CUDA
-#define LAL_USE_OLD_NEIGHBOR
-
 #include <string>
 #include <vector>
 #include <iostream>

lib/gpu/lal_neighbor_gpu.cu

@@ -115,7 +115,7 @@ __kernel void kernel_calc_cell_counts(const unsigned *restrict cell_id,
 #define tagint int
 #endif
 #ifdef LAMMPS_BIGBIG
-#define tagint long long int
+#define tagint long
 #endif
 #ifdef LAMMPS_SMALLSMALL
 #define tagint int