From 66eea6dedea0ce47c54ac6e9f7f4f960f5f78ef1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ole=20Sch=C3=BCtt?= Date: Fri, 11 Mar 2022 23:34:11 +0100 Subject: [PATCH] Fix HIP/ROCm compilation some more --- src/offload/offload_library.c | 3 --- src/offload/offload_operations.h | 2 +- src/pw/gpu/pw_gpu_kernels.h | 8 ++++---- src/pw/gpu/pw_gpu_kernels_cuda.cu | 10 +++++----- tools/toolchain/scripts/generate_arch_files.sh | 14 +++++++++++--- 5 files changed, 21 insertions(+), 16 deletions(-) diff --git a/src/offload/offload_library.c b/src/offload/offload_library.c index 8364d500dd..895dd21c1f 100644 --- a/src/offload/offload_library.c +++ b/src/offload/offload_library.c @@ -16,11 +16,8 @@ #if defined(__OFFLOAD_PROFILING) #if defined(__OFFLOAD_CUDA) -#include -#include #include #elif defined(__OFFLOAD_HIP) && defined(__HIP_PLATFORM_AMD__) -#include #include #endif #endif diff --git a/src/offload/offload_operations.h b/src/offload/offload_operations.h index 90a44ebdfa..98251298c2 100644 --- a/src/offload/offload_operations.h +++ b/src/offload/offload_operations.h @@ -18,7 +18,7 @@ #if defined(__OFFLOAD_CUDA) #include #elif defined(__OFFLOAD_HIP) -#include +#include #endif #ifdef __cplusplus diff --git a/src/pw/gpu/pw_gpu_kernels.h b/src/pw/gpu/pw_gpu_kernels.h index dd67b1e2de..b9b498ceff 100644 --- a/src/pw/gpu/pw_gpu_kernels.h +++ b/src/pw/gpu/pw_gpu_kernels.h @@ -16,21 +16,21 @@ extern "C" { /******************************************************************************* * \brief Launcher for pw_real_to_complex kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_real_to_complex(const double *din, double *zout, const int ngpts, offloadStream_t stream); /******************************************************************************* * \brief Launcher for pw_complex_to_real kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_complex_to_real(const double *zin, double *dout, const int ngpts, offloadStream_t stream); /******************************************************************************* * \brief Launcher for pw_gather_z kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_gather_z(double *pwcc, const double *c, const double scale, const int ngpts, const int *ghatmap, @@ -38,7 +38,7 @@ void pw_gpu_launch_gather_z(double *pwcc, const double *c, const double scale, /******************************************************************************* * \brief Launcher for pw_scatter_z kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_scatter_z(double *c, const double *pwcc, const double scale, const int ngpts, const int nmaps, diff --git a/src/pw/gpu/pw_gpu_kernels_cuda.cu b/src/pw/gpu/pw_gpu_kernels_cuda.cu index 97207776d6..6626c8a3cf 100644 --- a/src/pw/gpu/pw_gpu_kernels_cuda.cu +++ b/src/pw/gpu/pw_gpu_kernels_cuda.cu @@ -28,10 +28,10 @@ __global__ void pw_real_to_complex(const double *din, double *zout, /******************************************************************************* * \brief Launcher for pw_real_to_complex kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_real_to_complex(const double *din, double *zout, - const int ngpts, cudaStream_t stream) { + const int ngpts, offloadStream_t stream) { const int threadsPerBlock = 1024; const int numBlocks = (ngpts + threadsPerBlock - 1) / threadsPerBlock; pw_real_to_complex<<>>(din, zout, @@ -58,7 +58,7 @@ __global__ void pw_complex_to_real(const double *zin, double *dout, /******************************************************************************* * \brief Launcher for pw_complex_to_real kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_complex_to_real(const double *zin, double *dout, const int ngpts, offloadStream_t stream) { @@ -83,7 +83,7 @@ __global__ void pw_gather_z(double *pwcc, const double *c, const double scale, /******************************************************************************* * \brief Launcher for pw_gather_z kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_gather_z(double *pwcc, const double *c, const double scale, const int ngpts, const int *ghatmap, @@ -114,7 +114,7 @@ __global__ void pw_scatter_z(double *c, const double *pwcc, const double scale, /******************************************************************************* * \brief Launcher for pw_scatter_z kernel. - * \author Ole Schütt + * \author Ole Schuett ******************************************************************************/ void pw_gpu_launch_scatter_z(double *c, const double *pwcc, const double scale, const int ngpts, const int nmaps, diff --git a/tools/toolchain/scripts/generate_arch_files.sh b/tools/toolchain/scripts/generate_arch_files.sh index 2685b4cef8..c722d645de 100755 --- a/tools/toolchain/scripts/generate_arch_files.sh +++ b/tools/toolchain/scripts/generate_arch_files.sh @@ -161,17 +161,21 @@ if [ "${ENABLE_HIP}" = __TRUE__ ] && [ "${GPUVER}" != no ]; then add_lib_from_paths HIP_LDFLAGS "libhipfft.*" $LIB_PATHS PLATFORM_FLAGS='' - HIP_INCLUDES="-I${ROCM_PATH}/hip/include -I${ROCM_PATH}/hipblas/include -I${ROCM_PATH}/include" + HIP_INCLUDES="-I${ROCM_PATH}/include" case "${GPUVER}" in Mi50) check_lib -lamdhip64 "hip" add_lib_from_paths HIP_LDFLAGS "libamdhip64.*" $LIB_PATHS + check_lib -lhipfft "hip" + add_lib_from_paths HIP_LDFLAGS "libhipfft.*" $LIB_PATHS + check_lib -lrocblas "hip" + add_lib_from_paths HIP_LDFLAGS "librocblas.*" $LIB_PATHS check_lib -lroctx64 "hip" add_lib_from_paths HIP_LDFLAGS "libroctx64.*" $LIB_PATHS check_lib -lroctracer64 "hip" add_lib_from_paths HIP_LDFLAGS "libroctracer64.*" $LIB_PATHS HIP_FLAGS+="-fPIE -D__HIP_PLATFORM_AMD__ -g --offload-arch=gfx906 -O3 --std=c++11 \$(DFLAGS)" - LIBS+=" IF_HIP(-lhipblas -lamdhip64 IF_DEBUG(-lroctx64 -lroctracer64|)|)" + LIBS+=" IF_HIP(-lamdhip64 -lhipfft -lhipblas -lrocblas IF_DEBUG(-lroctx64 -lroctracer64|)|)" PLATFORM_FLAGS='-D__HIP_PLATFORM_AMD__' DFLAGS+=' IF_HIP(-D__GRID_HIP -D__HIP_PLATFORM_AMD__ -D__PW_HIP -D__PW_GPU IF_DEBUG(-D__OFFLOAD_PROFILING|)|) -D__DBCSR_ACC' CXXFLAGS+=" -fopenmp -std=c++11" @@ -179,12 +183,16 @@ if [ "${ENABLE_HIP}" = __TRUE__ ] && [ "${GPUVER}" != no ]; then Mi100) check_lib -lamdhip64 "hip" add_lib_from_paths HIP_LDFLAGS "libamdhip64.*" $LIB_PATHS + check_lib -lhipfft "hip" + add_lib_from_paths HIP_LDFLAGS "libhipfft.*" $LIB_PATHS + check_lib -lrocblas "hip" + add_lib_from_paths HIP_LDFLAGS "librocblas.*" $LIB_PATHS check_lib -lroctx64 "hip" add_lib_from_paths HIP_LDFLAGS "libroctx64.*" $LIB_PATHS check_lib -lroctracer64 "hip" add_lib_from_paths HIP_LDFLAGS "libroctracer64.*" $LIB_PATHS HIP_FLAGS+="-fPIE -D__HIP_PLATFORM_AMD__ -g --offload-arch=gfx908 -O3 --std=c++11 \$(DFLAGS)" - LIBS+=" IF_HIP( -lhipblas -lamdhip64 IF_DEBUG(-lroctx64 -lroctracer64|)|)" + LIBS+=" IF_HIP(-lamdhip64 -lhipfft -lhipblas -lrocblas IF_DEBUG(-lroctx64 -lroctracer64|)|)" PLATFORM_FLAGS='-D__HIP_PLATFORM_AMD__ ' DFLAGS+=' IF_HIP(-D__GRID_HIP -D__HIP_PLATFORM_AMD__ -D__PW_HIP -D__PW_GPU IF_DEBUG(-D__OFFLOAD_PROFILING|)|) -D__DBCSR_ACC' CXXFLAGS+=" -fopenmp -std=c++11"