From 4c2fb7a4312501d59ab530a09a5c5054272d53bf Mon Sep 17 00:00:00 2001 From: Nick Curtis Date: Thu, 1 Apr 2021 09:43:13 -0500 Subject: [PATCH 1/2] Porting to new default platform for AMD/HIP in ROCm 4.1 --- cmake/Modules/Packages/GPU.cmake | 12 +++++++++--- doc/src/Build_extras.rst | 13 ++++++++++--- lib/gpu/Makefile.hip | 6 +++++- lib/gpu/README | 5 ++--- lib/gpu/lal_pre_cuda_hip.h | 8 ++++---- 5 files changed, 30 insertions(+), 14 deletions(-) diff --git a/cmake/Modules/Packages/GPU.cmake b/cmake/Modules/Packages/GPU.cmake index 9aa917144b..68d74ea42e 100644 --- a/cmake/Modules/Packages/GPU.cmake +++ b/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") @@ -381,6 +381,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) diff --git a/doc/src/Build_extras.rst b/doc/src/Build_extras.rst index 2081dc4bcd..3af018c656 100644 --- a/doc/src/Build_extras.rst +++ b/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) diff --git a/lib/gpu/Makefile.hip b/lib/gpu/Makefile.hip index dbdef433ec..a736988596 100644 --- a/lib/gpu/Makefile.hip +++ b/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] \ diff --git a/lib/gpu/README b/lib/gpu/README index dfffe11b81..eb22839a59 100644 --- a/lib/gpu/README +++ b/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 -architecture (gfx803, gfx900, gfx906 etc.) +export HIP_PLATFORM=amd (ROCm >= 4.1), HIP_PLATFORM=hcc (ROCm <= 4.0) +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. - diff --git a/lib/gpu/lal_pre_cuda_hip.h b/lib/gpu/lal_pre_cuda_hip.h index d37b4a94c2..dfb6229bed 100644 --- a/lib/gpu/lal_pre_cuda_hip.h +++ b/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 From 160f2cc6307cb5c7e7b3179ca93e0d01de7b8b9f Mon Sep 17 00:00:00 2001 From: Richard Berger Date: Fri, 2 Apr 2021 16:11:23 -0400 Subject: [PATCH 2/2] Update ROCm container definitions --- tools/singularity/ubuntu18.04_amd_rocm.def | 2 +- tools/singularity/ubuntu20.04_amd_rocm.def | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/tools/singularity/ubuntu18.04_amd_rocm.def b/tools/singularity/ubuntu18.04_amd_rocm.def index bb04c738c4..7b06970110 100644 --- a/tools/singularity/ubuntu18.04_amd_rocm.def +++ b/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 .. diff --git a/tools/singularity/ubuntu20.04_amd_rocm.def b/tools/singularity/ubuntu20.04_amd_rocm.def index 28d57be341..4eee97ffec 100644 --- a/tools/singularity/ubuntu20.04_amd_rocm.def +++ b/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 ..