comp-chem/cp2k
1
0
Fork 0
mirror of https://github.com/cp2k/cp2k.git synced 2026-07-27 21:55:16 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

DBM: improvements and optimizations (#3937)

Browse source 
- 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.
This commit is contained in:
Hans Pabst 2025-02-05 12:28:58 +01:00 • committed by GitHub
parent 8824751c23
commit eeea73ec64
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
6 changed files with 75 additions and 70 deletions
Download patch file Download diff file Expand all files Collapse all files

42
src/dbm/Makefile
View file

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

9
src/dbm/dbm_multiply_comm.c
View file

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

30
src/dbm/dbm_multiply_cpu.c
View file

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

35
src/dbm/dbm_multiply_gpu.c
View file

@ -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);
}
}
/*******************************************************************************

15
src/dbm/dbm_multiply_opencl.c
View file

@ -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];

14
src/dbt/tas/dbt_tas_base.F
View file

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