From aa4adac34cf79714cddcdea931a90aba3810da90 Mon Sep 17 00:00:00 2001 From: Hans Pabst Date: Fri, 15 Mar 2024 16:41:35 +0100 Subject: [PATCH] DBM: introduced OpenCL support (#3321) --- INSTALL.md | 21 +-- src/CMakeLists.txt | 3 +- src/dbm/.gitignore | 2 + src/dbm/Makefile | 183 +++++++++++++++++------ src/dbm/dbm_multiply.c | 41 ++++- src/dbm/dbm_multiply_gpu.c | 13 +- src/dbm/dbm_multiply_gpu.h | 1 + src/dbm/dbm_multiply_gpu_kernel.cu | 11 +- src/dbm/dbm_multiply_gpu_kernel.h | 10 +- src/dbm/dbm_multiply_opencl.c | 119 +++++++++++++++ src/dbm/dbm_multiply_opencl.cl | 77 ++++++++++ tools/precommit/README.md | 2 +- tools/precommit/check_file_properties.py | 3 +- tools/precommit/precommit.py | 2 +- 14 files changed, 402 insertions(+), 86 deletions(-) create mode 100644 src/dbm/.gitignore create mode 100644 src/dbm/dbm_multiply_opencl.c create mode 100644 src/dbm/dbm_multiply_opencl.cl diff --git a/INSTALL.md b/INSTALL.md index e770b7c479..8a550f2909 100644 --- a/INSTALL.md +++ b/INSTALL.md @@ -400,37 +400,38 @@ serving DBM/DBT and other libraries depends on DBCSR's OpenCL backend. necessary if CP2K/DBCSR is compiled from source. Please note, some implementations ship with outdated OpenCL headers which can prevent using latest features (if an application discovers such features only at compile-time). When building from source, for instance `libOpenCL.so` is - sufficient (ICD loader) at link-time. However, an Installable Client Driver (ICD) is finally + sufficient at link-time (ICD loader). However, an Installable Client Driver (ICD) is finally necessary at runtime. - Nvidia CUDA, AMD HIP, and Intel OneAPI are fully equipped with an OpenCL runtime (if - `opencl-headers` package is not installed, CPATH can be needed to point into such an + `opencl-headers` package is not installed, CPATH can be needed to point into the former installation, similarly `LIBRARY_PATH` for finding `libOpenCL.so` at link-time). Installing a minimal or stand-alone OpenCL is also possible, e.g., following the instructions for Debian (or Ubuntu) as given for every [release](https://github.com/intel/compute-runtime/releases) of the [Intel Compute Runtime](https://github.com/intel/compute-runtime). - - The environment variable `ACC_OPENCL_VERBOSE=2` prints information at runtime of CP2K about - kernels generated and used (and thereby checks the installation). + - The environment variable `ACC_OPENCL_VERBOSE` prints information at runtime of CP2K about + kernels generated (`ACC_OPENCL_VERBOSE=2`) or executed (`ACC_OPENCL_VERBOSE=3`) which can be + used to check an installation. - Building CP2K with OpenCL-based DBCSR - CP2K's toolchain supports `--enable-opencl` to select DBCSR's OpenCL backend. This can be - combined with `--enable-cuda` (`--gpu-ver` is then imposed) to use a GPU for CP2K's grid and - DBM/DBT components (no OpenCL support yet). + combined with `--enable-cuda` (`--gpu-ver` is then imposed) to use a GPU for CP2K's GRID and PW + components (no OpenCL support yet) with DBM's CUDA implementation to be preferred. - For manually writing an ARCH-file, add `-D__OPENCL` and `-D__DBCSR_ACC` to `CFLAGS` and add `-lOpenCL` to the `LIBS` variable, i.e., `OFFLOAD_CC` and `OFFLOAD_FLAGS` can duplicate `CC` and `CFLAGS` (no special offload compiler needed). Please also set `OFFLOAD_TARGET = opencl` to enable the OpenCL backend in DBCSR. For OpenCL, it is not necessary to specify a GPU version (e.g., `GPUVER = V100` would map/limit to `exts/dbcsr/src/acc/opencl/smm/params/tune_multiply_V100.csv`). In fact, `GPUVER` limits tuned - parameters to the specified GPU, and by default all tuned parameters are embedded + parameters to the specified GPU, whereas by default all tuned parameters are embedded (`exts/dbcsr/src/acc/opencl/smm/params/*.csv`) and applied at runtime. If auto-tuned parameters are not available for DBCSR, well-chosen defaults will be used to populate kernels at runtime. - - Refer to the toolchain method (above) for an ARCH-file that blends, e.g., OpenCL and CUDA. - Auto-tuned parameters are embedded into the binary, i.e., CP2K does not rely on a hard-coded location. Setting `OPENCL_LIBSMM_SMM_PARAMS=/path/to/csv-file` environment variable can supply parameters for an already built application, or `OPENCL_LIBSMM_SMM_PARAMS=0` can disable using tuned parameters. Refer to on how to tune kernels (parameters). -- Building CP2K with OpenCL-based DBM/DBT libraries +- Building CP2K with OpenCL-based DBM library - For manually writing an ARCH-file, add `-D__OFFLOAD_OPENCL` to `CFLAGS` in addition to following - above instructions for "Building CP2K with OpenCL-based DBCSR". + above instructions for "Building CP2K with OpenCL-based DBCSR". An additional Makefile rule can + be necessary to transform OpenCL code into a ressource header file. ### 2y. matrix-matrix multiplication offloading on GPU using SPLA diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 828f2b7335..3200fae184 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -1372,7 +1372,8 @@ set(CP2K_GRID_SRCS_HIP grid/hip/grid_hip_context.cu) set(CP2K_DBM_SRCS_GPU dbm/dbm_multiply_gpu_kernel.cu) -set(CP2K_DBM_SRCS_GPU_C dbm/dbm_multiply_gpu.c) +set(CP2K_DBM_SRCS_GPU_C dbm/dbm_multiply_opencl.c dbm/dbm_multiply_gpu.c) +set(CP2K_DBM_SRCS_OPENCL dbm/dbm_multiply_opencl.cl) set(CP2K_PW_SRCS_GPU pw/gpu/pw_gpu_kernels.cu) diff --git a/src/dbm/.gitignore b/src/dbm/.gitignore new file mode 100644 index 0000000000..e8d7c52bcf --- /dev/null +++ b/src/dbm/.gitignore @@ -0,0 +1,2 @@ +dbm_multiply_opencl.cl.h +dbm_miniapp.x diff --git a/src/dbm/Makefile b/src/dbm/Makefile index 96e7f2f6e4..f7ed467e4c 100644 --- a/src/dbm/Makefile +++ b/src/dbm/Makefile @@ -1,44 +1,5 @@ -.PHONY : all clean - -all: dbm_miniapp.x - -clean: - rm -fv *.o */*.o ../offload/*.o - -realclean: clean - rm -fv *.x - -CFLAGS := -fopenmp -g -O3 -march=native -Wall -Wextra -Wno-vla-parameter - -NVCC := $(shell which nvcc) -NVARCH := sm_70 -NVFLAGS := -g -O3 -lineinfo -arch $(NVARCH) -Wno-deprecated-gpu-targets -Xcompiler "$(CFLAGS)" -D__OFFLOAD_CUDA - -HIPCC := $(shell which hipcc) -HIPARCH := gfx90a -HIPFLAGS := -fPIE -g -O3 --offload-arch=$(HIPARCH) -Wall -Wextra -Werror -I${ROCM_PATH}/include -D__OFFLOAD_HIP -D__HIP_PLATFORM_AMD__ -D__HIP_PLATFORM_AMD__ - -ifneq ($(LIBXSMMROOT),) -LIBS += \ - $(LIBXSMMROOT)/lib/libxsmm.a $(LIBXSMMROOT)/lib/libxsmmext.a \ - -lpthread -lrt -CFLAGS += -D__LIBXSMM -I$(LIBXSMMROOT)/include -endif - -ifneq ($(MKLROOT),) -LIBS += \ - -L$(MKLROOT)/lib/intel64 \ - -Wl,--start-group \ - -lmkl_gf_lp64 \ - -lmkl_core \ - -lmkl_sequential \ - -Wl,--end-group -else -LIBS += -lblas -endif - -LIBS += -fopenmp -ldl -lstdc++ -lc -lm - +PROJHOME := $(abspath $(dir $(lastword $(MAKEFILE_LIST)))) +CP2KHOME := $(abspath $(PROJHOME)/../..) ALL_HEADERS := $(shell find . -name "*.h") $(shell find ../offload/ -name "*.h") ALL_OBJECTS := ../offload/offload_library.o \ dbm_distribution.o \ @@ -51,12 +12,133 @@ ALL_OBJECTS := ../offload/offload_library.o \ dbm_multiply_cpu.o \ dbm_shard.o +# Optimization level +ifeq (,$(filter-out 0,$(DBG))) +OPT ?= 3 +else +OPT ?= 0 +endif + +# Optimization flag derived from OPT flag +OPTFLAG ?= -O$(patsubst O%,%,$(OPT)) + +CFLAGS := -fopenmp -g $(OPTFLAG) -march=native -Wall -Wextra -Wcast-qual + +# Some additional flags when relying on default-CC +ifeq ($(CC),) +CFLAGS += -Wno-vla-parameter +endif + +NVCC := $(shell which nvcc 2>/dev/null) +NVARCH := sm_70 +NVFLAGS := -g $(OPTFLAG) -lineinfo -arch $(NVARCH) -Wno-deprecated-gpu-targets -Xcompiler "$(CFLAGS)" -D__OFFLOAD_CUDA + +HIPCC := $(shell which hipcc 2>/dev/null) +HIPARCH := gfx90a +HIPFLAGS := -fPIE -g $(OPTFLAG) --offload-arch=$(HIPARCH) -Wall -Wextra -Werror -I${ROCM_PATH}/include -D__OFFLOAD_HIP -D__HIP_PLATFORM_AMD__ -D__HIP_PLATFORM_AMD__ + +# prefer NVCC or HIPCC over OpenCL +ifneq ($(NVCC)$(HIPCC),) +NVCC_PATH := $(if $(NVCC),$(wildcard $(dir $(NVCC))/..)) +CUDA_FILE := $(wildcard $(NVCC_PATH)/../cuda/include/cuda.h) +CUDA_PATH := $(if $(CUDA_FILE),$(NVCC_PATH)/../cuda,$(NVCC_PATH)) +CUDA_LIBS := $(if $(wildcard $(CUDA_PATH)/lib64),lib64,lib) +OPENCL := 0 +endif +ifneq ($(OPENCL),0) +OPENCL_OFFLOAD := $(firstword $(wildcard $(CP2KHOME)/lib/*/*/exts/dbcsr/libdbcsr.a)) +OPENCL_GENKRNL := $(PROJHOME)/dbm_multiply_opencl.cl.h +endif + +# Make foundational runtimes available +LIBS += -fopenmp -ldl -lstdc++ -lc -lm + + +# Make BLAS/LAPACK available +ifneq ($(MKLROOT),) +LIBS += \ + -L$(MKLROOT)/lib/intel64 \ + -Wl,--start-group \ + -lmkl_gf_lp64 \ + -lmkl_core \ + -lmkl_sequential \ + -Wl,--end-group +else +LIBS += -lblas +endif + +.PHONY : all clean + +all: dbm_miniapp.x + +clean: + rm -fv *.o ../offload/*.o $(OPENCL_GENKRNL) + +realclean: clean + rm -fv *.x + + +# Enable OpenCL when DBCSR library was prebuilt (assume __DBCSR_ACC). +ifneq ($(OPENCL_OFFLOAD),) +OPENCL_SRC := $(PROJHOME)/dbm_multiply_opencl.cl +OPENCL_CMN := $(wildcard $(CP2KHOME)/exts/dbcsr/src/acc/opencl/common/*.h) +OPENCL_GEN := $(CP2KHOME)/exts/dbcsr/src/acc/opencl/acc_opencl.sh +ALL_HEADERS += $(OPENCL_GENKRNL) + +ALL_OBJECTS += dbm_multiply_gpu.o dbm_multiply_opencl.o +CFLAGS += -I$(CP2KHOME)/exts/dbcsr -D__OFFLOAD_OPENCL -D__DBCSR_ACC +LIBS += $(OPENCL_OFFLOAD) -lgfortran + +ifneq (Darwin,$(shell uname)) +OPENCL_LIB := $(shell ldconfig -p 2>/dev/null | grep -m1 OpenCL | rev | cut -d' ' -f1 | rev) +ifeq (,$(OPENCL_LIB)) +OPENCL_LIB := $(wildcard /usr/lib/x86_64-linux-gnu/libOpenCL.so.1) +endif +ifneq (,$(CUDA_PATH)) +ifeq (,$(wildcard $(OPENCL_INC))) +CFLAGS += -I$(CUDA_PATH)/include +endif +ifeq (,$(wildcard $(OPENCL_LIB))) +LIBS += -L$(CUDA_PATH)/$(CUDA_LIBS) +LIBS += -Wl,-rpath=$(CUDA_PATH)/$(CUDA_LIBS) +endif +else ifneq (,$(wildcard $(OPENCL_ROOT)/include/CL/cl.h)) +ifeq (,$(wildcard $(OPENCL_INC))) +CFLAGS += -I$(OPENCL_ROOT)/include +endif +LIBS += -L$(OPENCL_ROOT)/lib64 +else ifneq (,$(wildcard $(OCL_ROOT)/include/CL/cl.h)) +ifeq (,$(wildcard $(OPENCL_INC))) +CFLAGS += -I$(OCL_ROOT)/include +endif +LIBS += -L$(OCL_ROOT)/lib64 +endif +# OPENCL_INC: directory containing CL/cl.h. +ifneq (,$(wildcard $(OPENCL_INC))) +CFLAGS += -I$(OPENCL_INC) +endif +# OPENCL_LIB: file/library to be linked +ifneq (,$(wildcard $(OPENCL_LIB))) +LIBS += $(OPENCL_LIB) +else +LIBS += -l:libOpenCL.so.1 +endif +else # macOS +LIBS += -framework OpenCL +endif + +$(OPENCL_GENKRNL): $(OPENCL_GEN) $(OPENCL_SRC) $(OPENCL_CMN) + $(OPENCL_GEN) -b 6 -p "" $(OPENCL_SRC) $@ + # Enable Cuda when nvcc compiler is present. -ifneq ($(NVCC),) -LIBS += -lcudart -lcuda -lcublas -L${CUDA_PATH}/lib64 -CFLAGS += -I${CUDA_PATH}/include -D__OFFLOAD_CUDA +else ifneq ($(NVCC),) ALL_OBJECTS += dbm_multiply_gpu.o dbm_multiply_gpu_kernel.o +CFLAGS += -I${CUDA_PATH}/include -D__OFFLOAD_CUDA +LIBS += -lcudart -lcuda -lcublas -L${CUDA_PATH}/$(CUDA_LIBS) +ifneq ($(wildcard $(NVCC_PATH)/../math_libs/$(CUDA_LIBS)),) +LIBS += -L$(NVCC_PATH)/../math_libs/$(CUDA_LIBS) +endif %.o: %.cu $(ALL_HEADERS) cd $(dir $<); $(NVCC) -c $(NVFLAGS) $(notdir $<) @@ -64,15 +146,24 @@ ALL_OBJECTS += dbm_multiply_gpu.o dbm_multiply_gpu_kernel.o # Enable HIP/ROCm when hipcc compiler is present. else ifneq ($(HIPCC),) -LIBS += -L${ROCM_PATH}/lib -lamdhip64 -lhipfft -lhipblas -lrocblas -CFLAGS += -I${ROCM_PATH}/include -D__OFFLOAD_HIP -D__HIP_PLATFORM_AMD__ ALL_OBJECTS += dbm_multiply_gpu.o +CFLAGS += -I${ROCM_PATH}/include -D__OFFLOAD_HIP -D__HIP_PLATFORM_AMD__ +LIBS += -L${ROCM_PATH}/lib -lamdhip64 -lhipfft -lhipblas -lrocblas %.o: %.cu $(ALL_HEADERS) cd $(dir $<); $(HIPCC) -c $(HIPFLAGS) $(notdir $<) endif +# Make LIBXSMM available +ifneq ($(LIBXSMMROOT),) +CFLAGS += -D__LIBXSMM -I$(LIBXSMMROOT)/include +LIBS += \ + $(LIBXSMMROOT)/lib/libxsmm.a $(LIBXSMMROOT)/lib/libxsmmext.a \ + -lpthread -lrt +endif + + # Build with MPI when mpicc compiler is present. ifneq ($(MPICC),) %.o: %.c $(ALL_HEADERS) diff --git a/src/dbm/dbm_multiply.c b/src/dbm/dbm_multiply.c index 8ea934f71c..b042385d17 100644 --- a/src/dbm/dbm_multiply.c +++ b/src/dbm/dbm_multiply.c @@ -6,6 +6,7 @@ /*----------------------------------------------------------------------------*/ #include +#include #include #include @@ -19,11 +20,24 @@ #include "dbm_multiply_internal.h" /******************************************************************************* - * \brief Returns the larger of two given integer (missing from the C standard) + * \brief Returns the larger of two given integer (missing from the C standard). * \author Ole Schuett ******************************************************************************/ static inline int imax(int x, int y) { return (x > y ? x : y); } +/******************************************************************************* + * \brief Updates the min/max of a range of values (initially {INT_MAX, 0}). + * \author Hans Pabst + ******************************************************************************/ +static inline void min_max(int result[2], int value) { + if (value < result[0]) { + result[0] = value; + } + if (result[1] < value) { + result[1] = value; + } +} + /******************************************************************************* * \brief Private routine for computing the max filter threshold for each row. * \author Ole Schuett @@ -110,7 +124,8 @@ static void backend_upload_packs(const dbm_pack_t *pack_a, * \author Ole Schuett ******************************************************************************/ static void backend_process_batch(const int ntasks, dbm_task_t batch[ntasks], - const double alpha, const dbm_pack_t *pack_a, + const int mnk_range[3][2], const double alpha, + const dbm_pack_t *pack_a, const dbm_pack_t *pack_b, const int kshard, dbm_shard_t *shard_c, backend_context_t *ctx) { @@ -118,9 +133,11 @@ static void backend_process_batch(const int ntasks, dbm_task_t batch[ntasks], (void)pack_a; // mark as used (void)pack_b; (void)shard_c; - dbm_multiply_gpu_process_batch(ntasks, batch, alpha, kshard, &ctx->gpu); + dbm_multiply_gpu_process_batch(ntasks, batch, mnk_range, alpha, kshard, + &ctx->gpu); #else - (void)kshard; // mark as used + (void)mnk_range; // mark as used + (void)kshard; (void)ctx; dbm_multiply_cpu_process_batch(ntasks, batch, alpha, pack_a, pack_b, shard_c); #endif @@ -208,6 +225,7 @@ static void multiply_packs(const bool transa, const bool transb, const int ishard = shard_row * nshard_cols + shard_col; dbm_shard_t *shard_c = &matrix_c->shards[ishard]; dbm_task_t batch[MAX_BATCH_SIZE]; + int mnk_range[][2] = {{INT_MAX, 0}, {INT_MAX, 0}, {INT_MAX, 0}}; int ntasks = 0; // Use a merge-join to find pairs of blocks with matching sum indices. @@ -276,15 +294,22 @@ static void multiply_packs(const bool transa, const bool transb, batch[ntasks].offset_c = blk_c->offset; ntasks++; + // track MxN-shape covering an entire batch + min_max(mnk_range[0], m); + min_max(mnk_range[1], n); + min_max(mnk_range[2], k); + if (ntasks == MAX_BATCH_SIZE) { - backend_process_batch(ntasks, batch, alpha, pack_a, pack_b, - ishard, shard_c, ctx); + backend_process_batch(ntasks, batch, mnk_range, alpha, pack_a, + pack_b, ishard, shard_c, ctx); + mnk_range[0][0] = mnk_range[1][0] = mnk_range[2][0] = INT_MAX; + mnk_range[0][1] = mnk_range[1][1] = mnk_range[2][1] = 0; ntasks = 0; } } } - backend_process_batch(ntasks, batch, alpha, pack_a, pack_b, ishard, - shard_c, ctx); + backend_process_batch(ntasks, batch, mnk_range, alpha, pack_a, pack_b, + ishard, shard_c, ctx); } } } diff --git a/src/dbm/dbm_multiply_gpu.c b/src/dbm/dbm_multiply_gpu.c index 6f8a40fe7f..c2ee8a08e7 100644 --- a/src/dbm/dbm_multiply_gpu.c +++ b/src/dbm/dbm_multiply_gpu.c @@ -102,6 +102,7 @@ void dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a, * \author Ole Schuett ******************************************************************************/ void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch, + const int mnk_range[3][2], const double alpha, const int kshard, dbm_multiply_gpu_context_t *ctx) { @@ -123,7 +124,7 @@ void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch, offloadEventCreate(&batch_uploaded); offloadEventRecord(batch_uploaded, shard_c_dev->stream); - // Reallocate shard_c_dev->data if nessecary. + // Reallocate shard_c_dev->data if necessary. if (shard_c_host->data_promised > shard_c_dev->data_allocated) { double *old_data_dev = shard_c_dev->data; shard_c_dev->data_allocated = @@ -138,7 +139,7 @@ void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch, dbm_mempool_free(old_data_dev); } - // Zero new blocks if nessecary. + // Zero new blocks if necessary. if (shard_c_host->data_promised > shard_c_dev->data_size) { const int tail = shard_c_host->data_promised - shard_c_dev->data_size; offloadMemsetAsync(&shard_c_dev->data[shard_c_dev->data_size], 0, @@ -147,9 +148,9 @@ void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch, } // Launch kernel. - dbm_multiply_gpu_launch_kernel(shard_c_dev->stream, alpha, ntasks, batch_dev, - ctx->pack_a_dev.data, ctx->pack_b_dev.data, - shard_c_dev->data); + dbm_multiply_gpu_launch_kernel(shard_c_dev->stream, mnk_range, alpha, ntasks, + batch_dev, ctx->pack_a_dev.data, + ctx->pack_b_dev.data, shard_c_dev->data); OFFLOAD_CHECK(offloadGetLastError()); // Wait for batch to be uploaded before refilling it. @@ -167,7 +168,7 @@ void dbm_multiply_gpu_download_results(dbm_multiply_gpu_context_t *ctx) { #pragma omp parallel for schedule(dynamic) for (int i = 0; i < ctx->nshards; i++) { - // Grow host buffer if nessecary. + // Grow host buffer if necessary. dbm_shard_t *shard_c_host = &ctx->shards_c_host[i]; dbm_shard_allocate_promised_blocks(shard_c_host); diff --git a/src/dbm/dbm_multiply_gpu.h b/src/dbm/dbm_multiply_gpu.h index 63c991a795..dbb4212c38 100644 --- a/src/dbm/dbm_multiply_gpu.h +++ b/src/dbm/dbm_multiply_gpu.h @@ -64,6 +64,7 @@ void dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a, * \author Ole Schuett ******************************************************************************/ void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch, + const int mnk_range[3][2], const double alpha, const int kshard, dbm_multiply_gpu_context_t *ctx); diff --git a/src/dbm/dbm_multiply_gpu_kernel.cu b/src/dbm/dbm_multiply_gpu_kernel.cu index ee8ae741c3..3cf97bbdf0 100644 --- a/src/dbm/dbm_multiply_gpu_kernel.cu +++ b/src/dbm/dbm_multiply_gpu_kernel.cu @@ -199,15 +199,14 @@ __global__ static void process_batch_kernel(const double alpha, * All arguments are assumed to be device pointers. * \author Ole Schuett ******************************************************************************/ -void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, - const double alpha, const int ntasks, - const dbm_task_t *batch, - const double *pack_a_data, - const double *pack_b_data, - double *shard_c_data) { +void dbm_multiply_gpu_launch_kernel( + const offloadStream_t stream, const int mnk_range[3][2], const double alpha, + const int ntasks, const dbm_task_t *batch, const double *pack_a_data, + const double *pack_b_data, double *shard_c_data) { const int nblocks = ntasks; // TODO tune launch parameters. const int threads_per_block = NUM_THREADS; const size_t smem_per_block = 0; + (void)mnk_range; // mark used process_batch_kernel<<>>( alpha, batch, pack_a_data, pack_b_data, shard_c_data); } diff --git a/src/dbm/dbm_multiply_gpu_kernel.h b/src/dbm/dbm_multiply_gpu_kernel.h index 136ed2310e..446fd98e1a 100644 --- a/src/dbm/dbm_multiply_gpu_kernel.h +++ b/src/dbm/dbm_multiply_gpu_kernel.h @@ -22,12 +22,10 @@ extern "C" { * All arguments are assumed to be device pointers. * \author Ole Schuett ******************************************************************************/ -void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, - const double alpha, const int ntasks, - const dbm_task_t *batch, - const double *pack_a_data, - const double *pack_b_data, - double *shard_c_data); +void dbm_multiply_gpu_launch_kernel( + const offloadStream_t stream, const int mnk_range[3][2], const double alpha, + const int ntasks, const dbm_task_t *batch, const double *pack_a_data, + const double *pack_b_data, double *shard_c_data); #ifdef __cplusplus } diff --git a/src/dbm/dbm_multiply_opencl.c b/src/dbm/dbm_multiply_opencl.c new file mode 100644 index 0000000000..d78b5d9d70 --- /dev/null +++ b/src/dbm/dbm_multiply_opencl.c @@ -0,0 +1,119 @@ +/*----------------------------------------------------------------------------*/ +/* CP2K: A general program to perform molecular dynamics simulations */ +/* Copyright 2000-2024 CP2K developers group */ +/* */ +/* SPDX-License-Identifier: BSD-3-Clause */ +/*----------------------------------------------------------------------------*/ + +#include "../offload/offload_runtime.h" +#if defined(__OFFLOAD_OPENCL) && !defined(__NO_OFFLOAD_DBM) + +#include "dbm_multiply_gpu_kernel.h" +#include "dbm_multiply_opencl.cl.h" + +void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, + const int mnk_range[3][2], double alpha, + int ntasks, const dbm_task_t *tasks, + const double *pack_a_data, + const double *pack_b_data, + double *shard_c_data) { + static cl_kernel kernel = NULL; + static int task_split = 0; + int result = EXIT_SUCCESS, verbosity = c_dbcsr_acc_opencl_config.verbosity; + cl_event event, *const perf_event = + ((0 <= verbosity && 2 >= verbosity) ? NULL : &event); + const c_dbcsr_acc_opencl_stream_t *const str = ACC_OPENCL_STREAM(stream); + const int max_m = mnk_range[0][1], max_n = mnk_range[1][1]; + const size_t work_tasks = ntasks, wgsize = 0; + size_t ibatch = 0, iadata = 0, ibdata = 0, icdata = 0, work_size = 0; + c_dbcsr_acc_opencl_info_memptr_t adata, bdata, cdata, batch; + assert(NULL != pack_a_data && NULL != pack_b_data && NULL != shard_c_data); + assert(0 < mnk_range[0][0] && 0 < max_m && mnk_range[0][0] <= max_m); + assert(0 < mnk_range[1][0] && 0 < max_n && mnk_range[1][0] <= max_n); + assert(0 < mnk_range[2][0] && 0 < mnk_range[2][1] && + mnk_range[2][0] <= mnk_range[2][1]); + assert(NULL != str && NULL != str->queue); + assert(0 < ntasks && NULL != tasks); + /* creating/calling kernel must be consistent across threads */ + ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_main); +#if defined(OPENCL_DBM_SOURCE_MULTIPLY_OPENCL) + if (NULL == kernel) { /* first-time check if kernel is present */ + char params[ACC_OPENCL_BUFFERSIZE] = "-DLU=0 -DBN=4"; + const char *const flags = "-cl-fast-relaxed-math -cl-denorms-are-zero"; + const char *extensions[] = {NULL, NULL}; + const size_t nextensions = sizeof(extensions) / sizeof(*extensions); + const size_t offset = strlen(params); + const libxsmm_timer_tickint start = libxsmm_timer_tick(); + const int nchar = c_dbcsr_acc_opencl_flags_atomics( + &c_dbcsr_acc_opencl_config.device, c_dbcsr_acc_opencl_atomic_fp_64, + extensions, nextensions, params + offset, sizeof(params) - offset); + if (0 < nchar && sizeof(params) > (offset + nchar)) { + const char *const task_split_env = getenv("LIBDBM_TASK_SPLIT"); + task_split = (NULL == task_split_env ? 1 : atoi(task_split_env)); + result |= c_dbcsr_acc_opencl_kernel( + 0 /*source_is_file*/, OPENCL_DBM_SOURCE_MULTIPLY_OPENCL, + "dbm_multiply", params, + 0 == c_dbcsr_acc_opencl_config.debug ? flags : NULL, NULL /*try*/, + NULL /*try_ok*/, extensions, nextensions, &kernel); + if (2 <= verbosity || 0 > verbosity) { + if (EXIT_SUCCESS == result) { + fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel ms=%.1f\n", + 1E3 * libxsmm_timer_duration(start, libxsmm_timer_tick())); + } else { + fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel failed to generate\n"); + } + } + } + } +#else +#error "OpenCL kernel code not found!" +#endif + result |= c_dbcsr_acc_opencl_info_devptr_lock(&adata, NULL /*lock*/, + pack_a_data, 1 /*esize*/, + NULL /*amount*/, &iadata); + result |= c_dbcsr_acc_opencl_info_devptr_lock(&bdata, NULL /*lock*/, + pack_b_data, 1 /*esize*/, + NULL /*amount*/, &ibdata); + result |= c_dbcsr_acc_opencl_info_devptr_lock(&cdata, NULL /*lock*/, + shard_c_data, 1 /*esize*/, + NULL /*amount*/, &icdata); + result |= c_dbcsr_acc_opencl_info_devptr_lock( + &batch, NULL /*lock*/, tasks /*batch*/, sizeof(dbm_task_t), &work_tasks, + &ibatch); + assert(0 == iadata && 0 == ibdata && 0 == icdata); + work_size = (0 != task_split ? (work_tasks * max_m) : work_tasks); + result |= clSetKernelArg(kernel, 0, sizeof(cl_double), &alpha); + result |= clSetKernelArg(kernel, 1, sizeof(cl_int), &ibatch); + result |= clSetKernelArg(kernel, 2, sizeof(cl_int), &ntasks); + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 3, batch.memory); + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 4, adata.memory); + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 5, bdata.memory); + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 6, cdata.memory); + result |= clEnqueueNDRangeKernel(str->queue, kernel, 1 /*work_dim*/, + NULL /*offset*/, &work_size, + 0 != wgsize ? &wgsize : NULL, 0 /*num_wait*/, + NULL /*wait_list*/, perf_event); + if (NULL != perf_event && EXIT_SUCCESS == result) { + cl_ulong begin = 0, end = 0; + clWaitForEvents(1, perf_event); + result |= clGetEventProfilingInfo(*perf_event, CL_PROFILING_COMMAND_START, + sizeof(cl_ulong), &begin, NULL); + result |= clGetEventProfilingInfo(*perf_event, CL_PROFILING_COMMAND_END, + sizeof(cl_ulong), &end, NULL); + if (EXIT_SUCCESS == result) { + const double duration_ms = 1E-6 * LIBXSMM_DELTA(begin, end); + const double gflops = + (1E-6 * max_m * max_n * mnk_range[2][1] * ntasks) / duration_ms; + fprintf(stderr, + "INFO ACC/LIBDBM: DBM-kernel mnk=%ix%ix%i " + "ntasks=%i gflops=%.1f ms=%.2g\n", + max_m, max_n, mnk_range[2][1], ntasks, gflops, duration_ms); + } + } + ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_main); + OFFLOAD_CHECK(result); +} + +#endif // defined(__OFFLOAD_OPENCL) && !defined(__NO_OFFLOAD_DBM) + +// EOF diff --git a/src/dbm/dbm_multiply_opencl.cl b/src/dbm/dbm_multiply_opencl.cl new file mode 100644 index 0000000000..3c97883a40 --- /dev/null +++ b/src/dbm/dbm_multiply_opencl.cl @@ -0,0 +1,77 @@ +/*----------------------------------------------------------------------------*/ +/* CP2K: A general program to perform molecular dynamics simulations */ +/* Copyright 2000-2024 CP2K developers group */ +/* */ +/* SPDX-License-Identifier: BSD-3-Clause */ +/*----------------------------------------------------------------------------*/ +#include "../../exts/dbcsr/src/acc/opencl/common/opencl_atomics.h" +#include "dbm_multiply_internal.h" + +#define IDX(I, J, K, M, N) ((I) * (N) + (J) + (K)) +#define IDT(I, J, K, M, N) IDX(J, I, K, N, M) + +#define DBM_MULTIPLY_KERNEL(TASK, A, B, CV, M, N0, N1, UNROLL_N, UNROLL_K) \ + do { \ + UNROLL_K \ + for (int k = 0; k < (TASK).k; ++k) { \ + const int ia = IDT(M, k, (TASK).offset_a, (TASK).m, (TASK).k); \ + const double a = (A)[ia]; \ + UNROLL_N \ + for (int n = 0; n < (N1); ++n) { \ + const int ib = IDX(k, n + (N0), (TASK).offset_b, (TASK).k, (TASK).n); \ + (CV)[n] = MAD(a, (B)[ib], (CV)[n]); \ + } \ + } \ + } while (0) + +#define DBM_MULTIPLY_ACCUMULATE(ALPHA, TASK, CV, C, M, BN, N0) \ + do { /* flush private accumulator to global memory using atomics */ \ + UNROLL_FORCE(BN) \ + for (int n = 0; n < (BN); ++n) { \ + const int ic = IDT(M, n + (N0), (TASK).offset_c, (TASK).m, (TASK).n); \ + ACCUMULATE((C) + ic, (ALPHA) * (CV)[n]); \ + cv[n] = ZERO; /* reset */ \ + } \ + } while (0) + +kernel void dbm_multiply(double alpha, int itask, int ntasks, + global const dbm_task_t *tasks, + global const double *restrict am, + global const double *restrict bm, + global double *restrict cm) { + double cv[BN] = {0}; /* private accumulator */ + const int size = (int)get_global_size(0), i = (int)get_global_id(0); + + if (size != ntasks) { /* LIBDBM_TASK_SPLIT */ + const int max_m = size / ntasks, tid = i / max_m; + const dbm_task_t task = tasks[itask + min(tid, ntasks - 1)]; /* copy */ + const int m = i - tid * max_m; + if (m < task.m) { + if ((BN) < task.n) { + UNROLL_AUTO + for (int n0 = 0; n0 < task.n; n0 += (BN)) { + const int n1 = min(BN, task.n - n0); + DBM_MULTIPLY_KERNEL(task, am, bm, cv, m, n0, n1, UNROLL_FORCE(BN), + UNROLL_AUTO); + DBM_MULTIPLY_ACCUMULATE(alpha, task, cv, cm, m, BN, n0); + } + } else { /* small */ + DBM_MULTIPLY_KERNEL(task, am, bm, cv, m, 0, task.n, UNROLL_FORCE(BN), + UNROLL_FORCE(BN)); + DBM_MULTIPLY_ACCUMULATE(alpha, task, cv, cm, m, BN, 0); + } + } + } else { /* full matrix multiplication */ + const dbm_task_t task = tasks[itask + i]; + UNROLL_OUTER(1) + for (int m = 0; m < task.m; ++m) { + UNROLL(1) + for (int n0 = 0; n0 < task.n; n0 += (BN)) { + const int n1 = min(BN, task.n - n0); + DBM_MULTIPLY_KERNEL(task, am, bm, cv, m, n0, n1, UNROLL_AUTO, + UNROLL_AUTO); + DBM_MULTIPLY_ACCUMULATE(alpha, task, cv, cm, m, BN, n0); + } + } + } +} diff --git a/tools/precommit/README.md b/tools/precommit/README.md index 543c26cb1f..1ed4574a47 100644 --- a/tools/precommit/README.md +++ b/tools/precommit/README.md @@ -6,7 +6,7 @@ The precommit system consists of the following tools that analyze and format the - [prettify](../prettify/) to format Fortran files. - [analyze_src](../conventions/analyze_src.py) to check copyright banners and a few other things. - [ast.parse](https://docs.python.org/3/library/ast.html) to check Python syntax. -- [clang-format](https://clang.llvm.org/docs/ClangFormat.html) to format C and Cuda files. +- [clang-format](https://clang.llvm.org/docs/ClangFormat.html) to format C, Cuda, and OpenCL files. - [black](https://github.com/psf/black) to format Python scripts. - [shfmt](https://github.com/mvdan/sh) to format Shell scripts. - [shellcheck](https://github.com/koalaman/shellcheck) to analyze Shell scripts. diff --git a/tools/precommit/check_file_properties.py b/tools/precommit/check_file_properties.py index 48310abfb1..2334082bba 100755 --- a/tools/precommit/check_file_properties.py +++ b/tools/precommit/check_file_properties.py @@ -21,7 +21,9 @@ FLAG_EXCEPTIONS = ( r"\$\{.*\}\$", r"__.*__", r"CUDA_VERSION", + r"DBM_LIBXSMM_PREFETCH", r"DBM_VALIDATE_AGAINST_DBCSR", + r"OPENCL_DBM_SOURCE_MULTIPLY_OPENCL", r"FD_DEBUG", r"GRID_DO_COLLOCATE", r"INTEL_MKL_VERSION", @@ -58,7 +60,6 @@ FLAG_EXCEPTIONS = ( r"LIBXSMM_VERSION2", r"LIBXSMM_VERSION3", r"LIBXSMM_VERSION4", - r"DBM_LIBXSMM_PREFETCH", ) FLAG_EXCEPTIONS_RE = re.compile(r"|".join(FLAG_EXCEPTIONS)) diff --git a/tools/precommit/precommit.py b/tools/precommit/precommit.py index a6768b9877..81991ba08b 100755 --- a/tools/precommit/precommit.py +++ b/tools/precommit/precommit.py @@ -194,7 +194,7 @@ def process_file(fn: str, allow_modifications: bool) -> None: run_local_tool("./tools/doxify/doxify.sh", fn) run_prettify(fn) - if re.match(r".*\.(c|cu|h)$", fn): + if re.match(r".*\.(c|cu|cl|h)$", fn): run_remote_tool("clangformat", fn) if re.match(r".*\.(cc|cpp|cxx|hcc|hpp|hxx)$", fn):