From b2cdc4091908478a3054b55d48d3e2516ccdf6ef Mon Sep 17 00:00:00 2001
From: Paulius Velesko <paulius.velesko@intel.com>
Date: Tue, 26 Jul 2022 16:00:43 +0000
Subject: [PATCH] Enable CHIP-SPV support

---
 cmake/Modules/Packages/GPU.cmake | 14 ++++++++++++--
 lib/gpu/lal_pre_cuda_hip.h       | 12 ++++++------
 2 files changed, 18 insertions(+), 8 deletions(-)

diff --git a/cmake/Modules/Packages/GPU.cmake b/cmake/Modules/Packages/GPU.cmake
index 75569aa55d..36409378be 100644
--- a/cmake/Modules/Packages/GPU.cmake
+++ b/cmake/Modules/Packages/GPU.cmake
@@ -233,7 +233,7 @@ elseif(GPU_API STREQUAL "OPENCL")
 elseif(GPU_API STREQUAL "HIP")
   if(NOT DEFINED HIP_PATH)
       if(NOT DEFINED ENV{HIP_PATH})
-          set(HIP_PATH "/opt/rocm/hip" CACHE PATH "Path to HIP installation")
+          message(FATAL_ERROR "GPU_API=HIP requires HIP_PATH to be defined")
       else()
           set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to HIP installation")
       endif()
@@ -261,6 +261,8 @@ elseif(GPU_API STREQUAL "HIP")
 
   if(HIP_PLATFORM STREQUAL "hcc" OR HIP_PLATFORM STREQUAL "amd")
     set(HIP_ARCH "gfx906" CACHE STRING "HIP target architecture")
+  elseif(HIP_PLATFORM STREQUAL "spirv")
+    set(HIP_ARCH "spirv" CACHE STRING "HIP target architecture")
   elseif(HIP_PLATFORM STREQUAL "nvcc")
     find_package(CUDA REQUIRED)
     set(HIP_ARCH "sm_50" CACHE STRING "HIP primary CUDA architecture (e.g. sm_60)")
@@ -321,7 +323,15 @@ elseif(GPU_API STREQUAL "HIP")
     set(CUBIN_FILE   "${LAMMPS_LIB_BINARY_DIR}/gpu/${CU_NAME}.cubin")
     set(CUBIN_H_FILE "${LAMMPS_LIB_BINARY_DIR}/gpu/${CU_NAME}_cubin.h")
 
-    if(HIP_PLATFORM STREQUAL "hcc" OR HIP_PLATFORM STREQUAL "amd")
+    if(HIP_PLATFORM STREQUAL "spirv")
+        configure_file(${CU_FILE} ${CU_CPP_FILE} COPYONLY)
+
+        add_custom_command(OUTPUT ${CUBIN_FILE}
+          VERBATIM COMMAND ${HIP_HIPCC_EXECUTABLE} -c -O3 -DUSE_HIP -D_${GPU_PREC_SETTING} -DLAMMPS_${LAMMPS_SIZES} -I${LAMMPS_LIB_SOURCE_DIR}/gpu -o ${CUBIN_FILE} ${CU_CPP_FILE}
+          DEPENDS ${CU_CPP_FILE}
+          COMMENT "Gerating ${CU_NAME}.cubin")
+
+    elseif(HIP_PLATFORM STREQUAL "hcc" OR HIP_PLATFORM STREQUAL "amd")
         configure_file(${CU_FILE} ${CU_CPP_FILE} COPYONLY)
 
         if(HIP_COMPILER STREQUAL "clang")
diff --git a/lib/gpu/lal_pre_cuda_hip.h b/lib/gpu/lal_pre_cuda_hip.h
index 47a005b998..f6ab1b5b6b 100644
--- a/lib/gpu/lal_pre_cuda_hip.h
+++ b/lib/gpu/lal_pre_cuda_hip.h
@@ -30,7 +30,7 @@
 // -------------------------------------------------------------------------
 
 
-#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
 #define CONFIG_ID 303
 #define SIMD_SIZE 64
 #else
@@ -112,7 +112,7 @@
 //                         KERNEL MACROS - TEXTURES
 // -------------------------------------------------------------------------
 
-#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
 #define _texture(name, type)  __device__ type* name
 #define _texture_2d(name, type)  __device__ type* name
 #else
@@ -135,8 +135,8 @@
     ans=__hiloint2double(qt.y, qt.x);                    \
   }
   #else
-  #define fetch4(ans,i,pos_tex) ans=tex1Dfetch(pos_tex, i);
-  #define fetch(ans,i,q_tex) ans=tex1Dfetch(q_tex,i);
+  #define fetch4(ans,i,pos_tex) tex1Dfetch(&ans, pos_tex, i);
+  #define fetch(ans,i,q_tex) tex1Dfetch(&ans, q_tex,i);
   #endif
 #else
   #define fetch4(ans,i,x) ans=x[i]
@@ -152,7 +152,7 @@
   #define mu_tex mu_
 #endif
 
-#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
 
 #undef fetch4
 #undef fetch
@@ -209,7 +209,7 @@
 #endif
 #endif
 
-#if defined(CUDA_PRE_NINE) || defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
+#if defined(CUDA_PRE_NINE) || defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
 
   #ifdef _SINGLE_SINGLE
     #define shfl_down __shfl_down