Merge pull request #2692 from arghdos/master

Porting to new default platform for AMD/HIP in ROCm 4.1

cmake/Modules/Packages/GPU.cmake

@@ -218,7 +218,7 @@ elseif(GPU_API STREQUAL "HIP")
 
   if(NOT DEFINED HIP_PLATFORM)
       if(NOT DEFINED ENV{HIP_PLATFORM})
-          set(HIP_PLATFORM "hcc" CACHE PATH "HIP Platform to be used during compilation")
+          set(HIP_PLATFORM "amd" CACHE PATH "HIP Platform to be used during compilation")
       else()
           set(HIP_PLATFORM $ENV{HIP_PLATFORM} CACHE PATH "HIP Platform used during compilation")
       endif()
@@ -226,7 +226,7 @@ elseif(GPU_API STREQUAL "HIP")
 
   set(ENV{HIP_PLATFORM} ${HIP_PLATFORM})
 
-  if(HIP_PLATFORM STREQUAL "hcc")
+  if(HIP_PLATFORM STREQUAL "hcc" OR HIP_PLATFORM STREQUAL "amd")
     set(HIP_ARCH "gfx906" CACHE STRING "HIP target architecture")
   elseif(HIP_PLATFORM STREQUAL "nvcc")
     find_package(CUDA REQUIRED)
@@ -284,7 +284,7 @@ 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")
+    if(HIP_PLATFORM STREQUAL "hcc" OR HIP_PLATFORM STREQUAL "amd")
         configure_file(${CU_FILE} ${CU_CPP_FILE} COPYONLY)
 
         if(HIP_COMPILER STREQUAL "clang")
@@ -386,6 +386,12 @@ elseif(GPU_API STREQUAL "HIP")
 
     target_compile_definitions(hip_get_devices PRIVATE -D__HIP_PLATFORM_HCC__)
     target_include_directories(hip_get_devices PRIVATE ${HIP_ROOT_DIR}/../include)
+  elseif(HIP_PLATFORM STREQUAL "amd")
+    target_compile_definitions(gpu PRIVATE -D__HIP_PLATFORM_AMD__)
+    target_include_directories(gpu PRIVATE ${HIP_ROOT_DIR}/../include)
+
+    target_compile_definitions(hip_get_devices PRIVATE -D__HIP_PLATFORM_AMD__)
+    target_include_directories(hip_get_devices PRIVATE ${HIP_ROOT_DIR}/../include)
   endif()
 
   target_link_libraries(lammps PRIVATE gpu)

doc/src/Build_extras.rst

@@ -125,7 +125,7 @@ CMake build
                                 # default is sm_50
    -D HIP_ARCH=value            # primary GPU hardware choice for GPU_API=hip
                                 # value depends on selected HIP_PLATFORM
-                                # default is 'gfx906' for HIP_PLATFORM=hcc and 'sm_50' for HIP_PLATFORM=nvcc
+                                # default is 'gfx906' for HIP_PLATFORM=amd and 'sm_50' for HIP_PLATFORM=nvcc
    -D HIP_USE_DEVICE_SORT=value # enables GPU sorting
                                 # value = yes (default) or no
    -D CUDPP_OPT=value           # use GPU binning on with CUDA (should be off for modern GPUs)
@@ -169,17 +169,24 @@ desired, you can set :code:`USE_STATIC_OPENCL_LOADER` to :code:`no`.
 
 If you are compiling with HIP, note that before running CMake you will have to
 set appropriate environment variables. Some variables such as
-:code:`HCC_AMDGPU_TARGET` or :code:`CUDA_PATH` are necessary for :code:`hipcc`
+:code:`HCC_AMDGPU_TARGET` (for ROCm <= 4.0) or :code:`CUDA_PATH` are necessary for :code:`hipcc`
 and the linker to work correctly.
 
 .. code:: bash
 
-   # AMDGPU target
+   # AMDGPU target (ROCm <= 4.0)
    export HIP_PLATFORM=hcc
    export HCC_AMDGPU_TARGET=gfx906
    cmake -D PKG_GPU=on -D GPU_API=HIP -D HIP_ARCH=gfx906 -D CMAKE_CXX_COMPILER=hipcc ..
    make -j 4
 
+.. code:: bash
+
+   # AMDGPU target (ROCm >= 4.1)
+   export HIP_PLATFORM=amd
+   cmake -D PKG_GPU=on -D GPU_API=HIP -D HIP_ARCH=gfx906 -D CMAKE_CXX_COMPILER=hipcc ..
+   make -j 4
+
 .. code:: bash
 
    # CUDA target (not recommended, use GPU_ARCH=cuda)

lib/gpu/Makefile.hip

@@ -1,6 +1,6 @@
 # /* ----------------------------------------------------------------------
 #  Generic Linux Makefile for HIP
-#     - export HIP_PLATFORM=hcc (or nvcc) before execution
+#     - export HIP_PLATFORM=amd (or nvcc) before execution
 #     - change HIP_ARCH for your GPU
 # ------------------------------------------------------------------------- */
 
