comp-chem/cp2k
1
0
Fork 0
mirror of https://github.com/cp2k/cp2k.git synced 2026-07-29 06:35:28 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

Fix HIP/ROCm compilation some more

Browse source
This commit is contained in:
Ole Schütt 2022-03-11 23:34:11 +01:00 • committed by Ole Schütt
parent e563fa0c91
commit 66eea6dede
5 changed files with 21 additions and 16 deletions
Download patch file Download diff file Expand all files Collapse all files

3
src/offload/offload_library.c
View file

@ -16,11 +16,8 @@
#if defined(__OFFLOAD_PROFILING)
#if defined(__OFFLOAD_CUDA)
#include <cuda.h>
#include <cuda_runtime.h>
#include <nvToolsExt.h>
#elif defined(__OFFLOAD_HIP) && defined(__HIP_PLATFORM_AMD__)
#include <hip/hip_runtime_api.h>
#include <roctracer/roctx.h>
#endif
#endif

2
src/offload/offload_operations.h
View file

@ -18,7 +18,7 @@
#if defined(__OFFLOAD_CUDA)
#include <cuda_runtime.h>
#elif defined(__OFFLOAD_HIP)
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>
#endif
#ifdef __cplusplus

8
src/pw/gpu/pw_gpu_kernels.h
View file

@ -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,

10
src/pw/gpu/pw_gpu_kernels_cuda.cu
View file

@ -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<<<numBlocks, threadsPerBlock, 0, stream>>>(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,

14
tools/toolchain/scripts/generate_arch_files.sh
View file

@ -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"