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

DBCSR/CUDA work.

Browse source 
Allows multiple GPUs per process; Fix for multiple stack binning-simplifies input; CUDA optimizations; DBCSR configuration output after setup;


svn-origin-rev: 11794
This commit is contained in:
Urban Borštnik 2011-09-26 13:19:07 +00:00
parent 270b55eb01
commit d2993ca2bb
26 changed files with 977 additions and 256 deletions
Download patch file Download diff file Expand all files Collapse all files

1
cuda_tools/NVOBJECTDEFS
View file

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

7
cuda_tools/README
View file

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

3
cuda_tools/dbcsr_cuda.h
View file

@ -13,3 +13,6 @@
int cuda_error_check (cudaError_t cudaError);
#define GROUPING 16

215
cuda_tools/dbcsr_cuda_calc.cu
View file

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

377
cuda_tools/dbcsr_cuda_calc_d.cu
View file

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

36
cuda_tools/dbcsr_cuda_dev.cu Normal file
View file

@ -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 <cuda_runtime.h>
#include <stdio.h>
#include <sm_11_atomic_functions.h>
#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;
}

12
src/cp2k_runs.F
View file

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

69
src/cp_dbcsr_interface.F
View file

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

1
src/dbcsr_lib/OBJECTDEFS
View file

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

39
src/dbcsr_lib/dbcsr_config.F
View file

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

236
src/dbcsr_lib/dbcsr_cuda_device.F Normal file
View file

@ -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
!>
!> <b>Modification history:</b>
!> - 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

57
src/dbcsr_lib/dbcsr_cuda_operations.F
View file

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

12
src/dbcsr_lib/dbcsr_cuda_operations__nametype1_.template
View file

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

12
src/dbcsr_lib/dbcsr_cuda_operations_c.F
View file

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

12
src/dbcsr_lib/dbcsr_cuda_operations_d.F
View file

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

12
src/dbcsr_lib/dbcsr_cuda_operations_i.F
View file

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

12
src/dbcsr_lib/dbcsr_cuda_operations_l.F
View file

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

12
src/dbcsr_lib/dbcsr_cuda_operations_r.F
View file

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

12
src/dbcsr_lib/dbcsr_cuda_operations_z.F
View file

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

33
src/dbcsr_lib/dbcsr_internal_operations.F
View file

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

19
src/dbcsr_lib/dbcsr_operations.F
View file

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

6
src/dbcsr_lib/dbcsr_util.F
View file

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

11
src/environment.F
View file

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

2
src/f77_interface.F
View file

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

21
src/input_cp2k.F
View file

@ -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)",&

4
tests/LIBTEST/dbcsr_multistack.inp
View file

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