From eeea73ec643a41ed6d74d67c204c1ce7c8d16dc6 Mon Sep 17 00:00:00 2001 From: Hans Pabst Date: Wed, 5 Feb 2025 12:28:58 +0100 Subject: [PATCH] DBM: improvements and optimizations (#3937) - Avoid stream-synchronization for copy-completion (GPU); at the expense of higher memory consumption (lifetime). - Implement SMM-threshold (CPU/LIBXSMM/BLAS). OpenCL backend: - Adjusted valid range of kernel-parameter (BN). - Adjusted default value (BN). Makefile/Miniapp - Consolidated MPI-compilation- - Ensure gcc rather than cc. - Adjusted CFLAGS (-fPIC). - Fixed clean-target. Other - Purified some getter-like functions. --- src/dbm/Makefile | 42 ++++++++++++++++------------------- src/dbm/dbm_multiply_comm.c | 9 ++++---- src/dbm/dbm_multiply_cpu.c | 30 +++++++++++++++---------- src/dbm/dbm_multiply_gpu.c | 35 +++++++++++++++++------------ src/dbm/dbm_multiply_opencl.c | 15 +++++-------- src/dbt/tas/dbt_tas_base.F | 14 ++++++------ 6 files changed, 75 insertions(+), 70 deletions(-) diff --git a/src/dbm/Makefile b/src/dbm/Makefile index 81e7932776..d15b9183ee 100644 --- a/src/dbm/Makefile +++ b/src/dbm/Makefile @@ -24,26 +24,32 @@ endif # Optimization flag derived from OPT flag OPTFLAG ?= -O$(patsubst O%,%,$(OPT)) -CFLAGS := -g $(OPTFLAG) -march=native -Wall -Wextra -Wcast-qual +CFLAGS := -fPIC -g $(OPTFLAG) -march=native -Wall -Wextra -Wcast-qual # Intel Compiler ICX := $(shell which icx 2>/dev/null) INTEL ?= $(if $(ICX),1,0) +# Build with MPI when MPICC is given +MPI ?= $(if $(MPICC),1,0) + ifeq ($(INTEL),0) MKL_FCRTL := gf LIBS += $(if $(OMPRT),-l$(OMPRT),-fopenmp) -CFLAGS += -fopenmp +CFLAGS += -fopenmp -Wno-vla-parameter +CC := $(if $(filter-out 0,$(MPI)),mpicc,gcc) else MKL_FCRTL := intel LIBS += $(if $(OMPRT),-l$(OMPRT),-qopenmp) CFLAGS += -qopenmp -CC := icx +CC := $(if $(filter-out 0,$(MPI)),mpiicx,icx) endif -# Some additional flags when relying on default-CC -ifeq ($(CC),) -CFLAGS += -Wno-vla-parameter +ifneq ($(MPI),0) +CFLAGS += -D__parallel +ifneq ($(MPICC),) +CC := $(MPICC) +endif endif NVCC := $(shell which nvcc 2>/dev/null) @@ -91,7 +97,11 @@ endif all: dbm_miniapp.x clean: - rm -fv $(ALL_OBJECTS) $(OPENCL_GENKRNL) + rm -fv $(ALL_OBJECTS) $(OPENCL_GENKRNL) \ + dbm_multiply_gpu.o \ + dbm_multiply_gpu_kernel.o \ + dbm_multiply_opencl.o \ + dbm_miniapp.o realclean: clean rm -fv dbm_miniapp.x @@ -177,32 +187,18 @@ LIBS += -L${ROCM_PATH}/lib -lamdhip64 -lhipfft -lhipblas -lrocblas 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 \ +LIBS += $(LIBXSMMROOT)/lib/libxsmm.a \ + $(LIBXSMMROOT)/lib/libxsmmext.a \ -lpthread -lrt endif - -# Build with MPI when mpicc compiler is present. -ifneq ($(MPICC),) -%.o: %.c $(OMP_TRACE) $(ALL_HEADERS) - cd $(dir $<); $(MPICC) -c -std=c11 $(CFLAGS) -D__parallel $(notdir $<) - -dbm_miniapp.x: dbm_miniapp.o $(ALL_OBJECTS) - $(MPICC) $(CFLAGS) -o $@ $^ $(LIBS) - - -# Build without MPI otherwise. -else %.o: %.c $(OMP_TRACE) $(ALL_HEADERS) cd $(dir $<); $(CC) -c -std=c11 $(CFLAGS) $(notdir $<) dbm_miniapp.x: dbm_miniapp.o $(ALL_OBJECTS) $(CC) $(CFLAGS) -o $@ $^ $(LIBS) -endif #EOF diff --git a/src/dbm/dbm_multiply_comm.c b/src/dbm/dbm_multiply_comm.c index e00c07450e..ec93a7e897 100644 --- a/src/dbm/dbm_multiply_comm.c +++ b/src/dbm/dbm_multiply_comm.c @@ -210,13 +210,14 @@ static void fill_send_buffers( const dbm_shard_t *shard = &matrix->shards[ishard]; const double *blk_data = &shard->data[blk->offset]; const int row_size = plan->row_size, col_size = plan->col_size; + const int plan_size = row_size * col_size; const int irank = plan->rank; // The blk_send_data is ordered by rank, thread, and block. // data_send_displ[irank]: Start of data for irank within blk_send_data. // ndata_mythread[irank]: Current threads offset within data for irank. nblks_mythread[irank] -= 1; - ndata_mythread[irank] -= row_size * col_size; + ndata_mythread[irank] -= plan_size; const int offset = data_send_displ[irank] + ndata_mythread[irank]; const int jblock = blks_send_displ[irank] + nblks_mythread[irank]; @@ -226,17 +227,17 @@ static void fill_send_buffers( for (int i = 0; i < row_size; i++) { for (int j = 0; j < col_size; j++) { const double element = blk_data[j * row_size + i]; - norm += element * element; data_send[offset + i * col_size + j] = element; + norm += element * element; } } blks_send[jblock].free_index = plan->blk->col; blks_send[jblock].sum_index = plan->blk->row; } else { - for (int i = 0; i < row_size * col_size; i++) { + for (int i = 0; i < plan_size; i++) { const double element = blk_data[i]; - norm += element * element; data_send[offset + i] = element; + norm += element * element; } blks_send[jblock].free_index = plan->blk->row; blks_send[jblock].sum_index = plan->blk->col; diff --git a/src/dbm/dbm_multiply_cpu.c b/src/dbm/dbm_multiply_cpu.c index eedf80e8df..19a5710923 100644 --- a/src/dbm/dbm_multiply_cpu.c +++ b/src/dbm/dbm_multiply_cpu.c @@ -120,21 +120,27 @@ void dbm_multiply_cpu_process_batch(const int ntasks, dbm_task_t batch[ntasks], task_next = batch[batch_order[(itask + 1) < ntasks ? (itask + 1) : itask]]; if (task.m != kernel_m || task.n != kernel_n || task.k != kernel_k) { + if (LIBXSMM_SMM(task.m, task.n, task.m, 1 /*assume in-$, no RFO*/, + sizeof(double))) { #if LIBXSMM_VERSION2(1, 17) < LIBXSMM_VERSION_NUMBER - const libxsmm_gemm_shape shape = libxsmm_create_gemm_shape( - task.m, task.n, task.k, task.m /*lda*/, task.n /*ldb*/, - task.m /*ldc*/, LIBXSMM_DATATYPE_F64 /*aprec*/, - LIBXSMM_DATATYPE_F64 /*bprec*/, LIBXSMM_DATATYPE_F64 /*cprec*/, - LIBXSMM_DATATYPE_F64 /*calcp*/); - kernel_func = (LIBXSMM_FEQ(1.0, alpha) - ? libxsmm_dispatch_gemm(shape, (libxsmm_bitfield)flags, - (libxsmm_bitfield)prefetch) - : NULL); + const libxsmm_gemm_shape shape = libxsmm_create_gemm_shape( + task.m, task.n, task.k, task.m /*lda*/, task.n /*ldb*/, + task.m /*ldc*/, LIBXSMM_DATATYPE_F64 /*aprec*/, + LIBXSMM_DATATYPE_F64 /*bprec*/, LIBXSMM_DATATYPE_F64 /*cprec*/, + LIBXSMM_DATATYPE_F64 /*calcp*/); + kernel_func = + (LIBXSMM_FEQ(1.0, alpha) + ? libxsmm_dispatch_gemm(shape, (libxsmm_bitfield)flags, + (libxsmm_bitfield)prefetch) + : NULL); #else - kernel_func = libxsmm_dmmdispatch(task.m, task.n, task.k, NULL /*lda*/, - NULL /*ldb*/, NULL /*ldc*/, &alpha, - &beta, &flags, &prefetch); + kernel_func = libxsmm_dmmdispatch(task.m, task.n, task.k, NULL /*lda*/, + NULL /*ldb*/, NULL /*ldc*/, &alpha, + &beta, &flags, &prefetch); #endif + } else { + kernel_func = NULL; + } kernel_m = task.m; kernel_n = task.n; kernel_k = task.k; diff --git a/src/dbm/dbm_multiply_gpu.c b/src/dbm/dbm_multiply_gpu.c index 7a8a8f682a..006402837e 100644 --- a/src/dbm/dbm_multiply_gpu.c +++ b/src/dbm/dbm_multiply_gpu.c @@ -106,7 +106,6 @@ 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) { - if (ntasks == 0) { return; // Nothing to do. } @@ -116,47 +115,55 @@ void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch, const dbm_shard_t *shard_c_host = &ctx->shards_c_host[kshard]; dbm_shard_gpu_t *shard_c_dev = &ctx->shards_c_dev[kshard]; + assert(NULL != shard_c_host && NULL != shard_c_dev); // Upload new batch. dbm_task_t *batch_dev = &ctx->batches_dev[kshard * ctx->max_batch_size]; const size_t size = ntasks * sizeof(dbm_task_t); offloadMemcpyAsyncHtoD(batch_dev, batch, size, shard_c_dev->stream); - offloadEvent_t batch_uploaded; - offloadEventCreate(&batch_uploaded); - offloadEventRecord(batch_uploaded, shard_c_dev->stream); + offloadEvent_t memsetup; + offloadEventCreate(&memsetup); // Reallocate shard_c_dev->data if necessary. + double *old_data_dev = NULL; if (shard_c_host->data_promised > shard_c_dev->data_allocated) { - double *old_data_dev = shard_c_dev->data; shard_c_dev->data_allocated = ALLOCATION_FACTOR * shard_c_host->data_promised; + old_data_dev = shard_c_dev->data; shard_c_dev->data = dbm_mempool_device_malloc(shard_c_dev->data_allocated * sizeof(double)); + // Omit to wait for copy before freeing old buffer. offloadMemcpyAsyncDtoD(shard_c_dev->data, old_data_dev, shard_c_dev->data_size * sizeof(double), shard_c_dev->stream); - // Wait for copy to complete before freeing old buffer. - offloadStreamSynchronize(shard_c_dev->stream); - dbm_mempool_free(old_data_dev); - } + offloadEventRecord(memsetup, shard_c_dev->stream); - // Zero new blocks if necessary. - if (shard_c_host->data_promised > shard_c_dev->data_size) { + // Zero new blocks if necessary. const int tail = shard_c_host->data_promised - shard_c_dev->data_size; offloadMemsetAsync(&shard_c_dev->data[shard_c_dev->data_size], 0, tail * sizeof(double), shard_c_dev->stream); shard_c_dev->data_size = shard_c_host->data_promised; + } else { // event only covers offloadMemcpyAsyncHtoD + offloadEventRecord(memsetup, shard_c_dev->stream); } // Launch kernel. + assert(0 != shard_c_dev->data_size); 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. - offloadEventSynchronize(batch_uploaded); - offloadEventDestroy(batch_uploaded); + // Wait for: + // - Batch to be uploaded (before refilling it). + // - Device memory buffer (if resized). + offloadEventSynchronize(memsetup); + offloadEventDestroy(memsetup); + + // Safely freeing old buffer. + if (NULL != old_data_dev) { + dbm_mempool_free(old_data_dev); + } } /******************************************************************************* diff --git a/src/dbm/dbm_multiply_opencl.c b/src/dbm/dbm_multiply_opencl.c index cf79c273f6..c8a1e9b49d 100644 --- a/src/dbm/dbm_multiply_opencl.c +++ b/src/dbm/dbm_multiply_opencl.c @@ -11,10 +11,6 @@ #include "dbm_multiply_gpu_kernel.h" #include "dbm_multiply_opencl.cl.h" -#if !defined(OPENCL_DBM_TLS) && 1 -#define OPENCL_DBM_TLS LIBXSMM_TLS -#endif - 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, @@ -23,11 +19,7 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double *shard_c_data) { /* creating/calling kernel must be consistent across threads */ static cl_kernel kernel_global = NULL; -#if defined(OPENCL_DBM_TLS) - static OPENCL_DBM_TLS cl_kernel kernel = NULL; -#else - cl_kernel kernel = NULL; -#endif + static LIBXSMM_TLS cl_kernel kernel = NULL; static int ndims = 1; static size_t wgsize[] = {0, 0, 0}; const libxsmm_timer_tickint start = libxsmm_timer_tick(); @@ -63,7 +55,10 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, const int gen = (NULL == gen_env ? 1 /*default*/ : atoi(gen_env)); const int xf = (NULL == xf_env ? -1 /*default*/ : atoi(xf_env)); const int lu = LIBXSMM_CLMP(NULL == lu_env ? 0 : atoi(lu_env), -2, 1); - const int bn = LIBXSMM_CLMP(NULL == bn_env ? 8 : atoi(bn_env), 4, 32); + const int bn0 = (0 == c_dbcsr_acc_opencl_config.device.nv + ? (0 == c_dbcsr_acc_opencl_config.device.amd ? 4 : 8) + : 2); + const int bn = LIBXSMM_CLMP(NULL == bn_env ? bn0 : atoi(bn_env), 1, 32); const char *extensions[] = {NULL, NULL}, *flags = NULL; size_t nextensions = sizeof(extensions) / sizeof(*extensions); const size_t wgsize0 = c_dbcsr_acc_opencl_config.device.wgsize[0]; diff --git a/src/dbt/tas/dbt_tas_base.F b/src/dbt/tas/dbt_tas_base.F index 79ac7a6c2a..87846c24b4 100644 --- a/src/dbt/tas/dbt_tas_base.F +++ b/src/dbt/tas/dbt_tas_base.F @@ -831,7 +831,7 @@ CONTAINS !> \return ... !> \author Patrick Seewald ! ************************************************************************************************** - FUNCTION dbt_tas_nblkrows_total(matrix) RESULT(nblkrows_total) + PURE FUNCTION dbt_tas_nblkrows_total(matrix) RESULT(nblkrows_total) TYPE(dbt_tas_type), INTENT(IN) :: matrix INTEGER(KIND=int_8) :: nblkrows_total @@ -844,7 +844,7 @@ CONTAINS !> \return ... !> \author Patrick Seewald ! ************************************************************************************************** - FUNCTION dbt_tas_nfullrows_total(matrix) RESULT(nfullrows_total) + PURE FUNCTION dbt_tas_nfullrows_total(matrix) RESULT(nfullrows_total) TYPE(dbt_tas_type), INTENT(IN) :: matrix INTEGER(KIND=int_8) :: nfullrows_total @@ -857,7 +857,7 @@ CONTAINS !> \return ... !> \author Patrick Seewald ! ************************************************************************************************** - FUNCTION dbt_tas_nblkcols_total(matrix) RESULT(nblkcols_total) + PURE FUNCTION dbt_tas_nblkcols_total(matrix) RESULT(nblkcols_total) TYPE(dbt_tas_type), INTENT(IN) :: matrix INTEGER(KIND=int_8) :: nblkcols_total @@ -870,7 +870,7 @@ CONTAINS !> \return ... !> \author Patrick Seewald ! ************************************************************************************************** - FUNCTION dbt_tas_nfullcols_total(matrix) RESULT(nfullcols_total) + PURE FUNCTION dbt_tas_nfullcols_total(matrix) RESULT(nfullcols_total) TYPE(dbt_tas_type), INTENT(IN) :: matrix INTEGER(KIND=int_8) :: nfullcols_total @@ -1003,11 +1003,11 @@ CONTAINS OPTIONAL :: local_rows, local_cols CLASS(dbt_tas_distribution), ALLOCATABLE, OPTIONAL, & - INTENT(OUT) :: proc_row_dist, proc_col_dist + INTENT(OUT) :: proc_row_dist, proc_col_dist CLASS(dbt_tas_rowcol_data), ALLOCATABLE, OPTIONAL, & - INTENT(OUT) :: row_blk_size, col_blk_size + INTENT(OUT) :: row_blk_size, col_blk_size TYPE(dbt_tas_distribution_type), OPTIONAL :: distribution - CHARACTER(len=*), INTENT(OUT), OPTIONAL :: name + CHARACTER(len=*), INTENT(OUT), OPTIONAL :: name TYPE(dbt_tas_split_info) :: info INTEGER :: irow, icol