@@ -42,6 +42,10 @@ ifeq (hcc,$(HIP_PLATFORM))
 	HIP_OPTS  += -ffast-math
 	# possible values: gfx803,gfx900,gfx906
 	HIP_ARCH = gfx906
+else ifeq (amd,$(HIP_PLATFORM))
+	HIP_OPTS  += -ffast-math
+	# possible values: gfx803,gfx900,gfx906
+	HIP_ARCH = gfx906
 else ifeq (nvcc,$(HIP_PLATFORM))
 	HIP_OPTS  += --use_fast_math
 	HIP_ARCH = -gencode arch=compute_30,code=[sm_30,compute_30] -gencode arch=compute_32,code=[sm_32,compute_32] -gencode arch=compute_35,code=[sm_35,compute_35] \

lib/gpu/README

@@ -212,8 +212,8 @@ additionally requires cub (https://nvlabs.github.io/cub). Download and
 extract the cub directory to lammps/lib/gpu/ or specify an appropriate
 path in lammps/lib/gpu/Makefile.hip.
 2. In Makefile.hip it is possible to specify the target platform via
-export HIP_PLATFORM=hcc or HIP_PLATFORM=nvcc as well as the target
+export HIP_PLATFORM=amd (ROCm >= 4.1), HIP_PLATFORM=hcc (ROCm <= 4.0)
-architecture (gfx803, gfx900, gfx906 etc.)
+or HIP_PLATFORM=nvcc as well as the target architecture (gfx803, gfx900, gfx906 etc.)
 3. If your MPI implementation does not support `mpicxx --showme` command,
 it is required to specify the corresponding MPI compiler and linker flags
 in lammps/lib/gpu/Makefile.hip and in lammps/src/MAKE/OPTIONS/Makefile.hip.
@@ -278,4 +278,3 @@ and
 Brown, W.M., Masako, Y. Implementing Molecular Dynamics on Hybrid High
 Performance Computers - Three-Body Potentials. Computer Physics Communications.
 2013. 184: p. 2785–2793.
-

lib/gpu/lal_pre_cuda_hip.h

@@ -30,7 +30,7 @@
 // -------------------------------------------------------------------------
 
 
-#ifdef __HIP_PLATFORM_HCC__
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
 #define CONFIG_ID 303
 #define SIMD_SIZE 64
 #else
@@ -161,7 +161,7 @@
 //                         KERNEL MACROS - TEXTURES
 // -------------------------------------------------------------------------
 
-#ifdef __HIP_PLATFORM_HCC__
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
 #define _texture(name, type)  __device__ type* name
 #define _texture_2d(name, type)  __device__ type* name
 #else
@@ -201,7 +201,7 @@
   #define mu_tex mu_
 #endif
 
-#ifdef __HIP_PLATFORM_HCC__
+#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
 
 #undef fetch4
 #undef fetch
@@ -266,7 +266,7 @@ typedef struct _double4 double4;
 #endif
 #endif
 
-#if defined(CUDA_PRE_NINE) || defined(__HIP_PLATFORM_HCC__)
+#if defined(CUDA_PRE_NINE) || defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
 
   #ifdef _SINGLE_SINGLE
     #define shfl_down __shfl_down

tools/singularity/ubuntu18.04_amd_rocm.def

@@ -94,7 +94,7 @@ From: ubuntu:18.04
     ###########################################################################
 
     export PATH=$PATH:/opt/rocm/bin:/opt/rocm/profiler/bin:/opt/rocm/opencl/bin/x86_64
-    git clone -b rocm-3.7.x https://github.com/ROCmSoftwarePlatform/hipCUB.git
+    git clone -b rocm-4.1.x https://github.com/ROCmSoftwarePlatform/hipCUB.git
     mkdir hipCUB/build
     cd hipCUB/build
     CXX=hipcc cmake -D BUILD_TEST=off ..

tools/singularity/ubuntu20.04_amd_rocm.def

@@ -2,7 +2,7 @@ BootStrap: docker
 From: ubuntu:20.04
 
 %environment
-    export PATH=/usr/lib/ccache:/usr/local/cuda-11.0/bin:${PATH}:/opt/rocm/bin:/opt/rocm/profiler/bin:/opt/rocm/opencl/bin/x86_64
+    export PATH=/usr/lib/ccache:${PATH}:/opt/rocm/bin:/opt/rocm/profiler/bin:/opt/rocm/opencl/bin/x86_64
 %post
     export DEBIAN_FRONTEND=noninteractive
     apt-get update
@@ -90,7 +90,7 @@ From: ubuntu:20.04
     ###########################################################################
 
     export PATH=$PATH:/opt/rocm/bin:/opt/rocm/profiler/bin:/opt/rocm/opencl/bin/x86_64
-    git clone -b rocm-3.7.x https://github.com/ROCmSoftwarePlatform/hipCUB.git
+    git clone -b rocm-4.1.x https://github.com/ROCmSoftwarePlatform/hipCUB.git
     mkdir hipCUB/build
     cd hipCUB/build
     CXX=hipcc cmake -D BUILD_TEST=off ..