From d2993ca2bb8ff2d0d7ed87b5a007553714f63374 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Urban=20Bor=C5=A1tnik?= Date: Mon, 26 Sep 2011 13:19:07 +0000 Subject: [PATCH] DBCSR/CUDA work. Allows multiple GPUs per process; Fix for multiple stack binning-simplifies input; CUDA optimizations; DBCSR configuration output after setup; svn-origin-rev: 11794 --- cuda_tools/NVOBJECTDEFS | 1 + cuda_tools/README | 7 +- cuda_tools/dbcsr_cuda.h | 3 + cuda_tools/dbcsr_cuda_calc.cu | 215 ++++++---- cuda_tools/dbcsr_cuda_calc_d.cu | 377 ++++++++++++++++-- cuda_tools/dbcsr_cuda_dev.cu | 36 ++ src/cp2k_runs.F | 12 +- src/cp_dbcsr_interface.F | 69 +++- src/dbcsr_lib/OBJECTDEFS | 1 + src/dbcsr_lib/dbcsr_config.F | 39 +- src/dbcsr_lib/dbcsr_cuda_device.F | 236 +++++++++++ src/dbcsr_lib/dbcsr_cuda_operations.F | 57 +-- ...dbcsr_cuda_operations__nametype1_.template | 12 +- src/dbcsr_lib/dbcsr_cuda_operations_c.F | 12 +- src/dbcsr_lib/dbcsr_cuda_operations_d.F | 12 +- src/dbcsr_lib/dbcsr_cuda_operations_i.F | 12 +- src/dbcsr_lib/dbcsr_cuda_operations_l.F | 12 +- src/dbcsr_lib/dbcsr_cuda_operations_r.F | 12 +- src/dbcsr_lib/dbcsr_cuda_operations_z.F | 12 +- src/dbcsr_lib/dbcsr_internal_operations.F | 33 +- src/dbcsr_lib/dbcsr_operations.F | 19 +- src/dbcsr_lib/dbcsr_util.F | 6 +- src/environment.F | 11 +- src/f77_interface.F | 2 +- src/input_cp2k.F | 21 +- tests/LIBTEST/dbcsr_multistack.inp | 4 +- 26 files changed, 977 insertions(+), 256 deletions(-) create mode 100644 cuda_tools/dbcsr_cuda_dev.cu create mode 100644 src/dbcsr_lib/dbcsr_cuda_device.F diff --git a/cuda_tools/NVOBJECTDEFS b/cuda_tools/NVOBJECTDEFS index 558769df84..6e83345b49 100644 --- a/cuda_tools/NVOBJECTDEFS +++ b/cuda_tools/NVOBJECTDEFS @@ -7,5 +7,6 @@ dbcsr_cuda_calc_r.o \ dbcsr_cuda_calc_d.o \ dbcsr_cuda_calc_c.o \ dbcsr_cuda_calc_z.o \ +dbcsr_cuda_dev.o \ dbcsr_cuda_mem.o \ dbcsr_cuda_timing.o diff --git a/cuda_tools/README b/cuda_tools/README index f5f2e39ab1..f315830427 100644 --- a/cuda_tools/README +++ b/cuda_tools/README @@ -9,10 +9,15 @@ BASIC: 1) Add a line "NVCC = nvcc" pointing the makefile to the nvcc compiler. -2) Add -D__CUDAPW and -D__FFTSGL to the DFLAGS and NVFLAGS environmental variables +2) Add -D__CUDAPW and -D__FFTCU to the DFLAGS and NVFLAGS environmental variables; + optionally add -D__FFTSGL for faster single precision FFTs. 3) Add libcudart.so, libcufft.so, and cublas to the LIBS variable. +4) If using in conjunction with DBCSR (see next section) then specify + an appropriate size for the GLOBAL / CUDA / MEMORY option in the + input file. + DBCSR (experimental, with limitations): 1. Add __DBCSR_CUDA to DFLAGS. diff --git a/cuda_tools/dbcsr_cuda.h b/cuda_tools/dbcsr_cuda.h index c7dc9d625c..ce58dfec8a 100644 --- a/cuda_tools/dbcsr_cuda.h +++ b/cuda_tools/dbcsr_cuda.h @@ -13,3 +13,6 @@ int cuda_error_check (cudaError_t cudaError); + +#define GROUPING 16 + diff --git a/cuda_tools/dbcsr_cuda_calc.cu b/cuda_tools/dbcsr_cuda_calc.cu index 06a71baa0d..671c1cfd18 100644 --- a/cuda_tools/dbcsr_cuda_calc.cu +++ b/cuda_tools/dbcsr_cuda_calc.cu @@ -76,91 +76,158 @@ __global__ void stack_mm_z double *__restrict__ c_data, int *__restrict__ c_locks); +__global__ void stack_mm_mnk_d ( + const int *__restrict__ param_stack, + const int careful, const int nruns, + const int m, const int n, const int k, +// const int mn, const int mk, const int nk, const int maxb, + const int liter, + const double *__restrict__ a_data, + const double *__restrict__ b_data, + double *__restrict__ c_data, + int *__restrict__ c_locks); + +__global__ void stack_mm_mnk_d_direct ( + const int *__restrict__ param_stack, + const int careful, const int nruns, + const int m, const int n, const int k, const int mn, + const double *__restrict__ a_data, + const double *__restrict__ b_data, + double *__restrict__ c_data, + int *__restrict__ c_locks); + +__global__ void stack_mm_mnk_vec_d ( + const int *__restrict__ param_stack, + const int stack_size, const int nmat, + const int m, const int n, const int k, const int mn, + const double *__restrict__ a_data, + const double *__restrict__ b_data, + double *__restrict__ c_data, + int *__restrict__ c_locks); + /** * \brief Bridge routine to call appropriate CUDA kernel. */ -extern "C" int dc_do_stack_cu(int *param_stack, int stack_size, int nparams, - int which_data, - void *a_data, void *b_data, void *c_data, - int *c_locks, - int m_max, int n_max, int k_max) { - int maxt; - cudaError_t cErr; - int myDevice; - size_t shared_size; - struct cudaDeviceProp devProperties; +extern "C" int dc_do_stack_cu( + int *param_stack, int stack_size, int nparams, + int which_data, + void *a_data, void *b_data, void *c_data, + int *c_locks, + int m_max, int n_max, int k_max, int def_mnk) +{ - if (verbose_print) - printf("Locks address %p.\n", c_locks); + int maxt, nmat, careful, nruns; + cudaError_t cErr; + int myDevice; + size_t shared_size; + struct cudaDeviceProp devProperties; + int mn, mk, nk, maxb, liter; - maxt = m_max * n_max; + if (verbose_print) + printf("Locks address %p.\n", c_locks); - cErr = cudaGetDevice(&myDevice); - if (cuda_error_check (cErr)) return 1; + maxt = m_max * n_max; - cErr = cudaGetDeviceProperties(&devProperties, myDevice); - if (cuda_error_check (cErr)) return 1; + cErr = cudaGetDevice(&myDevice); + if (cuda_error_check (cErr)) return 1; - if (maxt > devProperties.maxThreadsPerBlock) - return 3; + cErr = cudaGetDeviceProperties(&devProperties, myDevice); + if (cuda_error_check (cErr)) return 1; -// zerolocks <<< (stack_size+127)/128, 128 >>> ( -// c_locks, param_stack, stack_size -// ); + if (maxt > devProperties.maxThreadsPerBlock) + return 3; - switch (which_data) { - /* The data type identifier numbers correspond to the values - defined in dbcsr_types.F. */ - case 1: - /* Real, single precision */ - shared_size = (m_max*k_max + k_max*n_max)*sizeof(float); - if (shared_size > devProperties.sharedMemPerBlock) return 4; - stack_mm_r <<< stack_size, maxt, shared_size >>> - (param_stack, stack_size, nparams, - (float *) a_data, (float *) b_data, (float *) c_data, - c_locks); - break; - case 3: - /* Real, double precision */ - shared_size = (m_max*k_max + k_max*n_max)*sizeof(double); - if (shared_size > devProperties.sharedMemPerBlock) return 4; - stack_mm_d <<< stack_size, maxt, shared_size >>> - (param_stack, stack_size, nparams, - (double *) a_data, (double *) b_data, (double *) c_data, - c_locks); - break; - case 5: - /* Complex, single precision */ - shared_size = (m_max*k_max + k_max*n_max)*sizeof(float)*2; - if (shared_size > devProperties.sharedMemPerBlock) return 4; - stack_mm_c <<< stack_size, maxt, shared_size >>> - (param_stack, stack_size, nparams, - (float *) a_data, (float *) b_data, (float *) c_data, - c_locks); - break; - case 7: - /* Complex, single precision */ - shared_size = (m_max*k_max + k_max*n_max)*sizeof(double)*2; - if (shared_size > devProperties.sharedMemPerBlock) return 4; - stack_mm_z <<< stack_size, maxt, shared_size >>> - (param_stack, stack_size, nparams, - (double *) a_data, (double *) b_data, (double *) c_data, - c_locks); - break; - default: - return 2; - } - if (cuda_error_check (cudaGetLastError())) return 1; - return 0; + switch (which_data) { + /* The data type identifier numbers correspond to the values + defined in dbcsr_types.F. */ + case 1: + /* Real, single precision */ + shared_size = (m_max*k_max + k_max*n_max)*sizeof(float); + if (shared_size > devProperties.sharedMemPerBlock) return 4; + stack_mm_r <<< stack_size, maxt, shared_size >>> + (param_stack, stack_size, nparams, + (float *) a_data, (float *) b_data, (float *) c_data, + c_locks); + break; + case 3: + /* Real, double precision */ + shared_size = (m_max*k_max + k_max*n_max)*sizeof(double); + //shared_size = 0*sizeof(double); + if (shared_size > devProperties.sharedMemPerBlock) return 4; + if (def_mnk) { + //printf("Defined m,n,k: %d %d %d; %d.\n", m_max, n_max, k_max, stack_size); + if (0) { + mn = m_max * n_max; + mk = m_max * k_max; + nk = n_max * k_max; + nmat = 1; + if (m_max > 0) + nmat = 128 / m_max; + careful = (stack_size/GROUPING); + nruns = stack_size - careful * GROUPING; + shared_size = 0; + stack_mm_mnk_vec_d <<< (stack_size+nmat-1)/nmat, nmat*m_max, shared_size >>> + (param_stack, stack_size, nmat, + m_max, n_max, k_max, mn, + (double *) a_data, (double *) b_data, (double *) c_data, + c_locks); + } else if (0) { + careful = (stack_size/GROUPING); + nruns = stack_size - careful * GROUPING; + stack_mm_mnk_d_direct <<< (stack_size+GROUPING-1)/GROUPING, maxt, 0 >>> + (param_stack, careful, nruns, + m_max, n_max, k_max, m_max*n_max, + (double *) a_data, (double *) b_data, (double *) c_data, + c_locks); + } else { + mn = m_max * n_max; + mk = m_max * k_max; + nk = n_max * k_max; + careful = (stack_size/GROUPING); + nruns = stack_size - careful * GROUPING; + maxb = MAX(mk, nk); + liter = (maxb-1) / maxt; + stack_mm_mnk_d <<< (stack_size+GROUPING-1)/GROUPING, maxt, shared_size >>> + (param_stack, careful, nruns, + m_max, n_max, k_max, + //mn, mk, nk, maxb, liter, + liter, + (double *) a_data, (double *) b_data, (double *) c_data, + c_locks); + } + } else { + //printf("Generic m,n,k: %d %d %d; %d.\n", m_max, n_max, k_max, stack_size); + stack_mm_d <<< stack_size, maxt, shared_size >>> + (param_stack, stack_size, nparams, + (double *) a_data, (double *) b_data, (double *) c_data, + c_locks); + } + break; + case 5: + /* Complex, single precision */ + shared_size = (m_max*k_max + k_max*n_max)*sizeof(float)*2; + if (shared_size > devProperties.sharedMemPerBlock) return 4; + stack_mm_c <<< stack_size, maxt, shared_size >>> + (param_stack, stack_size, nparams, + (float *) a_data, (float *) b_data, (float *) c_data, + c_locks); + break; + case 7: + /* Complex, single precision */ + shared_size = (m_max*k_max + k_max*n_max)*sizeof(double)*2; + if (shared_size > devProperties.sharedMemPerBlock) return 4; + stack_mm_z <<< stack_size, maxt, shared_size >>> + (param_stack, stack_size, nparams, + (double *) a_data, (double *) b_data, (double *) c_data, + c_locks); + break; + default: + return 2; + } + if (cuda_error_check (cudaGetLastError())) return 1; + + return 0; }; - -extern "C" int dc_thread_sync_cu() { - cudaError_t cErr; - - cErr = cudaThreadSynchronize (); - if (cuda_error_check (cErr)) return 1; - return 0; -} diff --git a/cuda_tools/dbcsr_cuda_calc_d.cu b/cuda_tools/dbcsr_cuda_calc_d.cu index a98149dd49..1823189699 100644 --- a/cuda_tools/dbcsr_cuda_calc_d.cu +++ b/cuda_tools/dbcsr_cuda_calc_d.cu @@ -7,6 +7,130 @@ extern __shared__ double cache[]; + +__global__ void stack_mm_mnk_d ( + const int *__restrict__ param_stack, + const int careful, const int nruns, + const int m, const int n, const int k, + //const int mn, const int mk, const int kn, const int maxb, + const int liter, + const double *__restrict__ a_data, + const double *__restrict__ b_data, + double *__restrict__ c_data, + int *__restrict__ c_locks) { + + /** + * \var sp which stack member this thread block is processing + (= CUDA thread block) + * \var psp pointer to first element of parameters + * \var c_loc pointer to C data + * \var run run number + * \var nrun number of runs + * \var my_id my ID for locking + * \var tn thread number (of CUDA thread block) + * \var mn product of the block dimensions + * \var l multiplication loop index + * \var c, r C matrix row, column of this thread + * \var myc C matrix accumulator + * \var buff_l cache for A data + * \var buff_r cache for B data + * \var c_id translated C block number (used in locking) + * \var lock_owner current C block owner (used in locking) + */ + + int lock_owner, c_id, my_id; + const int mn = m * n; + const int mk = m * k; + const int kn = n * k; + const int r = threadIdx.x % m; + const int c = threadIdx.x / m; + int l, i; + double myc, tmp; + const double * __restrict__ buff_l, * __restrict__ buff_r; + + int psp, c_loc; + + int run, nrun; + + __shared__ int our_params[7]; + double *buff; + + buff = (double *) cache; + buff_l = buff; + buff_r = &(buff[mk]); + + nrun = GROUPING; + if (blockIdx.x == careful) + nrun = nruns; + + for (run = 0; run < nrun; run ++) { + psp = 7*(blockIdx.x*GROUPING + run); + +// for (l = 0; l <= (mk-1) / blockDim.x; l++) { +// i = threadIdx.x+blockDim.x*l; +// if (i < mk) +// buff[i] = a_data[our_params[3]-1+i]; +// } +// for (l = 0; l <= (kn-1) / blockDim.x; l++) { +// i = threadIdx.x+blockDim.x*l; +// if (i < kn) +// buff[mk+i] = b_data[our_params[4]-1+i]; +// } + for (l = 0; l <= liter; l++) { + i = threadIdx.x+blockDim.x*l; + if (i < mk) + buff[i] = a_data[param_stack[psp+3]-1+i]; + if (i < kn) + buff[mk+i] = b_data[param_stack[psp+4]-1+i]; + } + + syncthreads(); + + /* Do actual multiplication. */ + if (threadIdx.x < mn) { + myc = 0.0l; + + for (l = 0; l < k; l++) { + myc = myc + + buff_l[ l*m + r] * + buff_r[ c*k + l]; + } + } + + /* Lock the C block. */ + c_loc = param_stack[psp+5]-1; + //c_loc = our_params[5]-1; + syncthreads(); + c_id = param_stack[psp+6]-1; + //c_id = our_params[6]-1; + + if (threadIdx.x == 0) { + my_id = blockIdx.x+1; + lock_owner = 0; + while ((lock_owner != my_id)) + lock_owner = atomicCAS (&(c_locks[c_id]), 0, my_id); + } else if (threadIdx.x==1) { + tmp = c_data[c_loc]; + } + + + /* Add our results to the C block. */ + syncthreads(); + if (threadIdx.x < mn) { + c_data[c_loc+threadIdx.x] += myc; + } + + /* Release the lock on the C block. */ + syncthreads(); + if (threadIdx.x == 0) { + c_locks[c_id] = 0; + } + } + + +}; + + __global__ void stack_mm_d (const int *__restrict__ param_stack, int stack_size, int nparams, @@ -21,7 +145,6 @@ __global__ void stack_mm_d * \var sp_one translated stack (=sp+1) * \var tn thread number (of CUDA thread block) * \var nt number of threads (size of CUDA thread block) - * \var our_params cache for this thread block's multiplication parameters * \var m, n, k dimensions of the blocks (C is m*n, A is m*k, B is k*n) * \var mn, mk, kn product of the block dimensions * \var l multiplication loop index @@ -33,51 +156,32 @@ __global__ void stack_mm_d */ int sp, lock_owner, c_id, sp_one; - int tn, nt; + int tn; int r, c, l; int m, n, k; - int mn, mk, kn; + int mn; double myc; - __shared__ int our_params[7]; - double *buff; + const double *buff_l, *buff_r; + + int psp, c_loc; /* Setup shared memory. */ - buff = (double *) cache; + //buff = (double *) cache; /* Determine who I am. */ sp = blockIdx.x; tn = threadIdx.x; - nt = blockDim.x; - /* Load in the parameters. */ - for (l = 0; l <= 6/nt; l++) { - r = tn+nt*l; - if (r < 7) - our_params[r] = param_stack[7*sp+r]; - } + psp = 7*sp; + m = param_stack[psp]; + n = param_stack[psp+1]; + k = param_stack[psp+2]; - syncthreads(); - m = our_params[0]; - n = our_params[1]; - k = our_params[2]; - - /* Load in the buffers. */ - mk = m*k; - kn = k*n; - for (l = 0; l <= (mk-1) / nt; l++) { - r = tn+nt*l; - if (r < mk) - buff[r] = a_data[our_params[3]-1+r]; - } - for (l = 0; l <= (kn-1) / nt; l++) { - r = tn+nt*l; - if (r < kn) - buff[mk+r] = b_data[our_params[4]-1+r]; - } + buff_l = &(a_data[param_stack[psp+3]-1]); + buff_r = &(b_data[param_stack[psp+4]-1]); /* Calculate who I am. */ - syncthreads(); mn = m*n; @@ -89,16 +193,17 @@ __global__ void stack_mm_d for (l = 0; l < k; l++) { myc = myc + - buff[ l*m+r] * - buff[mk+c*k+l]; + buff_l[ l*m+r] * + buff_r[ c*k+l]; } } /* Lock the C block. */ + c_id = param_stack[psp+6]-1; + c_loc = param_stack[psp+5]-1; syncthreads(); if (tn == 0) { sp_one = sp + 1; - c_id = our_params[6]-1; lock_owner = 0; while ((lock_owner != sp_one)) lock_owner = atomicCAS (&(c_locks[c_id]), 0, sp_one); @@ -107,7 +212,7 @@ __global__ void stack_mm_d /* Add our results to the C block. */ syncthreads(); if (tn < mn) { - c_data[our_params[5]-1+tn] += myc; + c_data[c_loc+tn] += myc; } /* Release the lock on the C block. */ @@ -119,3 +224,203 @@ __global__ void stack_mm_d }; + +__global__ void stack_mm_mnk_d_direct ( + const int *__restrict__ param_stack, + const int careful, const int nruns, + const int m, const int n, const int k, const int mn, + const double *__restrict__ a_data, + const double *__restrict__ b_data, + double *__restrict__ c_data, + int *__restrict__ c_locks) { + + /** + * \var sp which stack member this thread block is processing + (= CUDA thread block) + * \var psp pointer to first element of parameters + * \var c_loc pointer to C data + * \var run run number + * \var nrun number of runs + * \var my_id my ID for locking + * \var tn thread number (of CUDA thread block) + * \var mn product of the block dimensions + * \var l multiplication loop index + * \var c, r C matrix row, column of this thread + * \var myc C matrix accumulator + * \var buff_l cache for A data + * \var buff_r cache for B data + * \var c_id translated C block number (used in locking) + * \var lock_owner current C block owner (used in locking) + */ + + int lock_owner, c_id, my_id; + int l; + const int r = threadIdx.x % m; + const int c = threadIdx.x / m; + double myc, tmp; + const double *buff_l, *buff_r; + + int psp, c_loc; + + int run, nrun; + + nrun = GROUPING; + if (blockIdx.x == careful) + nrun = nruns; + + for (run = 0; run < nrun; run ++) { + psp = 7*(blockIdx.x*GROUPING + run); + + buff_l = &(a_data[param_stack[psp+3]-1]); + buff_r = &(b_data[param_stack[psp+4]-1]); + /* Do actual multiplication. */ + if (threadIdx.x < mn) { + myc = 0.0l; + + for (l = 0; l < k; l++) { + myc = myc + + buff_l[ l*m+r] * + buff_r[ c*k+l]; + } + } + + /* Lock the C block. */ + c_loc = param_stack[psp+5]-1; + syncthreads(); + c_id = param_stack[psp+6]-1; + + if (threadIdx.x == 0) { + my_id = blockIdx.x+1; + lock_owner = 0; + while ((lock_owner != my_id)) + lock_owner = atomicCAS (&(c_locks[c_id]), 0, my_id); + } else if (threadIdx.x==1) { + tmp = c_data[c_loc]; + } + + + /* Add our results to the C block. */ + syncthreads(); + if (threadIdx.x < mn) { + c_data[c_loc+threadIdx.x] += myc; + } + + /* Release the lock on the C block. */ + syncthreads(); + if (threadIdx.x == 0) { + c_locks[c_id] = 0; + //threadfence(); + } + } + + +}; + + +__global__ void stack_mm_mnk_vec_d ( + const int *__restrict__ param_stack, + const int stack_size, const int nmat, + const int m, const int n, const int k, const int mn, + const double *__restrict__ a_data, + const double *__restrict__ b_data, + double *__restrict__ c_data, + int *__restrict__ c_locks) { + + /** + * \var sp which stack member this thread block is processing + (= CUDA thread block) + * \var psp pointer to first element of parameters + * \var c_loc pointer to C data + * \var run run number + * \var nrun number of runs + * \var my_id translated stack (=sp+1) + * \var tn thread number (of CUDA thread block) + * \var mn product of the block dimensions + * \var l multiplication loop index + * \var c, r C matrix row, column of this thread + * \var myc C matrix accumulator + * \var buff_l cache for A data + * \var buff_r cache for B data + * \var c_id translated C block number (used in locking) + * \var lock_owner current C block owner (used in locking) + */ + + int lock_owner, c_id, my_id; + const int tn = threadIdx.x; + int nmat_used; + int nt; + const int r = threadIdx.x % m; + int c, l; + double myc[32]; + double mya[32]; + __shared__ int our_b[32]; + const double *buff_l, *buff_r; + + int psp, c_loc; + int run, nrun; + const int my_mat_num = threadIdx.x / m; + int imat; + + //nrun = GROUPING; + //if ((blockIdx.x+1) * GROUPING > stack_size) + // nrun = stack_size - (blockIdx.x)*GROUPING; + + nmat_used = nmat; + if ((blockIdx.x+1)*nmat > stack_size) + nmat_used = stack_size - (blockIdx.x)*nmat; + nt = m * nmat_used; + + //for (run = 0; run < nrun; run ++) { + //sp = blockIdx.x*GROUPING + run; + + psp = 7*(blockIdx.x*nmat + my_mat_num); + + buff_l = &(a_data[param_stack[psp+3]-1]); + buff_r = &(b_data[param_stack[psp+4]-1]); + + /* Do actual multiplication. */ + if (tn < nt) { + for (l = 0; l < k; l++) { + mya[l] = buff_l[ l*m + r ]; + } + for (c = 0; c < n; c++) { + if (tn < k) + our_b[l] = buff_r[c*k+tn]; + syncthreads(); + myc[c] = 0.0l; + + for (l = 0; l < k; l++) { + myc[c] = myc[c] + + mya [ l ] * + our_b [ l ]; + //buff_r[ c*k+l]; + } + } + } + /* Lock the C block. */ + c_id = param_stack[psp+6]-1; + syncthreads(); + c_loc = param_stack[psp+5]-1; + my_id = blockIdx.x + 1; + for (imat = 0; imat < nmat_used; imat++) { + if (r == 0 && imat == my_mat_num) { + lock_owner = 0; + while ((lock_owner != my_id)) + lock_owner = atomicCAS (&(c_locks[c_id]), 0, my_id); + } + + /* Add our results to the C block. */ + syncthreads(); + if (tn < nt && imat == my_mat_num) { + for (c = 0; c < n; c++) { + c_data[c_loc+r+c*m] += myc[c]; + } + } + + /* Release the lock on the C block. */ + syncthreads(); + if (r == 0 && imat == my_mat_num) { + c_locks[c_id] = 0; + } + } +}; diff --git a/cuda_tools/dbcsr_cuda_dev.cu b/cuda_tools/dbcsr_cuda_dev.cu new file mode 100644 index 0000000000..d3d40586cf --- /dev/null +++ b/cuda_tools/dbcsr_cuda_dev.cu @@ -0,0 +1,36 @@ +/****************************************************************************** + * CP2K: A general program to perform molecular dynamics simulations + * Copyright (C) 2000 - 2011 Urban Borstnik and the CP2K developers group + *****************************************************************************/ + +#include +#include +#include + +#include "dbcsr_cuda.h" + +static const int verbose_print = 0; + +extern "C" int dc_thread_sync_cu() { + cudaError_t cErr; + + cErr = cudaThreadSynchronize (); + if (cuda_error_check (cErr)) return 1; + return 0; +} + +extern "C" int dc_set_device_cu(int device_id) { + cudaError_t cErr; + + cErr = cudaSetDevice(device_id); + if (cuda_error_check (cErr)) return 1; + return 0; +} + +extern "C" int dc_get_ndevices_cu(int *n_devices) { + cudaError_t cErr; + + cErr = cudaGetDeviceCount(n_devices); + if (cuda_error_check (cErr)) return 1; + return 0; +} diff --git a/src/cp2k_runs.F b/src/cp2k_runs.F index d95d2006e9..026e9bdbfc 100644 --- a/src/cp2k_runs.F +++ b/src/cp2k_runs.F @@ -17,7 +17,8 @@ MODULE cp2k_runs enable_color_tags USE cp_dbcsr_interface, ONLY: cp_dbcsr_config,& cp_dbcsr_finalize_lib,& - cp_dbcsr_init_lib + cp_dbcsr_init_lib,& + cp_dbcsr_print_config USE cp_files, ONLY: close_file,& open_file USE cp_ma_interface, ONLY: cp_ma_finalize_lib,& @@ -178,7 +179,7 @@ CONTAINS CALL cp_para_env_create(para_env, group=mpi_comm,& owns_group=.FALSE.,error=error) - CALL cp_dbcsr_init_lib (error=error) + CALL cp_dbcsr_init_lib (group=mpi_comm, error=error) IF (has_ma) THEN CALL cp_ma_init_lib(para_env,error=error) @@ -287,6 +288,13 @@ CONTAINS CALL cuda_device_mem_init(root_section, error) CALL cp_dbcsr_config (root_section, error) + IF (output_unit > 0 .AND.& + cp_logger_would_log(logger, cp_note_level)) THEN + CALL cp_dbcsr_print_config (unit_nr=output_unit, error=error) + WRITE (UNIT=output_unit,FMT='()') + ENDIF + + SELECT CASE (prog_name_id) diff --git a/src/cp_dbcsr_interface.F b/src/cp_dbcsr_interface.F index 352fa14e29..cfd9890ac3 100644 --- a/src/cp_dbcsr_interface.F +++ b/src/cp_dbcsr_interface.F @@ -54,15 +54,15 @@ MODULE cp_dbcsr_interface dbcsr_get_conf_combtypes, dbcsr_get_conf_comm_thread_load, & dbcsr_get_conf_cuda_mem, dbcsr_get_conf_mm_driver, & dbcsr_get_conf_mm_stacksize, dbcsr_get_conf_mpi_mem, & - dbcsr_get_conf_subcomm, dbcsr_get_conf_use_comm_thread, & - dbcsr_set_conf_combtypes, dbcsr_set_conf_comm_thread_load, & - dbcsr_set_conf_cuda_mem, dbcsr_set_conf_mm_driver, & - dbcsr_set_conf_mm_stacksize, dbcsr_set_conf_mpi_mem, & - dbcsr_set_conf_nstacks, dbcsr_set_conf_subcomm, & - dbcsr_set_conf_use_comm_thread, detailed_timing, has_cuda, has_mpi, & - kernel_timing, mm_driver_blas, mm_driver_cuda, mm_driver_matmul, & - mm_driver_plasma, mm_driver_smm, mm_name_blas, mm_name_cuda, & - mm_name_matmul, mm_name_plasma, mm_name_smm + dbcsr_get_conf_nstacks, dbcsr_get_conf_subcomm, & + dbcsr_get_conf_use_comm_thread, dbcsr_set_conf_combtypes, & + dbcsr_set_conf_comm_thread_load, dbcsr_set_conf_cuda_mem, & + dbcsr_set_conf_mm_driver, dbcsr_set_conf_mm_stacksize, & + dbcsr_set_conf_mpi_mem, dbcsr_set_conf_nstacks, & + dbcsr_set_conf_subcomm, dbcsr_set_conf_use_comm_thread, & + detailed_timing, has_cuda, has_mpi, kernel_timing, mm_driver_blas, & + mm_driver_cuda, mm_driver_matmul, mm_driver_plasma, mm_driver_smm, & + mm_name_blas, mm_name_cuda, mm_name_matmul, mm_name_plasma, mm_name_smm USE dbcsr_data_methods, ONLY: & dbcsr_data_clear_pointer, dbcsr_data_init, dbcsr_data_new, & dbcsr_data_release, dbcsr_data_set_pointer, dbcsr_scalar, & @@ -434,7 +434,8 @@ CONTAINS ! ***************************************************************************** !> \brief Initializes DBCSR ! ***************************************************************************** - SUBROUTINE cp_dbcsr_init_lib (error) + SUBROUTINE cp_dbcsr_init_lib (group, error) + INTEGER, INTENT(IN) :: group TYPE(cp_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'cp_dbcsr_init_lib', & @@ -442,7 +443,7 @@ CONTAINS TYPE(dbcsr_error_type) :: dbcsr_error - CALL dbcsr_init_lib (dbcsr_error) + CALL dbcsr_init_lib (group, dbcsr_error) END SUBROUTINE cp_dbcsr_init_lib ! ***************************************************************************** @@ -502,11 +503,8 @@ CONTAINS CALL section_vals_val_get(dbcsr_section,& "kernel_timing", l_val=kernel_timing, error=error) CALL section_vals_val_get(dbcsr_section,& - "n_size_m_stacks", i_val=nstacks(1), error=error) - CALL section_vals_val_get(dbcsr_section,& - "n_size_n_stacks", i_val=nstacks(2), error=error) - CALL section_vals_val_get(dbcsr_section,& - "n_size_k_stacks", i_val=nstacks(3), error=error) + "n_size_mnk_stacks", i_val=nstacks(1), error=error) + nstacks(2:3) = nstacks(1) CALL section_vals_val_get(dbcsr_section,& "n_stack_buffers", i_val=n_stack_buffers, error=error) CALL section_vals_val_get(dbcsr_section,& @@ -543,14 +541,18 @@ CONTAINS CHARACTER(len=default_string_length) :: comm_thread_load_str, mm_name, & mm_ss_str, use_combtypes_str, use_comm_thread_str, use_cuda_mem_str, & - use_mpi_mem_str, use_subcomms_str + use_k_stacks, use_m_stacks, use_mem_regions, use_mpi_mem_str, & + use_n_stacks, use_stack_buffers, use_subcomms_str INTEGER :: comm_thread_load, mm_driver, & - mm_ss, unit_num + mm_ss, nbuffers, nmemregions, & + unit_num + INTEGER, DIMENSION(3) :: n_mnk_stacks LOGICAL :: use_combtypes, & use_comm_thread, & use_cuda_mem, use_mpi_mem, & use_subcomms TYPE(cp_logger_type), POINTER :: logger + TYPE(dbcsr_error_type) :: dbcsr_error mm_driver = dbcsr_get_conf_mm_driver () use_subcomms = dbcsr_get_conf_subcomm () @@ -578,6 +580,14 @@ CONTAINS use_cuda_mem_str = l2str_r(use_cuda_mem,info_len) use_comm_thread_str = l2str_r(use_comm_thread,info_len) comm_thread_load_str = int2str_r(comm_thread_load,info_len) + + CALL dbcsr_get_conf_nstacks (n_mnk_stacks, nbuffers, nmemregions,& + error=dbcsr_error) + use_stack_buffers = int2str_r (nbuffers, info_len) + use_mem_regions = int2str_r (nmemregions, info_len) + use_m_stacks = int2str_r (n_mnk_stacks(1), info_len) + use_n_stacks = int2str_r (n_mnk_stacks(2), info_len) + use_k_stacks = int2str_r (n_mnk_stacks(3), info_len) logger => cp_error_get_logger(error) IF (PRESENT (unit_nr)) THEN @@ -589,6 +599,29 @@ CONTAINS WRITE(UNIT=unit_num, FMT=o_fmt) & plabel, "Multiplication driver", ADJUSTR(mm_name(1:info_len)),& plabel, "Multiplication stack size", mm_ss_str(1:info_len) + IF (nmemregions .NE. 1) & + WRITE(UNIT=unit_num, FMT=o_fmt) & + plabel, "Multiplication stack memory regions",& + use_mem_regions(1:info_len) + IF (nbuffers .NE. 1) & + WRITE(UNIT=unit_num, FMT=o_fmt) & + plabel, "Multiplication stack buffers",& + use_stack_buffers(1:info_len) + IF (ALL(n_mnk_stacks .EQ. n_mnk_stacks(1))) THEN + WRITE(UNIT=unit_num, FMT=o_fmt) & + plabel, "Multiplication size stacks",& + use_m_stacks(1:info_len) + ELSE + WRITE(UNIT=unit_num, FMT=o_fmt) & + plabel, "Multiplication size m stacks",& + use_m_stacks(1:info_len) + WRITE(UNIT=unit_num, FMT=o_fmt) & + plabel, "Multiplication size n stacks",& + use_n_stacks(1:info_len) + WRITE(UNIT=unit_num, FMT=o_fmt) & + plabel, "Multiplication size k stacks",& + use_k_stacks(1:info_len) + ENDIF IF (has_mpi) & WRITE(UNIT=unit_num, FMT=o_fmt) & plabel, "Use subcommunicators", use_subcomms_str(1:info_len),& diff --git a/src/dbcsr_lib/OBJECTDEFS b/src/dbcsr_lib/OBJECTDEFS index 75036c17d5..391f1f9330 100644 --- a/src/dbcsr_lib/OBJECTDEFS +++ b/src/dbcsr_lib/OBJECTDEFS @@ -9,6 +9,7 @@ LIB3_OBJECTS = array_types.o\ dbcsr_block_operations.o\ dbcsr_c_mpi_calls.o\ dbcsr_config.o\ + dbcsr_cuda_device.o\ dbcsr_cuda_memory.o\ dbcsr_cuda_methods.o\ dbcsr_cuda_operations.o\ diff --git a/src/dbcsr_lib/dbcsr_config.F b/src/dbcsr_lib/dbcsr_config.F index 28f5c47e85..790e71f8f0 100644 --- a/src/dbcsr_lib/dbcsr_config.F +++ b/src/dbcsr_lib/dbcsr_config.F @@ -65,6 +65,7 @@ MODULE dbcsr_config ! PUBLIC :: detailed_timing, kernel_timing + PUBLIC :: is_configured ! First the constants are declared. @@ -116,37 +117,39 @@ MODULE dbcsr_config ! by calling the dbcsr_init_conf() subroutine. ! Allocates subcommunicators for process rows and columns. - LOGICAL :: use_subcommunicators = .FALSE. + LOGICAL, SAVE :: use_subcommunicators = .FALSE. ! Use combined data types for MPI transfers. - LOGICAL :: use_combined_types = .FALSE. + LOGICAL, SAVE :: use_combined_types = .FALSE. ! Use MPI-allocated memory. - LOGICAL :: use_MPI_memory = has_MPI + LOGICAL, SAVE :: use_MPI_memory = has_MPI ! Use CUDA host-pinned memory. - LOGICAL :: use_CUDA_host_pinned_memory = .FALSE. + LOGICAL, SAVE :: use_CUDA_host_pinned_memory = .FALSE. ! Which driver to use for matrix multiplications. - INTEGER :: mm_driver = mm_driver_smm + INTEGER, SAVE :: mm_driver = mm_driver_smm ! Stack size to use for multiplication parameters - INTEGER :: mm_stack_size = 1000 + INTEGER, SAVE :: mm_stack_size = 1000 ! Whether to print extra timing - LOGICAL :: detailed_timing = .FALSE. - LOGICAL :: kernel_timing = .FALSE. + LOGICAL, SAVE :: detailed_timing = .FALSE. + LOGICAL, SAVE :: kernel_timing = .FALSE. ! Number of stacks to use - INTEGER :: nm_stacks = 0 - INTEGER :: nn_stacks = 0 - INTEGER :: nk_stacks = 0 - INTEGER :: nstackbuffers = 1 - INTEGER :: nstackmemregions = 1 + INTEGER, SAVE :: nm_stacks = 0 + INTEGER, SAVE :: nn_stacks = 0 + INTEGER, SAVE :: nk_stacks = 0 + INTEGER, SAVE :: nstackbuffers = 1 + INTEGER, SAVE :: nstackmemregions = 1 ! Configuration of an MPI progress thread - LOGICAL :: use_comm_thread = .TRUE. - INTEGER :: comm_thread_load = 100 + LOGICAL, SAVE :: use_comm_thread = .TRUE. + INTEGER, SAVE :: comm_thread_load = 100 + + LOGICAL, SAVE :: is_configured = .FALSE. CONTAINS @@ -171,7 +174,11 @@ CONTAINS mm_stack_size = 1000 IF (has_cuda) THEN mm_driver = mm_driver_cuda - mm_stack_size = 10000 + mm_stack_size = 30000 + nm_stacks = 2 + nn_stacks = 2 + nk_stacks = 2 + nstackbuffers = 4 ENDIF ! use_comm_thread = .TRUE. diff --git a/src/dbcsr_lib/dbcsr_cuda_device.F b/src/dbcsr_lib/dbcsr_cuda_device.F new file mode 100644 index 0000000000..116b3b7583 --- /dev/null +++ b/src/dbcsr_lib/dbcsr_cuda_device.F @@ -0,0 +1,236 @@ +!-----------------------------------------------------------------------------! +! CP2K: A general program to perform molecular dynamics simulations ! +! Copyright (C) 2000 - 2011 Urban Borstnik and the CP2K developers group ! +!-----------------------------------------------------------------------------! + +! ***************************************************************************** +!> \brief CUDA device support for DBCSR +!> \author Urban Borstnik +!> \date 2011-09-22 +!> \version 1.0 +!> +!> Modification history: +!> - Created 2011-09-22 +! ***************************************************************************** +MODULE dbcsr_cuda_device +#if !defined (__HAS_NO_ISO_C_BINDING) + USE ISO_C_BINDING +#endif + USE dbcsr_cuda_methods, ONLY: dbcsr_cuda_dev_mem_get_type + USE dbcsr_cuda_types, ONLY: dbcsr_cuda_mem_type,& + dbcsr_cuda_mem_type_c4,& + dbcsr_cuda_mem_type_c8,& + dbcsr_cuda_mem_type_i4,& + dbcsr_cuda_mem_type_i8,& + dbcsr_cuda_mem_type_r4,& + dbcsr_cuda_mem_type_r8 + USE dbcsr_data_methods, ONLY: dbcsr_data_get_size_referenced,& + dbcsr_data_get_type,& + dbcsr_get_data,& + dbcsr_get_data_p_c,& + dbcsr_get_data_p_d,& + dbcsr_get_data_p_s,& + dbcsr_get_data_p_z + USE dbcsr_error_handling + USE dbcsr_kinds, ONLY: int_4,& + int_4_size,& + int_8,& + int_8_size,& + real_4,& + real_4_size,& + real_8,& + real_8_size + USE dbcsr_types, ONLY: dbcsr_data_obj,& + dbcsr_type_complex_4,& + dbcsr_type_complex_8,& + dbcsr_type_int_4,& + dbcsr_type_int_8,& + dbcsr_type_real_4,& + dbcsr_type_real_8 + USE dummy_c_bindings + + IMPLICIT NONE + + PRIVATE + + CHARACTER(len=*), PARAMETER, PRIVATE :: moduleN = 'dbcsr_cuda_device' + + LOGICAL, PARAMETER :: careful_mod = .TRUE. + + + PUBLIC :: dbcsr_cuda_init, dbcsr_cuda_thread_sync,& + dbcsr_cuda_get_n_devices + +#if defined (__DBCSR_CUDA) + INTERFACE + FUNCTION cuda_memcpy_h2d_cu(host, dev, count, async_type) RESULT (istat) & + BIND(C, name="dc_memcpy_h2d_cu") + USE ISO_C_BINDING + TYPE(C_PTR), INTENT(IN), VALUE :: host + TYPE(C_PTR), VALUE :: dev + INTEGER(KIND=C_SIZE_T), INTENT(IN), & + VALUE :: count + INTEGER(KIND=C_INT), INTENT(IN), VALUE :: async_type + INTEGER(KIND=C_INT) :: istat + + END FUNCTION cuda_memcpy_h2d_cu + END INTERFACE + + INTERFACE + FUNCTION cuda_memcpy_d2h_cu(dev, host, count, async_type) RESULT (istat) & + BIND(C, name="dc_memcpy_d2h_cu") + USE ISO_C_BINDING + TYPE(C_PTR), INTENT(IN), VALUE :: dev + TYPE(C_PTR), VALUE :: host + INTEGER(KIND=C_SIZE_T), INTENT(IN), & + VALUE :: count + INTEGER(KIND=C_INT), INTENT(IN), VALUE :: async_type + INTEGER(KIND=C_INT) :: istat + + END FUNCTION cuda_memcpy_d2h_cu + END INTERFACE + + + INTERFACE + FUNCTION cuda_do_stack_cu(param_stack, stack_size, nparams,& + data_type,& + a_data, b_data, c_data, c_locks, m_max, n_max, k_max)& + RESULT (istat) & + BIND(C, name="dc_do_stack_cu") + USE ISO_C_BINDING + TYPE(C_PTR), INTENT(IN), VALUE :: param_stack + INTEGER(KIND=C_INT), INTENT(IN), VALUE :: stack_size, nparams, data_type + TYPE(C_PTR), INTENT(IN), VALUE :: a_data, b_data + TYPE(C_PTR), VALUE :: c_data, c_locks + INTEGER(KIND=C_INT), VALUE :: m_max, n_max, k_max + INTEGER(KIND=C_INT) :: istat + + END FUNCTION cuda_do_stack_cu + END INTERFACE + + + INTERFACE + FUNCTION cuda_set_device_cu (device_id) RESULT (istat) & + BIND(C, name="dc_set_device_cu") + USE ISO_C_BINDING + INTEGER(KIND=C_INT), INTENT(IN), VALUE :: device_id + INTEGER(KIND=C_INT) :: istat + + END FUNCTION cuda_set_device_cu + END INTERFACE + + INTERFACE + FUNCTION cuda_get_ndevices_cu (n_devices) RESULT (istat) & + BIND(C, name="dc_get_ndevices_cu") + USE ISO_C_BINDING + INTEGER(KIND=C_INT), INTENT(OUT) :: n_devices + INTEGER(KIND=C_INT) :: istat + + END FUNCTION cuda_get_ndevices_cu + END INTERFACE + + + INTERFACE + FUNCTION cuda_thread_sync_cu() RESULT (istat) BIND(C, name="dc_thread_sync_cu") + USE ISO_C_BINDING + INTEGER(KIND=C_INT) :: istat + + END FUNCTION cuda_thread_sync_cu + END INTERFACE +#endif + +CONTAINS + + SUBROUTINE dbcsr_cuda_init (card_num, error) + INTEGER, INTENT(IN), OPTIONAL :: card_num + TYPE(dbcsr_error_type), INTENT(INOUT) :: error + + CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_cuda_init', & + routineP = moduleN//':'//routineN + + INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: icard + +! --------------------------------------------------------------------------- + + CALL dbcsr_error_set (routineN, error_handle, error) +#if defined (__DBCSR_CUDA) + IF (PRESENT (card_num)) THEN + icard = card_num + ELSE + + ENDIF + istat = cuda_set_device_cu (icard) +#else + istat = -1 +#endif + IF (istat /= 0) THEN + CALL dbcsr_assert (istat, "EQ", 0,& + dbcsr_fatal_level, dbcsr_internal_error, routineN,& + "Error selecting GPU device.",& + __LINE__, error=error) + ENDIF + CALL dbcsr_error_stop (error_handle, error) + END SUBROUTINE dbcsr_cuda_init + + + FUNCTION dbcsr_cuda_get_n_devices (error) RESULT (n_devices) + TYPE(dbcsr_error_type), INTENT(INOUT) :: error + INTEGER :: n_devices + + CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_cuda_get_n_devices', & + routineP = moduleN//':'//routineN + + INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: ndev + +! --------------------------------------------------------------------------- + + CALL dbcsr_error_set (routineN, error_handle, error) +#if defined (__DBCSR_CUDA) + istat = cuda_get_ndevices_cu (ndev) + n_devices = INT (ndev) +#else + istat = -1 + n_devices = 0 +#endif + IF (istat /= 0) THEN + CALL dbcsr_assert (istat, "EQ", 0,& + dbcsr_fatal_level, dbcsr_internal_error, routineN,& + "Error getting device count",& + __LINE__, error=error) + ENDIF + CALL dbcsr_error_stop (error_handle, error) + END FUNCTION dbcsr_cuda_get_n_devices + + + SUBROUTINE dbcsr_cuda_thread_sync(error) + TYPE(dbcsr_error_type), INTENT(INOUT), & + OPTIONAL :: error + + CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_cuda_thread_sync', & + routineP = moduleN//':'//routineN + + INTEGER :: error_handle, istat + +! --------------------------------------------------------------------------- + + CALL dbcsr_error_set (routineN, error_handle, error) +#if defined (__DBCSR_CUDA) + istat = cuda_thread_sync_cu(); +#else + istat = -1 +#endif + IF (istat /= 0) THEN + IF (PRESENT(error)) THEN + CALL dbcsr_assert (istat, "EQ", 0,& + dbcsr_fatal_level, dbcsr_internal_error, routineN,& + "Could not synchronize all threads",& + __LINE__, error=error) + ENDIF + ENDIF + CALL dbcsr_error_stop (error_handle, error) + END SUBROUTINE dbcsr_cuda_thread_sync + + +END MODULE dbcsr_cuda_device diff --git a/src/dbcsr_lib/dbcsr_cuda_operations.F b/src/dbcsr_lib/dbcsr_cuda_operations.F index 2a16418bc4..7c5a1a65bc 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations.F +++ b/src/dbcsr_lib/dbcsr_cuda_operations.F @@ -62,8 +62,6 @@ MODULE dbcsr_cuda_operations PUBLIC :: dbcsr_cuda_do_mm_stack - PUBLIC :: dbcsr_cuda_thread_sync - INTERFACE dbcsr_cuda_do_mm_stack MODULE PROCEDURE do_mm_stack_any @@ -125,7 +123,7 @@ MODULE dbcsr_cuda_operations INTERFACE FUNCTION cuda_do_stack_cu(param_stack, stack_size, nparams,& data_type,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max)& + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk)& RESULT (istat) & BIND(C, name="dc_do_stack_cu") USE ISO_C_BINDING @@ -133,66 +131,29 @@ MODULE dbcsr_cuda_operations INTEGER(KIND=C_INT), INTENT(IN), VALUE :: stack_size, nparams, data_type TYPE(C_PTR), INTENT(IN), VALUE :: a_data, b_data TYPE(C_PTR), VALUE :: c_data, c_locks - INTEGER(KIND=C_INT), VALUE :: m_max, n_max, k_max + INTEGER(KIND=C_INT), INTENT(IN), VALUE :: m_max, n_max, k_max, def_mnk INTEGER(KIND=C_INT) :: istat END FUNCTION cuda_do_stack_cu END INTERFACE - - - INTERFACE - FUNCTION cuda_thread_sync_cu() RESULT (istat) BIND(C, name="dc_thread_sync_cu") - USE ISO_C_BINDING - INTEGER(KIND=C_INT) :: istat - - END FUNCTION cuda_thread_sync_cu - END INTERFACE #endif + CONTAINS - SUBROUTINE dbcsr_cuda_thread_sync(error) - TYPE(dbcsr_error_type), INTENT(INOUT), & - OPTIONAL :: error - - CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_cuda_thread_sync', & - routineP = moduleN//':'//routineN - - INTEGER :: error_handle, istat - -! --------------------------------------------------------------------------- - - CALL dbcsr_error_set (routineN, error_handle, error) -#if defined (__DBCSR_CUDA) - istat = cuda_thread_sync_cu(); -#else - istat = -1 -#endif - IF (istat /= 0) THEN - IF (PRESENT(error)) THEN - CALL dbcsr_assert (istat, "EQ", 0,& - dbcsr_fatal_level, dbcsr_internal_error, routineN,& - "Could not synchronize all threads",& - __LINE__, error=error) - ENDIF - ENDIF - CALL dbcsr_error_stop (error_handle, error) - END SUBROUTINE dbcsr_cuda_thread_sync - - - !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!1111!!!!!!!!!!!!!!!!!! ! Encapsulated data routines SUBROUTINE do_mm_stack_any (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type), INTENT(IN) :: a_data, b_data TYPE(dbcsr_cuda_mem_type), INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_any', & @@ -217,19 +178,19 @@ CONTAINS CASE (dbcsr_type_real_4) CALL dbcsr_cuda_do_mm_stack (param_stack, stack_size, nparams,& a_data%d_r, b_data%d_r, c_data%d_r,& - c_locks, m_max, n_max, k_max, error=error) + c_locks, m_max, n_max, k_max, def_mnk, error=error) CASE (dbcsr_type_real_8) CALL dbcsr_cuda_do_mm_stack (param_stack, stack_size, nparams,& a_data%d_d, b_data%d_d, c_data%d_d,& - c_locks, m_max, n_max, k_max, error=error) + c_locks, m_max, n_max, k_max, def_mnk, error=error) CASE (dbcsr_type_complex_4) CALL dbcsr_cuda_do_mm_stack (param_stack, stack_size, nparams,& a_data%d_c, b_data%d_c, c_data%d_c,& - c_locks, m_max, n_max, k_max, error=error) + c_locks, m_max, n_max, k_max, def_mnk, error=error) CASE (dbcsr_type_complex_8) CALL dbcsr_cuda_do_mm_stack (param_stack, stack_size, nparams,& a_data%d_z, b_data%d_z, c_data%d_z,& - c_locks, m_max, n_max, k_max, error=error) + c_locks, m_max, n_max, k_max, def_mnk, error=error) CASE default CALL dbcsr_assert (.FALSE.,& dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,& diff --git a/src/dbcsr_lib/dbcsr_cuda_operations__nametype1_.template b/src/dbcsr_lib/dbcsr_cuda_operations__nametype1_.template index 63c2aaff38..62086158c9 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations__nametype1_.template +++ b/src/dbcsr_lib/dbcsr_cuda_operations__nametype1_.template @@ -4,7 +4,7 @@ !-----------------------------------------------------------------------------! SUBROUTINE do_mm_stack_[nametype1] (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type_[shorttype1]), INTENT(IN) :: a_data, b_data @@ -12,22 +12,30 @@ INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_[nametype1]', & routineP = moduleN//':'//routineN INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: mnk ! --------------------------------------------------------------------------- CALL dbcsr_error_set (routineN, error_handle, error) #if defined (__DBCSR_CUDA) + IF (def_mnk) THEN + mnk = 1 + ELSE + mnk = 0 + ENDIF istat = cuda_do_stack_cu(param_stack%ref, INT(stack_size, KIND=C_INT),& INT(nparams, KIND=C_INT),& INT(dbcsr_type_[fulltype1], KIND=C_INT),& a_data%ref, b_data%ref, c_data%ref, c_locks%ref,& - INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT)) + INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT),& + mnk) #else istat = -1 #endif diff --git a/src/dbcsr_lib/dbcsr_cuda_operations_c.F b/src/dbcsr_lib/dbcsr_cuda_operations_c.F index 7c38cb89a6..5d6eac6d52 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations_c.F +++ b/src/dbcsr_lib/dbcsr_cuda_operations_c.F @@ -4,7 +4,7 @@ !-----------------------------------------------------------------------------! SUBROUTINE do_mm_stack_c (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type_c4), INTENT(IN) :: a_data, b_data @@ -12,22 +12,30 @@ INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_c', & routineP = moduleN//':'//routineN INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: mnk ! --------------------------------------------------------------------------- CALL dbcsr_error_set (routineN, error_handle, error) #if defined (__DBCSR_CUDA) + IF (def_mnk) THEN + mnk = 1 + ELSE + mnk = 0 + ENDIF istat = cuda_do_stack_cu(param_stack%ref, INT(stack_size, KIND=C_INT),& INT(nparams, KIND=C_INT),& INT(dbcsr_type_complex_4, KIND=C_INT),& a_data%ref, b_data%ref, c_data%ref, c_locks%ref,& - INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT)) + INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT),& + mnk) #else istat = -1 #endif diff --git a/src/dbcsr_lib/dbcsr_cuda_operations_d.F b/src/dbcsr_lib/dbcsr_cuda_operations_d.F index 2dceeaf957..c1078f6aa2 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations_d.F +++ b/src/dbcsr_lib/dbcsr_cuda_operations_d.F @@ -4,7 +4,7 @@ !-----------------------------------------------------------------------------! SUBROUTINE do_mm_stack_d (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type_r8), INTENT(IN) :: a_data, b_data @@ -12,22 +12,30 @@ INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_d', & routineP = moduleN//':'//routineN INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: mnk ! --------------------------------------------------------------------------- CALL dbcsr_error_set (routineN, error_handle, error) #if defined (__DBCSR_CUDA) + IF (def_mnk) THEN + mnk = 1 + ELSE + mnk = 0 + ENDIF istat = cuda_do_stack_cu(param_stack%ref, INT(stack_size, KIND=C_INT),& INT(nparams, KIND=C_INT),& INT(dbcsr_type_real_8, KIND=C_INT),& a_data%ref, b_data%ref, c_data%ref, c_locks%ref,& - INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT)) + INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT),& + mnk) #else istat = -1 #endif diff --git a/src/dbcsr_lib/dbcsr_cuda_operations_i.F b/src/dbcsr_lib/dbcsr_cuda_operations_i.F index 88b6fb89c1..47dc5dc445 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations_i.F +++ b/src/dbcsr_lib/dbcsr_cuda_operations_i.F @@ -4,7 +4,7 @@ !-----------------------------------------------------------------------------! SUBROUTINE do_mm_stack_i (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: a_data, b_data @@ -12,22 +12,30 @@ INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_i', & routineP = moduleN//':'//routineN INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: mnk ! --------------------------------------------------------------------------- CALL dbcsr_error_set (routineN, error_handle, error) #if defined (__DBCSR_CUDA) + IF (def_mnk) THEN + mnk = 1 + ELSE + mnk = 0 + ENDIF istat = cuda_do_stack_cu(param_stack%ref, INT(stack_size, KIND=C_INT),& INT(nparams, KIND=C_INT),& INT(dbcsr_type_int_4, KIND=C_INT),& a_data%ref, b_data%ref, c_data%ref, c_locks%ref,& - INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT)) + INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT),& + mnk) #else istat = -1 #endif diff --git a/src/dbcsr_lib/dbcsr_cuda_operations_l.F b/src/dbcsr_lib/dbcsr_cuda_operations_l.F index 046613b6db..98a2cf0bdf 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations_l.F +++ b/src/dbcsr_lib/dbcsr_cuda_operations_l.F @@ -4,7 +4,7 @@ !-----------------------------------------------------------------------------! SUBROUTINE do_mm_stack_l (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type_i8), INTENT(IN) :: a_data, b_data @@ -12,22 +12,30 @@ INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_l', & routineP = moduleN//':'//routineN INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: mnk ! --------------------------------------------------------------------------- CALL dbcsr_error_set (routineN, error_handle, error) #if defined (__DBCSR_CUDA) + IF (def_mnk) THEN + mnk = 1 + ELSE + mnk = 0 + ENDIF istat = cuda_do_stack_cu(param_stack%ref, INT(stack_size, KIND=C_INT),& INT(nparams, KIND=C_INT),& INT(dbcsr_type_int_8, KIND=C_INT),& a_data%ref, b_data%ref, c_data%ref, c_locks%ref,& - INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT)) + INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT),& + mnk) #else istat = -1 #endif diff --git a/src/dbcsr_lib/dbcsr_cuda_operations_r.F b/src/dbcsr_lib/dbcsr_cuda_operations_r.F index 46ffcabb3d..9f4ff81e4f 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations_r.F +++ b/src/dbcsr_lib/dbcsr_cuda_operations_r.F @@ -4,7 +4,7 @@ !-----------------------------------------------------------------------------! SUBROUTINE do_mm_stack_r (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type_r4), INTENT(IN) :: a_data, b_data @@ -12,22 +12,30 @@ INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_r', & routineP = moduleN//':'//routineN INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: mnk ! --------------------------------------------------------------------------- CALL dbcsr_error_set (routineN, error_handle, error) #if defined (__DBCSR_CUDA) + IF (def_mnk) THEN + mnk = 1 + ELSE + mnk = 0 + ENDIF istat = cuda_do_stack_cu(param_stack%ref, INT(stack_size, KIND=C_INT),& INT(nparams, KIND=C_INT),& INT(dbcsr_type_real_4, KIND=C_INT),& a_data%ref, b_data%ref, c_data%ref, c_locks%ref,& - INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT)) + INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT),& + mnk) #else istat = -1 #endif diff --git a/src/dbcsr_lib/dbcsr_cuda_operations_z.F b/src/dbcsr_lib/dbcsr_cuda_operations_z.F index e9f4bd054e..6767b83ebd 100644 --- a/src/dbcsr_lib/dbcsr_cuda_operations_z.F +++ b/src/dbcsr_lib/dbcsr_cuda_operations_z.F @@ -4,7 +4,7 @@ !-----------------------------------------------------------------------------! SUBROUTINE do_mm_stack_z (param_stack, stack_size, nparams,& - a_data, b_data, c_data, c_locks, m_max, n_max, k_max, error) + a_data, b_data, c_data, c_locks, m_max, n_max, k_max, def_mnk, error) TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: param_stack INTEGER, INTENT(IN) :: stack_size, nparams TYPE(dbcsr_cuda_mem_type_c8), INTENT(IN) :: a_data, b_data @@ -12,22 +12,30 @@ INTENT(INOUT) :: c_data TYPE(dbcsr_cuda_mem_type_i4), INTENT(IN) :: c_locks INTEGER, INTENT(IN) :: m_max, n_max, k_max + LOGICAL, INTENT(IN) :: def_mnk TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'do_mm_stack_z', & routineP = moduleN//':'//routineN INTEGER :: error_handle, istat + INTEGER(KIND=C_INT) :: mnk ! --------------------------------------------------------------------------- CALL dbcsr_error_set (routineN, error_handle, error) #if defined (__DBCSR_CUDA) + IF (def_mnk) THEN + mnk = 1 + ELSE + mnk = 0 + ENDIF istat = cuda_do_stack_cu(param_stack%ref, INT(stack_size, KIND=C_INT),& INT(nparams, KIND=C_INT),& INT(dbcsr_type_complex_8, KIND=C_INT),& a_data%ref, b_data%ref, c_data%ref, c_locks%ref,& - INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT)) + INT(m_max,KIND=C_INT), INT(n_max,KIND=C_INT), INT(k_max,KIND=C_INT),& + mnk) #else istat = -1 #endif diff --git a/src/dbcsr_lib/dbcsr_internal_operations.F b/src/dbcsr_lib/dbcsr_internal_operations.F index 35d6e7032b..d4c90437da 100644 --- a/src/dbcsr_lib/dbcsr_internal_operations.F +++ b/src/dbcsr_lib/dbcsr_internal_operations.F @@ -29,10 +29,19 @@ MODULE dbcsr_internal_operations mm_driver_blas, mm_driver_cuda, mm_driver_matmul, mm_driver_plasma, & mm_driver_smm, mm_stack_size, use_CUDA_host_pinned_memory, & use_MPI_memory, use_combined_types, use_comm_thread - USE dbcsr_cuda_memory - USE dbcsr_cuda_methods - USE dbcsr_cuda_operations - USE dbcsr_cuda_types + USE dbcsr_cuda_device, ONLY: dbcsr_cuda_thread_sync + USE dbcsr_cuda_memory, ONLY: dbcsr_cuda_dev_mem_alloc,& + dbcsr_cuda_dev_mem_dealloc,& + dbcsr_cuda_dev_mem_hold,& + dbcsr_cuda_dev_mem_new,& + dbcsr_cuda_dev_mem_realloc,& + dbcsr_cuda_dev_mem_release,& + dbcsr_cuda_dev_mem_zero + USE dbcsr_cuda_methods, ONLY: dbcsr_cuda_dev_mem_get_alloc + USE dbcsr_cuda_operations, ONLY: dbcsr_cuda_cp_dev_to_host,& + dbcsr_cuda_cp_host_to_dev,& + dbcsr_cuda_do_mm_stack + USE dbcsr_cuda_types, ONLY: dbcsr_cuda_mem_type USE dbcsr_data_methods, ONLY: & dbcsr_data_clear_pointer, dbcsr_data_ensure_size, dbcsr_data_get_size, & dbcsr_data_get_size_referenced, dbcsr_data_get_type, dbcsr_data_init, & @@ -3325,6 +3334,7 @@ CONTAINS param_stack%state, param_stack%t%t%stack_state_dev,& param_stack%m, param_stack%n, param_stack%k,& param_stack%max_m, param_stack%max_n, param_stack%max_k,& + param_stack%defined_mnk,& param_stack%t%t%c_locks_dev, param_stack%t%t%params_dev,& error=error) IF (careful_mod) & @@ -3376,7 +3386,7 @@ CONTAINS IF (zero_last .GT. c_size) THEN !WRITE(*,*)routineN//" reallocating c_dev",& ! c_size, zero_last - IF (.FALSE.) THEN + IF (detailed_timing) THEN t_dev_sync = t_dev_sync - m_walltime() CALL dbcsr_cuda_thread_sync(error=error) t_dev_sync = t_dev_sync + m_walltime() @@ -3399,7 +3409,7 @@ CONTAINS ! Resize block count. IF (nblks .GT. dbcsr_cuda_dev_mem_get_alloc(c_locks_dev)) THEN maxs = dbcsr_cuda_dev_mem_get_alloc(c_locks_dev) - IF (.FALSE.) THEN + IF (detailed_timing) THEN t_dev_sync = t_dev_sync - m_walltime() CALL dbcsr_cuda_thread_sync(error=error) t_dev_sync = t_dev_sync + m_walltime() @@ -3452,7 +3462,7 @@ CONTAINS product_data_card, has_product_data_card,& a_dev, b_dev,& zero_first, zero_last, nblks, state, stack_state_dev,& - m, n, k, max_m, max_n, max_k,& + m, n, k, max_m, max_n, max_k, defined_mnk,& c_locks_dev, params_dev, & error) INTEGER, INTENT(INOUT) :: stack_size @@ -3470,6 +3480,7 @@ CONTAINS INTEGER, POINTER :: state TYPE(dbcsr_cuda_mem_type), POINTER :: stack_state_dev INTEGER, INTENT(IN) :: m, n, k, max_m, max_n, max_k + LOGICAL, INTENT(IN) :: defined_mnk TYPE(dbcsr_cuda_mem_type), POINTER :: c_locks_dev, params_dev TYPE(dbcsr_error_type), INTENT(inout) :: error @@ -3634,7 +3645,7 @@ CONTAINS a_dev, b_dev, product_data_card,& c_locks_dev,& params_dev,& - m, n, k, max_m, max_n, max_k,& + m, n, k, max_m, max_n, max_k, defined_mnk,& state, stack_state_dev,& error=error) CASE default @@ -4226,7 +4237,7 @@ CONTAINS data_a_dev, data_b_dev, data_c_dev,& c_locks,& params_dev,& - m, n, k, max_m, max_n, max_k,& + m, n, k, max_m, max_n, max_k, defined_mnk,& state, stack_state_dev,& error) INTEGER, INTENT(IN) :: stack_size @@ -4237,6 +4248,7 @@ CONTAINS TYPE(dbcsr_cuda_mem_type), INTENT(INOUT) :: data_c_dev, c_locks, & params_dev INTEGER, INTENT(IN) :: m, n, k, max_m, max_n, max_k + LOGICAL, INTENT(IN) :: defined_mnk INTEGER, POINTER :: state TYPE(dbcsr_cuda_mem_type), INTENT(IN) :: stack_state_dev TYPE(dbcsr_error_type), INTENT(INOUT) :: error @@ -4307,7 +4319,7 @@ CONTAINS IF (kernel_timing) kt = -m_walltime() CALL dbcsr_cuda_do_mm_stack (params_dev%d_i, stack_size, n_mult_params,& data_a_dev, data_b_dev, data_c_dev,& - c_locks%d_i, ABS(m_max), ABS(n_max), ABS(k_max),& + c_locks%d_i, ABS(m_max), ABS(n_max), ABS(k_max), defined_mnk,& error=error) IF (kernel_timing) THEN CALL dbcsr_cuda_thread_sync (error=error) @@ -4456,6 +4468,7 @@ CONTAINS param_stack%state, param_stack%t%t%stack_state_dev,& param_stack%m, param_stack%n, param_stack%k,& param_stack%max_m, param_stack%max_n, param_stack%max_k,& + param_stack%defined_mnk,& param_stack%t%t%c_locks_dev, param_stack%t%t%params_dev, & error=error) !$OMP CRITICAL (crit_target) diff --git a/src/dbcsr_lib/dbcsr_operations.F b/src/dbcsr_lib/dbcsr_operations.F index 65ca9b9a23..b7873fd64f 100644 --- a/src/dbcsr_lib/dbcsr_operations.F +++ b/src/dbcsr_lib/dbcsr_operations.F @@ -36,9 +36,13 @@ MODULE dbcsr_operations set_block2d_diagonal USE dbcsr_config, ONLY: dbcsr_init_conf,& detailed_timing,& + has_cuda,& + is_configured,& mm_driver,& mm_driver_cuda,& mm_driver_plasma + USE dbcsr_cuda_device, ONLY: dbcsr_cuda_get_n_devices,& + dbcsr_cuda_init USE dbcsr_cuda_memory, ONLY: dbcsr_cuda_dev_mem_alloc,& dbcsr_cuda_dev_mem_dealloc USE dbcsr_cuda_methods, ONLY: dbcsr_cuda_dev_mem_setup @@ -81,6 +85,7 @@ MODULE dbcsr_operations USE dbcsr_message_passing, ONLY: dmp_max,& mp_allgather,& mp_bcast,& + mp_environ,& mp_recv,& mp_send,& mp_sum @@ -4455,16 +4460,19 @@ CONTAINS !> Prepares the DBCSR library for use. !> \param[in,out] error error ! ***************************************************************************** - SUBROUTINE dbcsr_init_lib (error) + SUBROUTINE dbcsr_init_lib (group, error) + INTEGER, INTENT(IN) :: group TYPE(dbcsr_error_type), INTENT(INOUT) :: error CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_init_lib', & routineP = moduleN//':'//routineN - INTEGER :: error_handle + INTEGER :: error_handle, mynode, ngpus, & + numnode ! --------------------------------------------------------------------------- + IF (is_configured) RETURN CALL dbcsr_error_set(routineN, error_handle, error) ! CALL dbcsr_assert (int_1_size, "EQ", 1,& @@ -4486,6 +4494,13 @@ CONTAINS ! CALL dbcsr_init_conf (error) ! + IF (has_cuda) THEN + CALL mp_environ (numnode, mynode, group) + ngpus = dbcsr_cuda_get_n_devices (error) + CALL dbcsr_cuda_init (card_num=MOD(mynode, ngpus), error=error) + ENDIF + ! + is_configured = .TRUE. CALL dbcsr_error_stop (error_handle, error) END SUBROUTINE dbcsr_init_lib diff --git a/src/dbcsr_lib/dbcsr_util.F b/src/dbcsr_lib/dbcsr_util.F index c1080e23ff..b78a498009 100644 --- a/src/dbcsr_lib/dbcsr_util.F +++ b/src/dbcsr_lib/dbcsr_util.F @@ -1432,7 +1432,7 @@ CONTAINS size_counts(array(i)) = size_counts(array(i)) - 1 END DO IF (SIZE(array) .GT. 0) THEN - CALL sort (size_counts, max_val_l, permutation) + CALL sort (size_counts, max_val_l+1, permutation) ENDIF nmc = MIN(nmost_common, SIZE(size_counts)) ! Determine the biggest block size and allocate the map. @@ -1440,11 +1440,11 @@ CONTAINS ! Create the mapping from block size to order. most_common_map = nmost_common+1 DO i = 1, nmc - most_common_map(permutation(i-1)) = i + most_common_map(permutation(i-1)-1) = i END DO ! Copy the most common elements most_common_elements(:) = 0 - most_common_elements(1:nmc) = permutation(0:nmc-1) + most_common_elements(1:nmc) = permutation(0:nmc-1)-1 DEALLOCATE (size_counts) DEALLOCATE (permutation) END SUBROUTINE map_most_common diff --git a/src/environment.F b/src/environment.F index 55306e1109..1c1a741397 100644 --- a/src/environment.F +++ b/src/environment.F @@ -18,8 +18,7 @@ MODULE environment compile_arch, compile_date, compile_host, compile_lastcvs, cp2k_home, & cp2k_version, cp2k_year, get_runtime_info, id_cp2k_version, & r_host_name, r_pid, r_user_name - USE cp_dbcsr_interface, ONLY: cp_dbcsr_print_config,& - trash_map + USE cp_dbcsr_interface, ONLY: trash_map USE cp_files, ONLY: close_file,& open_file USE cp_ma_interface @@ -819,14 +818,6 @@ CONTAINS !$omp end parallel ENDIF - IF (output_unit > 0) THEN - IF (.TRUE.) THEN - WRITE (UNIT=output_unit,FMT='()') - CALL cp_dbcsr_print_config (unit_nr=output_unit, error=error) - WRITE (UNIT=output_unit,FMT='()') - ENDIF - ENDIF - CALL cp_print_key_finished_output(output_unit,logger,global_section,& "PROGRAM_RUN_INFO", error=error) diff --git a/src/f77_interface.F b/src/f77_interface.F index ce33c9c231..e0abaae20f 100644 --- a/src/f77_interface.F +++ b/src/f77_interface.F @@ -227,7 +227,7 @@ CONTAINS CALL add_all_references() ! Initialize the DBCSR configuration - CALL cp_dbcsr_init_lib (error=error) + CALL cp_dbcsr_init_lib (group=mpi_comm_default, error=error) ELSE ierr=cp_failure_level END IF diff --git a/src/input_cp2k.F b/src/input_cp2k.F index 7463d6c0ec..5a8189325e 100644 --- a/src/input_cp2k.F +++ b/src/input_cp2k.F @@ -841,28 +841,13 @@ CONTAINS CALL keyword_release(keyword,error=error) ! CALL dbcsr_get_conf_nstacks (nstacks, n_buffers, n_mem_regions, dbcsr_error) - CALL keyword_create(keyword, name="n_size_m_stacks",& - description="Number of stacks to use for distinct m (row) sizes" & + CALL keyword_create(keyword, name="n_size_mnk_stacks",& + description="Number of stacks to use for distinct atomic sizes" & // " (e.g., 2 for a system of mostly waters).",& - usage="n_size_m_stacks 2",& + usage="n_size_mnk_stacks 2",& default_i_val=nstacks(1),error=error) CALL section_add_keyword(section,keyword,error=error) CALL keyword_release(keyword,error=error) - CALL keyword_create(keyword, name="n_size_n_stacks",& - description="Number of stacks to use for distinct n (col) sizes" & - // " (e.g., 2 for a system of mostly waters).",& - usage="n_size_n_stacks 2",& - default_i_val=nstacks(2),error=error) - CALL section_add_keyword(section,keyword,error=error) - CALL keyword_release(keyword,error=error) - CALL keyword_create(keyword, name="n_size_k_stacks",& - description="Number of stacks to use for distinct k" & - // " (common row and column) sizes" & - // " (e.g., 2 for a system of mostly waters).",& - usage="n_size_k_stacks 2",& - default_i_val=nstacks(2),error=error) - CALL section_add_keyword(section,keyword,error=error) - CALL keyword_release(keyword,error=error) CALL keyword_create(keyword, name="n_stack_buffers",& description="Number of stack buffers to use" & // " (e.g., 2 when using GPUs)",& diff --git a/tests/LIBTEST/dbcsr_multistack.inp b/tests/LIBTEST/dbcsr_multistack.inp index 0ce19574ff..936af3d16b 100644 --- a/tests/LIBTEST/dbcsr_multistack.inp +++ b/tests/LIBTEST/dbcsr_multistack.inp @@ -6,9 +6,7 @@ THRESHOLD 0.00000000001 &END &DBCSR - n_size_m_stacks 1 - n_size_n_stacks 3 - n_size_k_stacks 2 + n_size_mnk_stacks 2 n_stack_buffers 4 n_stack_memory_regions 5 &END DBCSR