Merge pull request #3365 from pvelesko/chip_spv
Add CMake and GNU make support for HIP via CHIP-SPV
This commit is contained in:
Axel Kohlmeyer
@ -1,6 +1,9 @@
|
||||
# /* ----------------------------------------------------------------------
|
||||
# Generic Linux Makefile for HIP
|
||||
# - export HIP_PLATFORM=amd (or nvcc) before execution
|
||||
# - export HIP_PATH=/path/to/HIP/install path to the HIP implementation
|
||||
# such as hipamd or CHIP-SPV.
|
||||
# - export HIP_PLATFORM=<amd/nvcc/spirv> specify the HIP platform to use.
|
||||
# Optional. If not set, will be determined by ${HIP_PATH}/bin/hipconfig.
|
||||
# - change HIP_ARCH for your GPU
|
||||
# ------------------------------------------------------------------------- */
|
||||
|
||||
@ -20,41 +23,62 @@ HIP_OPTS = -O3
|
||||
HIP_HOST_OPTS = -Wno-deprecated-declarations -fopenmp
|
||||
HIP_HOST_INCLUDE =
|
||||
|
||||
ifndef HIP_PATH
|
||||
$(error HIP_PATH is not set)
|
||||
endif
|
||||
|
||||
ifndef HIP_PLATFORM
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
|
||||
endif
|
||||
HIP_COMPILER=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
|
||||
|
||||
# use device sort
|
||||
# requires linking with hipcc and hipCUB + (rocPRIM or CUB for AMD or Nvidia respectively)
|
||||
ifneq (spirv,$(HIP_PLATFORM))
|
||||
# hipCUB not aviable for CHIP-SPV
|
||||
HIP_HOST_OPTS += -DUSE_HIP_DEVICE_SORT
|
||||
endif
|
||||
# path to cub
|
||||
HIP_HOST_INCLUDE += -I./
|
||||
# path to hipcub
|
||||
HIP_HOST_INCLUDE += -I$(HIP_PATH)/../include
|
||||
|
||||
ifeq (amd,$(HIP_PLATFORM))
|
||||
# newer version of ROCm (5.1+) require c++14 for rocprim
|
||||
HIP_OPTS += -std=c++14
|
||||
# newer version of ROCm (5.1+) require c++14 for rocprim
|
||||
HIP_OPTS += -std=c++14
|
||||
endif
|
||||
|
||||
# use mpi
|
||||
HIP_HOST_OPTS += -DMPI_GERYON -DUCL_NO_EXIT
|
||||
# this settings should match LAMMPS Makefile
|
||||
MPI_COMP_OPTS = $(shell mpicxx --showme:compile)
|
||||
# automatic flag detection for OpenMPI
|
||||
ifeq ($(shell mpicxx --showme:compile >/dev/null 2>&1; echo $$?), 0)
|
||||
MPI_COMP_OPTS = $(shell mpicxx --showme:compile) -DOMPI_SKIP_MPICXX=1
|
||||
MPI_LINK_OPTS = $(shell mpicxx --showme:link)
|
||||
|
||||
HIP_PATH ?= $(wildcard /opt/rocm/hip)
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
|
||||
HIP_COMPILER=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
|
||||
# automatic flag detection for MPICH
|
||||
else ifeq ($(shell mpicxx -compile_info >/dev/null 2>&1; echo $$?),0)
|
||||
MPI_COMP_OPTS = $(filter -I%,$(shell mpicxx -compile_info)) -DMPICH_IGNORE_CXX_SEEK
|
||||
MPI_LINK_OPTS = $(filter -Wl%,$(shell mpicxx -link_info)) $(filter -L%,$(shell mpicxx -link_info)) $(filter -l%,$(shell mpicxx -link_info))
|
||||
# for other MPI libs: must set flags manually, if needed
|
||||
else
|
||||
MPI_COMP_OPTS =
|
||||
MPI_LINK_OPTS =
|
||||
endif
|
||||
|
||||
ifeq (hcc,$(HIP_PLATFORM))
|
||||
# possible values: gfx803,gfx900,gfx906
|
||||
HIP_ARCH = gfx906
|
||||
# possible values: gfx803,gfx900,gfx906
|
||||
HIP_ARCH = gfx906
|
||||
else ifeq (amd,$(HIP_PLATFORM))
|
||||
# possible values: gfx803,gfx900,gfx906
|
||||
HIP_ARCH = gfx906
|
||||
# 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] \
|
||||
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] \
|
||||
-gencode arch=compute_50,code=[sm_50,compute_50] -gencode arch=compute_52,code=[sm_52,compute_52] -gencode arch=compute_53,code=[sm_53,compute_53]\
|
||||
-gencode arch=compute_60,code=[sm_60,compute_60] -gencode arch=compute_61,code=[sm_61,compute_61] -gencode arch=compute_62,code=[sm_62,compute_62]\
|
||||
-gencode arch=compute_70,code=[sm_70,compute_70] -gencode arch=compute_72,code=[sm_72,compute_72] -gencode arch=compute_75,code=[sm_75,compute_75]
|
||||
else ifeq (spirv,$(HIP_PLATFORM))
|
||||
HIP_ARCH = spirv
|
||||
endif
|
||||
|
||||
BIN_DIR = .
|
||||
@ -71,7 +95,15 @@ BSH = /bin/sh
|
||||
HIP_OPTS += -DUSE_HIP $(HIP_PRECISION)
|
||||
HIP_GPU_OPTS += $(HIP_OPTS) -I./
|
||||
|
||||
ifeq (clang,$(HIP_COMPILER))
|
||||
ifeq (spirv,$(HIP_PLATFORM))
|
||||
HIP_HOST_OPTS += -fPIC
|
||||
HIP_GPU_CC = $(HIP_PATH)/bin/hipcc -c
|
||||
HIP_GPU_OPTS_S =
|
||||
HIP_GPU_OPTS_E =
|
||||
HIP_KERNEL_SUFFIX = .cpp
|
||||
HIP_LIBS_TARGET =
|
||||
export HCC_AMDGPU_TARGET := $(HIP_ARCH)
|
||||
else ifeq (clang,$(HIP_COMPILER))
|
||||
HIP_HOST_OPTS += -fPIC
|
||||
HIP_GPU_CC = $(HIP_PATH)/bin/hipcc --genco
|
||||
HIP_GPU_OPTS_S = --offload-arch=$(HIP_ARCH)
|
||||
|
||||
@ -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
|
||||
@ -134,9 +134,12 @@
|
||||
int2 qt = tex1Dfetch(q_tex,i); \
|
||||
ans=__hiloint2double(qt.y, qt.x); \
|
||||
}
|
||||
#elseif defined(__HIP_PLATFORM_SPIRV__)
|
||||
#define fetch4(ans,i,pos_tex) tex1Dfetch(&ans, pos_tex, i);
|
||||
#define fetch(ans,i,q_tex) tex1Dfetch(&ans, q_tex,i);
|
||||
#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) ans=tex1Dfetch(pos_tex, i);
|
||||
#define fetch(ans,i,q_tex) ans=tex1Dfetch(q_tex,i);
|
||||
#endif
|
||||
#else
|
||||
#define fetch4(ans,i,x) ans=x[i]
|
||||
@ -152,7 +155,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 +212,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
|
||||
|
||||
Reference in New Issue
Block a user