refactoring of the lower level of the DBCSR multiply (Ole Schuett).

svn-origin-rev: 12737
This commit is contained in:
Joost VandeVondele 2013-03-04 11:24:50 +00:00
parent 860e88d4cf
commit 235b7094af
36 changed files with 2565 additions and 2879 deletions

View file

@ -39,25 +39,25 @@ __global__ void stack_mm_r
int stack_size, int nparams,
const float *__restrict__ a_data,
const float *__restrict__ b_data,
float *__restrict__ c_data, int *__restrict__ c_locks);
float *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset);
__global__ void stack_mm_d
(const int *__restrict__ param_stack,
int stack_size, int nparams,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data, int *__restrict__ c_locks);
double *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset);
__global__ void stack_mm_c
(const int *__restrict__ param_stack,
int stack_size, int nparams,
const float *__restrict__ a_data,
const float *__restrict__ b_data,
float *__restrict__ c_data, int *__restrict__ c_locks);
float *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset);
__global__ void stack_mm_z
(const int *__restrict__ param_stack,
int stack_size, int nparams,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data, int *__restrict__ c_locks);
double *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset);
__global__ void stack_mm_mnk_d (const int *__restrict__ param_stack,
const int careful, const int nruns,
@ -66,7 +66,7 @@ __global__ void stack_mm_mnk_d (const int *__restrict__ param_stack,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data,
int *__restrict__ c_locks);
int *__restrict__ c_locks, int lock_offset);
__global__ void stack_mm_mnk_sq23_d (const int *__restrict__ param_stack,
const int careful, const int nruns,
@ -75,7 +75,8 @@ __global__ void stack_mm_mnk_sq23_d (const int *__restrict__ param_stack,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data,
int *__restrict__ c_locks);
int *__restrict__ c_locks,
int lock_offset);
__global__ void stack_mm_mnk_sq5_d (const int *__restrict__ param_stack,
const int careful, const int nruns,
@ -84,7 +85,8 @@ __global__ void stack_mm_mnk_sq5_d (const int *__restrict__ param_stack,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data,
int *__restrict__ c_locks);
int *__restrict__ c_locks,
int lock_offset);
/**
* \brief Bridge routine to call appropriate CUDA kernel.
@ -129,7 +131,8 @@ dc_do_stack_cu (int *param_stack, int stack_size, int nparams,
return 4;
stack_mm_r <<< stack_size, maxt, shared_size, stream >>>
(param_stack, stack_size, nparams,
(float *) a_data, (float *) b_data, (float *) c_data, c_locks);
(float *) a_data, (float *) b_data, (float *) c_data, c_locks,
pow (2, 15) * stream_id);
break;
case 3:
/* Real, double precision */
@ -164,7 +167,9 @@ dc_do_stack_cu (int *param_stack, int stack_size, int nparams,
shared_size, stream >>> (param_stack, careful, nruns, m_max,
n_max, k_max, liter,
(double *) a_data, (double *) b_data,
(double *) c_data, c_locks);
(double *) c_data, c_locks, pow (2,
15)
* stream_id);
}
else if (m_max == 5 && n_max == 5 && k_max == 5)
{
@ -177,7 +182,8 @@ dc_do_stack_cu (int *param_stack, int stack_size, int nparams,
m_max, n_max, k_max, liter,
(double *) a_data,
(double *) b_data,
(double *) c_data, c_locks);
(double *) c_data, c_locks,
pow (2, 15) * stream_id);
}
else
{
@ -187,7 +193,9 @@ dc_do_stack_cu (int *param_stack, int stack_size, int nparams,
//mn, mk, nk, maxb, liter,
liter,
(double *) a_data, (double *) b_data,
(double *) c_data, c_locks);
(double *) c_data, c_locks, pow (2,
15)
* stream_id);
}
}
else
@ -200,7 +208,7 @@ dc_do_stack_cu (int *param_stack, int stack_size, int nparams,
stack_mm_d <<< stack_size, maxt, shared_size, stream >>>
(param_stack, stack_size, nparams,
(double *) a_data, (double *) b_data, (double *) c_data,
c_locks);
c_locks, pow (2, 15) * stream_id);
}
break;
case 5:
@ -210,7 +218,8 @@ dc_do_stack_cu (int *param_stack, int stack_size, int nparams,
return 4;
stack_mm_c <<< stack_size, maxt, shared_size, stream >>>
(param_stack, stack_size, nparams,
(float *) a_data, (float *) b_data, (float *) c_data, c_locks);
(float *) a_data, (float *) b_data, (float *) c_data, c_locks,
pow (2, 15) * stream_id);
break;
case 7:
/* Complex, double precision */
@ -219,7 +228,8 @@ dc_do_stack_cu (int *param_stack, int stack_size, int nparams,
return 4;
stack_mm_z <<< stack_size, maxt, shared_size, stream >>>
(param_stack, stack_size, nparams,
(double *) a_data, (double *) b_data, (double *) c_data, c_locks);
(double *) a_data, (double *) b_data, (double *) c_data, c_locks,
pow (2, 15) * stream_id);
break;
default:
return 2;

View file

@ -12,7 +12,7 @@ __global__ void stack_mm_c
int stack_size, int nparams,
const float *__restrict__ a_data,
const float *__restrict__ b_data,
float *__restrict__ c_data, int *__restrict__ c_locks)
float *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset)
{
/**
@ -117,7 +117,7 @@ __global__ void stack_mm_c
// /* Lock the C block. */
// syncthreads();
// if (tn == 0) {
// sp_one = sp + 1;
// sp_one = lock_offset + sp + 1;
// c_id = our_params[6]-1;
// lock_owner = 0;
// while ((lock_owner != sp_one))

View file

@ -50,7 +50,8 @@ stack_mm_mnk_sq23_d (const int *__restrict__ param_stack,
const int liter,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data, int *__restrict__ c_locks)
double *__restrict__ c_data, int *__restrict__ c_locks,
int lock_offset)
{
/**
@ -227,7 +228,7 @@ stack_mm_mnk_sq23_d (const int *__restrict__ param_stack,
c_loc = param_stack[psp + 5] - 1;
// c_id = param_stack[psp+6]-1;
// if (threadIdx.x == 0) {
// my_id = blockIdx.x+1;
// my_id = lock_offset + blockIdx.x+1;
// lock_owner = 0;
// while ((lock_owner != my_id))
// lock_owner = atomicCAS (&(c_locks[c_id]), 0, my_id);
@ -347,7 +348,8 @@ stack_mm_mnk_sq5_d (const int *__restrict__ param_stack,
const int liter,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data, int *__restrict__ c_locks)
double *__restrict__ c_data, int *__restrict__ c_locks,
int lock_offset)
{
/**
@ -519,7 +521,8 @@ stack_mm_mnk_sq5_d (const int *__restrict__ param_stack,
/*
if (lid == 0) {
my_id = 4*blockIdx.x + wid + 1;
// TODO: make shure 2^16 is enough as lock_offset stepsize
my_id = lock_offset + 4*blockIdx.x + wid + 1;
int lock_owner = 0;
while ((lock_owner != my_id))
lock_owner = atomicCAS (&(c_locks[c_id]), 0, my_id);
@ -556,7 +559,8 @@ stack_mm_mnk_d (const int *__restrict__ param_stack,
const int liter,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data, int *__restrict__ c_locks)
double *__restrict__ c_data, int *__restrict__ c_locks,
int lock_offset)
{
/**
@ -646,7 +650,7 @@ stack_mm_mnk_d (const int *__restrict__ param_stack,
// c_id = param_stack[psp+6]-1;
//
// if (threadIdx.x == 0) {
// my_id = blockIdx.x+1;
// my_id = lock_offset + blockIdx.x+1;
// lock_owner = 0;
// while ((lock_owner != my_id))
// lock_owner = atomicCAS (&(c_locks[c_id]), 0, my_id);
@ -682,7 +686,7 @@ __global__ void stack_mm_d
int stack_size, int nparams,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data, int *__restrict__ c_locks)
double *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset)
{
/**
@ -754,7 +758,7 @@ __global__ void stack_mm_d
// c_loc = param_stack[psp+5]-1;
// syncthreads();
// if (tn == 0) {
// sp_one = sp + 1;
// sp_one = lock_offset + sp + 1;
// lock_owner = 0;
// while ((lock_owner != sp_one))
// lock_owner = atomicCAS (&(c_locks[c_id]), 0, sp_one);

View file

@ -12,7 +12,7 @@ __global__ void stack_mm_r
int stack_size, int nparams,
const float *__restrict__ a_data,
const float *__restrict__ b_data,
float *__restrict__ c_data, int *__restrict__ c_locks)
float *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset)
{
/**
@ -104,7 +104,7 @@ __global__ void stack_mm_r
// /* Lock the C block. */
// syncthreads();
// if (tn == 0) {
// sp_one = sp + 1;
// sp_one = lock_offset + sp + 1;
// c_id = our_params[6]-1;
// lock_owner = 0;
// while ((lock_owner != sp_one))

View file

@ -12,7 +12,7 @@ __global__ void stack_mm_z
int stack_size, int nparams,
const double *__restrict__ a_data,
const double *__restrict__ b_data,
double *__restrict__ c_data, int *__restrict__ c_locks)
double *__restrict__ c_data, int *__restrict__ c_locks, int lock_offset)
{
/**
@ -121,7 +121,7 @@ __global__ void stack_mm_z
// /* Lock the C block. */
// syncthreads();
// if (tn == 0) {
// sp_one = sp + 1;
// sp_one = lock_offset + sp + 1;
// c_id = our_params[6]-1;
// lock_owner = 0;
// while ((lock_owner != sp_one))

View file

@ -56,6 +56,9 @@ dc_set_device_cu (int device_id)
if (cuda_error_check (cErr))
return 1;
if (myDevice != device_id)
return 1;
cErr = cudaGetDeviceProperties (&devProperties, myDevice);
if (cuda_error_check (cErr))
return 1;

View file

@ -280,3 +280,15 @@ dc_memzero_cu (void *dev_mem, size_t offset, size_t length)
if (cuda_error_check (cudaGetLastError())) return 1; */
return 0;
}
extern "C" int
dc_dev_mem_info_cu (size_t * free, size_t * avail)
{
cudaError_t cErr;
cErr = cudaMemGetInfo (free, avail);
if (cuda_error_check (cErr))
return 1;
if (cuda_error_check (cudaGetLastError ()))
return 1;
return 0;
}

View file

@ -46,8 +46,7 @@ MODULE cp_dbcsr_interface
dbcsr_set_conf_nstacks, dbcsr_set_conf_subcomm, &
dbcsr_set_conf_use_comm_thread, has_cuda, has_mpi, 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, mm_thread_workshare
mm_name_blas, mm_name_cuda, mm_name_matmul, mm_name_plasma, mm_name_smm
USE dbcsr_data_methods, ONLY: dbcsr_scalar,&
dbcsr_scalar_fill_all,&
dbcsr_scalar_get_value,&
@ -433,8 +432,7 @@ CONTAINS
INTEGER :: comm_thread_load, &
error_handle, mm_driver, &
mm_ss, n_mem_regions, &
n_stack_buffers
mm_ss, n_stack_buffers
INTEGER, DIMENSION(3) :: nstacks
LOGICAL :: use_combtypes, &
use_comm_thread, &
@ -463,15 +461,11 @@ CONTAINS
"mm_stack_size", i_val=mm_ss, error=error)
CALL section_vals_val_get(dbcsr_section,&
"mm_driver", i_val=mm_driver, error=error)
CALL section_vals_val_get(dbcsr_section,&
"mm_thread_workshare", l_val=mm_thread_workshare, error=error)
CALL section_vals_val_get(dbcsr_section,&
"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,&
"n_stack_memory_regions", i_val=n_mem_regions, error=error)
CALL section_vals_val_get(dbcsr_section,&
"use_comm_thread", l_val=use_comm_thread, error=error)
CALL section_vals_val_get(dbcsr_section,&
@ -482,8 +476,7 @@ CONTAINS
CALL dbcsr_set_conf_mpi_mem (use_mpi_mem, error=dbcsr_error)
CALL dbcsr_set_conf_cuda_mem (use_cuda_mem, error=dbcsr_error)
CALL dbcsr_set_conf_mm_stacksize (mm_ss, error=dbcsr_error)
CALL dbcsr_set_conf_nstacks(nstacks, n_stack_buffers, n_mem_regions,&
error=dbcsr_error)
CALL dbcsr_set_conf_nstacks(nstacks, n_stack_buffers, error=dbcsr_error)
CALL dbcsr_set_conf_use_comm_thread(use_comm_thread, error=dbcsr_error)
CALL dbcsr_set_conf_comm_thread_load(comm_thread_load, error=dbcsr_error)
CALL dbcsr_error_stop (error_handle, dbcsr_error)
@ -504,11 +497,10 @@ 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_k_stacks, use_m_stacks, use_mem_regions, use_mpi_mem_str, &
use_n_stacks, use_stack_buffers, use_subcomms_str
use_k_stacks, use_m_stacks, use_mpi_mem_str, use_n_stacks, &
use_stack_buffers, use_subcomms_str
INTEGER :: comm_thread_load, mm_driver, &
mm_ss, nbuffers, nmemregions, &
unit_num
mm_ss, nbuffers, unit_num
INTEGER, DIMENSION(3) :: n_mnk_stacks
LOGICAL :: use_combtypes, &
use_comm_thread, &
@ -544,10 +536,8 @@ CONTAINS
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)
CALL dbcsr_get_conf_nstacks (n_mnk_stacks, nbuffers, 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)
@ -562,10 +552,6 @@ CONTAINS
WRITE(UNIT=unit_num, FMT=o_fmt) &
plabel, "Multiplication driver", 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",&

View file

@ -34,8 +34,6 @@ LIB3_OBJECTS =\
dbcsr_operations.o\
dbcsr_performance_multiply.o\
dbcsr_plasma_interface.o\
dbcsr_pq_methods.o\
dbcsr_pq_types.o\
dbcsr_ptr_util.o\
dbcsr_test_add.o\
dbcsr_test_methods.o\

View file

@ -69,12 +69,11 @@ MODULE dbcsr_config
mm_name_smm,&
mm_name_plasma,&
mm_name_cuda
PUBLIC :: nm_stacks, nn_stacks, nk_stacks, nstackbuffers, nstackmemregions
PUBLIC :: nm_stacks, nn_stacks, nk_stacks, nstackbuffers
PUBLIC :: use_comm_thread, comm_thread_load
!
PUBLIC :: is_configured
PUBLIC :: mm_thread_workshare
! First the constants are declared.
@ -156,16 +155,13 @@ MODULE dbcsr_config
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, SAVE :: use_comm_thread = .TRUE.
INTEGER, SAVE :: comm_thread_load = 100
LOGICAL, SAVE :: is_configured = .FALSE.
! Worksharing between threads
LOGICAL, SAVE :: mm_thread_workshare = .TRUE.
CONTAINS
@ -201,16 +197,12 @@ CONTAINS
!
use_comm_thread = .TRUE.
comm_thread_load = 100
mm_thread_workshare = .FALSE.
!$ mm_thread_workshare = .TRUE.
mm_async = driver_is_async (mm_driver)
IF (mm_async) THEN
mm_host_driver = mm_driver_smm
mm_thread_workshare = .FALSE.
ELSE
mm_host_driver = mm_driver
ENDIF
IF (has_cuda) mm_thread_workshare = .FALSE.
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE dbcsr_init_conf
@ -390,8 +382,8 @@ CONTAINS
END FUNCTION dbcsr_get_conf_mm_stacksize
SUBROUTINE dbcsr_set_conf_nstacks (nstacks, nbuffers, nmemregs, error)
INTEGER, INTENT(IN) :: nstacks(:), nbuffers, nmemregs
SUBROUTINE dbcsr_set_conf_nstacks (nstacks, nbuffers, error)
INTEGER, INTENT(IN) :: nstacks(:), nbuffers
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_set_conf_nstacks', &
@ -407,12 +399,11 @@ CONTAINS
nn_stacks = nstacks(2)
nk_stacks = nstacks(3)
nstackbuffers = nbuffers
nstackmemregions = nmemregs
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE dbcsr_set_conf_nstacks
SUBROUTINE dbcsr_get_conf_nstacks (nstacks, nbuffers, nmemregs, error)
INTEGER, INTENT(OUT) :: nstacks(:), nbuffers, nmemregs
SUBROUTINE dbcsr_get_conf_nstacks (nstacks, nbuffers, error)
INTEGER, INTENT(OUT) :: nstacks(:), nbuffers
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_get_conf_nstacks', &
@ -428,7 +419,6 @@ CONTAINS
nstacks(2) = nn_stacks
nstacks(3) = nk_stacks
nbuffers = nstackbuffers
nmemregs = nstackmemregions
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE dbcsr_get_conf_nstacks

View file

@ -149,7 +149,7 @@ CONTAINS
IF (PRESENT (card_num)) THEN
icard = card_num
ELSE
icard = 0
ENDIF
istat = cuda_set_device_cu (icard)
#else
@ -212,14 +212,7 @@ CONTAINS
#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
IF (istat /= 0) STOP "dbcsr_cuda_device_sync: Could not synchronize all threads"
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dbcsr_cuda_device_sync
@ -241,14 +234,7 @@ CONTAINS
#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 stream",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) STOP "dbcsr_cuda_stream_sync: Could not synchronize stream"
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dbcsr_cuda_stream_sync
@ -271,14 +257,7 @@ CONTAINS
#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 create streams",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) STOP "dbcsr_cuda_create_streams: Could not create streams"
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dbcsr_cuda_create_streams
@ -299,14 +278,7 @@ CONTAINS
#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 destroy streams",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) STOP "dbcsr_cuda_destroy_streams: Could not destroy streams"
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dbcsr_cuda_destroy_streams

View file

@ -61,7 +61,7 @@ MODULE dbcsr_cuda_memory
PUBLIC :: dbcsr_cuda_dev_mem_init
PUBLIC :: dbcsr_cuda_dev_mem_new, dbcsr_cuda_dev_mem_release,&
dbcsr_cuda_dev_mem_hold
PUBLIC :: dbcsr_cuda_dev_mem_info
INTERFACE dbcsr_cuda_dev_mem_alloc
MODULE PROCEDURE dev_mem_alloc_any
@ -88,12 +88,24 @@ MODULE dbcsr_cuda_memory
MODULE PROCEDURE host_mem_alloc_i, host_mem_alloc_l
MODULE PROCEDURE host_mem_alloc_r, host_mem_alloc_d
MODULE PROCEDURE host_mem_alloc_c, host_mem_alloc_z
MODULE PROCEDURE host_mem_alloc_i_2D, host_mem_alloc_l_2D
MODULE PROCEDURE host_mem_alloc_r_2D, host_mem_alloc_d_2D
MODULE PROCEDURE host_mem_alloc_c_2D, host_mem_alloc_z_2D
MODULE PROCEDURE host_mem_alloc_i_4D, host_mem_alloc_l_4D
MODULE PROCEDURE host_mem_alloc_r_4D, host_mem_alloc_d_4D
MODULE PROCEDURE host_mem_alloc_c_4D, host_mem_alloc_z_4D
END INTERFACE
INTERFACE dbcsr_cuda_host_mem_dealloc
MODULE PROCEDURE host_mem_dealloc_i, host_mem_dealloc_l
MODULE PROCEDURE host_mem_dealloc_r, host_mem_dealloc_d
MODULE PROCEDURE host_mem_dealloc_c, host_mem_dealloc_z
MODULE PROCEDURE host_mem_dealloc_i_2D, host_mem_dealloc_l_2D
MODULE PROCEDURE host_mem_dealloc_r_2D, host_mem_dealloc_d_2D
MODULE PROCEDURE host_mem_dealloc_c_2D, host_mem_dealloc_z_2D
MODULE PROCEDURE host_mem_dealloc_i_4D, host_mem_dealloc_l_4D
MODULE PROCEDURE host_mem_dealloc_r_4D, host_mem_dealloc_d_4D
MODULE PROCEDURE host_mem_dealloc_c_4D, host_mem_dealloc_z_4D
END INTERFACE
INTERFACE dbcsr_cuda_dev_mem_zero
@ -106,6 +118,17 @@ MODULE dbcsr_cuda_memory
#if defined (__DBCSR_CUDA)
INTERFACE
FUNCTION dbcsr_cuda_dev_mem_info(free, avail) RESULT (istat) BIND(C, name="dc_dev_mem_info_cu")
USE ISO_C_BINDING
INTEGER(KIND=C_INT), INTENT(OUT) :: free, avail
INTEGER(KIND=C_INT) :: istat
END FUNCTION dbcsr_cuda_dev_mem_info
END INTERFACE
INTERFACE
FUNCTION cuda_dev_mem_alloc_cu(mem, n) RESULT (istat) BIND(C, name="dc_dev_mem_alloc")
USE ISO_C_BINDING
@ -177,6 +200,17 @@ MODULE dbcsr_cuda_memory
CONTAINS
#if ! defined (__DBCSR_CUDA)
FUNCTION dbcsr_cuda_dev_mem_info(free, avail) RESULT (istat)
INTEGER, INTENT(OUT) :: free, avail
INTEGER :: istat
free = 0
avail = 0
istat = -1
END FUNCTION dbcsr_cuda_dev_mem_info
#endif
#include "dbcsr_cuda_memory_i.F"
#include "dbcsr_cuda_memory_l.F"
@ -189,6 +223,35 @@ CONTAINS
!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
! Encapsulated memory routines for the device
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_raw (host_mem_c_ptr, n_bytes, error)
TYPE(C_PTR), INTENT(OUT) :: host_mem_c_ptr
INTEGER, INTENT(IN) :: n_bytes
TYPE(dbcsr_error_type), INTENT(INOUT), &
OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_raw', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr, &
INT(n_bytes, KIND=C_SIZE_T), 0_c_int, 0_c_int)
IF (istat /= 0) STOP "host_mem_alloc_raw: Could not allocate host pinned memory"
#else
STOP "host_mem_alloc_raw: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_alloc_raw
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dbcsr_cuda_dev_mem_init (dev_mem, error)
TYPE(dbcsr_cuda_mem_type), POINTER :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
@ -196,14 +259,16 @@ CONTAINS
NULLIFY (dev_mem)
END SUBROUTINE dbcsr_cuda_dev_mem_init
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dbcsr_cuda_dev_mem_new (dev_mem, data_type, error)
TYPE(dbcsr_cuda_mem_type), POINTER :: dev_mem
INTEGER, INTENT(IN), OPTIONAL :: data_type
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
INTEGER :: stat
ALLOCATE (dev_mem, stat=stat)
ALLOCATE (dev_mem)
dev_mem%data_type = 0
dev_mem%refcount = 1
IF (PRESENT (data_type)) THEN
@ -211,6 +276,10 @@ CONTAINS
ENDIF
END SUBROUTINE dbcsr_cuda_dev_mem_new
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dbcsr_cuda_dev_mem_hold (dev_mem, error)
TYPE(dbcsr_cuda_mem_type), POINTER :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
@ -219,21 +288,25 @@ CONTAINS
END SUBROUTINE dbcsr_cuda_dev_mem_hold
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dbcsr_cuda_dev_mem_release (dev_mem, error)
TYPE(dbcsr_cuda_mem_type), POINTER :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
INTEGER :: stat
dev_mem%refcount = dev_mem%refcount - 1
IF (dev_mem%refcount == 0) THEN
CALL dev_mem_dealloc_any (dev_mem, error=error)
DEALLOCATE (dev_mem, stat=stat)
DEALLOCATE (dev_mem)
NULLIFY (dev_mem)
ENDIF
END SUBROUTINE dbcsr_cuda_dev_mem_release
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_any (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type), INTENT(INOUT) :: dev_mem
INTEGER, INTENT(IN) :: n
@ -245,13 +318,10 @@ CONTAINS
routineP = moduleN//':'//routineN
INTEGER :: error_handle
TYPE(dbcsr_error_type) :: my_error
! ---------------------------------------------------------------------------
IF (PRESENT (error)) my_error = error
CALL dbcsr_error_set (routineN, error_handle, my_error)
IF (PRESENT (error)) error = my_error
IF (PRESENT (error)) CALL dbcsr_error_set(routineN, error_handle, error)
SELECT CASE (dbcsr_cuda_dev_mem_get_type (dev_mem))
CASE (dbcsr_type_int_4)
@ -273,19 +343,16 @@ CONTAINS
CALL dbcsr_cuda_dev_mem_alloc (dev_mem%d_z, n,&
error=error, stat=stat)
CASE default
IF (PRESENT (error)) my_error = error
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Invalid data type", __LINE__, error=my_error)
IF (PRESENT (error)) error = my_error
STOP "dev_mem_alloc_any: Invalid data type"
END SELECT
IF (PRESENT (error)) my_error = error
CALL dbcsr_error_stop (error_handle, my_error)
IF (PRESENT (error)) error = my_error
IF (PRESENT (error)) CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE dev_mem_alloc_any
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_any (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type), INTENT(INOUT) :: dev_mem
INTEGER, INTENT(IN) :: n
@ -297,13 +364,10 @@ CONTAINS
routineP = moduleN//':'//routineN
INTEGER :: error_handle
TYPE(dbcsr_error_type) :: my_error
! ---------------------------------------------------------------------------
IF (PRESENT (error)) my_error = error
CALL dbcsr_error_set (routineN, error_handle, my_error)
IF (PRESENT (error)) error = my_error
IF (PRESENT (error)) CALL dbcsr_error_set(routineN, error_handle, error)
SELECT CASE (dbcsr_cuda_dev_mem_get_type (dev_mem))
CASE (dbcsr_type_int_4)
@ -325,19 +389,16 @@ CONTAINS
CALL dbcsr_cuda_dev_mem_realloc (dev_mem%d_z, n,&
error=error, stat=stat)
CASE default
IF (PRESENT (error)) my_error = error
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Invalid data type", __LINE__, error=my_error)
IF (PRESENT (error)) error = my_error
STOP "dev_mem_realloc_any: Invalid data type"
END SELECT
IF (PRESENT (error)) my_error = error
CALL dbcsr_error_stop (error_handle, my_error)
IF (PRESENT (error)) error = my_error
IF (PRESENT (error)) CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE dev_mem_realloc_any
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_any (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), &
@ -348,13 +409,10 @@ CONTAINS
routineP = moduleN//':'//routineN
INTEGER :: error_handle
TYPE(dbcsr_error_type) :: my_error
! ---------------------------------------------------------------------------
IF (PRESENT (error)) my_error = error
CALL dbcsr_error_set (routineN, error_handle, my_error)
IF (PRESENT (error)) error = my_error
IF (PRESENT (error)) CALL dbcsr_error_set(routineN, error_handle, error)
SELECT CASE (dbcsr_cuda_dev_mem_get_type (dev_mem))
CASE (dbcsr_type_int_4)
@ -376,19 +434,16 @@ CONTAINS
CALL dbcsr_cuda_dev_mem_dealloc (dev_mem%d_z,&
error=error, stat=stat)
CASE default
IF (PRESENT (error)) my_error = error
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Invalid data type", __LINE__, error=my_error)
IF (PRESENT (error)) error = my_error
STOP "dev_mem_dealloc_any: Invalid data type"
END SELECT
IF (PRESENT (error)) my_error = error
CALL dbcsr_error_stop (error_handle, my_error)
IF (PRESENT (error)) error = my_error
IF (PRESENT (error)) CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE dev_mem_dealloc_any
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_any (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type), INTENT(INOUT) :: dev_mem
INTEGER, INTENT(IN) :: first, last

View file

@ -3,6 +3,9 @@
! Copyright (C) 2000 - 2013 Urban Borstnik and the CP2K developers group !
!-----------------------------------------------------------------------------!
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_[nametype1] (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type_[shorttype1]), &
INTENT(OUT) :: dev_mem
@ -28,18 +31,16 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_alloc_[nametype1]: Could not allocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_alloc_[nametype1]
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_[nametype1] (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type_[shorttype1]), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
@ -63,19 +64,16 @@
ELSE
istat = 0
ENDIF
IF (istat /= 0) THEN
IF (PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not deallocate GPU device memory",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) &
STOP "dev_mem_dealloc_[nametype1]: Could not deallocate GPU device memory"
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_dealloc_[nametype1]
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_[nametype1] (dev_mem, n, error, stat, memory_crunch)
TYPE(dbcsr_cuda_mem_type_[shorttype1]), &
INTENT(INOUT) :: dev_mem
@ -112,108 +110,134 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not reallocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_realloc_[nametype1]: Could not reallocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_realloc_[nametype1]
SUBROUTINE host_mem_alloc_[nametype1] (host_mem, n, write_combined, all_contexts, error, stat)
[type1], DIMENSION(:), &
POINTER :: host_mem
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_[nametype1] (host_mem, n, error)
[type1], DIMENSION(:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n
LOGICAL, INTENT(IN), OPTIONAL :: write_combined, all_contexts
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_[nametype1]', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=c_int) :: all_flag, wc_flag
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
IF (PRESENT (write_combined)) THEN
wc_flag = 1
ELSE
wc_flag = 0
ENDIF
IF (PRESENT (all_contexts)) THEN
all_flag = 1
ELSE
all_flag = 0
ENDIF
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr,&
INT(MAX(1,n)*[bytes1], KIND=C_SIZE_T),&
wc_flag, all_flag)
#else
istat = -1
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate host pinned memory",&
__LINE__, error=error)
NULLIFY (host_mem)
ELSE
#if defined (__DBCSR_CUDA)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
#endif
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
CALL host_mem_alloc_raw(host_mem_c_ptr, MAX(1,n)*[bytes1], error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
END SUBROUTINE host_mem_alloc_[nametype1]
SUBROUTINE host_mem_dealloc_[nametype1] (host_mem, error, stat)
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_[nametype1]_2D (host_mem, n1, n2, error)
[type1], DIMENSION(:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*[bytes1]
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n1),MAX(1,n2) /))
END SUBROUTINE host_mem_alloc_[nametype1]_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_[nametype1]_4D (host_mem, n1, n2, n3, n4, error)
[type1], DIMENSION(:,:,:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2, n3, n4
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*MAX(1,n3)*MAX(1,n4)*[bytes1]
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, &
(/ MAX(1,n1),MAX(1,n2),MAX(1,n3),MAX(1,n4) /))
END SUBROUTINE host_mem_alloc_[nametype1]_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_[nametype1] (host_mem, error)
[type1], DIMENSION(:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_[nametype1]', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
[type1], POINTER :: host_first_el
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
IF (SIZE (host_mem) .GT. 0) THEN
host_first_el => host_mem(1)
host_mem_c_ptr = C_LOC (host_first_el)
ENDIF
IF (SIZE (host_mem) .GT. 0) THEN
istat = cuda_host_mem_dealloc_cu (host_mem_c_ptr)
ENDIF
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_[nametype1]: Error deallocating host pinned memory"
#else
istat = -1
STOP "host_mem_dealloc_[nametype1]: DBCSR_CUDA not compiled in."
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Error deallocating host pinned memory",&
__LINE__, error=error)
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_[nametype1]
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_[nametype1]_2D (host_mem, error)
[type1], DIMENSION(:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_[nametype1]_2D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_[nametype1]_2D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_[nametype1]: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_[nametype1]_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_[nametype1]_4D (host_mem, error)
[type1], DIMENSION(:,:,:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_[nametype1]_4D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1,1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_[nametype1]_4D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_[nametype1]: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_[nametype1]_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_[nametype1] (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type_[shorttype1]), &
INTENT(INOUT) :: dev_mem

View file

@ -3,6 +3,9 @@
! Copyright (C) 2000 - 2013 Urban Borstnik and the CP2K developers group !
!-----------------------------------------------------------------------------!
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_c (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type_c4), &
INTENT(OUT) :: dev_mem
@ -28,18 +31,16 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_alloc_c: Could not allocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_alloc_c
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_c (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type_c4), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
@ -63,19 +64,16 @@
ELSE
istat = 0
ENDIF
IF (istat /= 0) THEN
IF (PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not deallocate GPU device memory",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) &
STOP "dev_mem_dealloc_c: Could not deallocate GPU device memory"
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_dealloc_c
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_c (dev_mem, n, error, stat, memory_crunch)
TYPE(dbcsr_cuda_mem_type_c4), &
INTENT(INOUT) :: dev_mem
@ -112,108 +110,134 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not reallocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_realloc_c: Could not reallocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_realloc_c
SUBROUTINE host_mem_alloc_c (host_mem, n, write_combined, all_contexts, error, stat)
COMPLEX(kind=real_4), DIMENSION(:), &
POINTER :: host_mem
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_c (host_mem, n, error)
COMPLEX(kind=real_4), DIMENSION(:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n
LOGICAL, INTENT(IN), OPTIONAL :: write_combined, all_contexts
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_c', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=c_int) :: all_flag, wc_flag
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
IF (PRESENT (write_combined)) THEN
wc_flag = 1
ELSE
wc_flag = 0
ENDIF
IF (PRESENT (all_contexts)) THEN
all_flag = 1
ELSE
all_flag = 0
ENDIF
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr,&
INT(MAX(1,n)*(2*real_4_size), KIND=C_SIZE_T),&
wc_flag, all_flag)
#else
istat = -1
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate host pinned memory",&
__LINE__, error=error)
NULLIFY (host_mem)
ELSE
#if defined (__DBCSR_CUDA)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
#endif
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
CALL host_mem_alloc_raw(host_mem_c_ptr, MAX(1,n)*(2*real_4_size), error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
END SUBROUTINE host_mem_alloc_c
SUBROUTINE host_mem_dealloc_c (host_mem, error, stat)
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_c_2D (host_mem, n1, n2, error)
COMPLEX(kind=real_4), DIMENSION(:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*(2*real_4_size)
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n1),MAX(1,n2) /))
END SUBROUTINE host_mem_alloc_c_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_c_4D (host_mem, n1, n2, n3, n4, error)
COMPLEX(kind=real_4), DIMENSION(:,:,:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2, n3, n4
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*MAX(1,n3)*MAX(1,n4)*(2*real_4_size)
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, &
(/ MAX(1,n1),MAX(1,n2),MAX(1,n3),MAX(1,n4) /))
END SUBROUTINE host_mem_alloc_c_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_c (host_mem, error)
COMPLEX(kind=real_4), DIMENSION(:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_c', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
COMPLEX(kind=real_4), POINTER :: host_first_el
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
IF (SIZE (host_mem) .GT. 0) THEN
host_first_el => host_mem(1)
host_mem_c_ptr = C_LOC (host_first_el)
ENDIF
IF (SIZE (host_mem) .GT. 0) THEN
istat = cuda_host_mem_dealloc_cu (host_mem_c_ptr)
ENDIF
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_c: Error deallocating host pinned memory"
#else
istat = -1
STOP "host_mem_dealloc_c: DBCSR_CUDA not compiled in."
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Error deallocating host pinned memory",&
__LINE__, error=error)
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_c
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_c_2D (host_mem, error)
COMPLEX(kind=real_4), DIMENSION(:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_c_2D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_c_2D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_c: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_c_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_c_4D (host_mem, error)
COMPLEX(kind=real_4), DIMENSION(:,:,:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_c_4D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1,1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_c_4D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_c: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_c_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_c (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type_c4), &
INTENT(INOUT) :: dev_mem

View file

@ -3,6 +3,9 @@
! Copyright (C) 2000 - 2013 Urban Borstnik and the CP2K developers group !
!-----------------------------------------------------------------------------!
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_d (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type_r8), &
INTENT(OUT) :: dev_mem
@ -28,18 +31,16 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_alloc_d: Could not allocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_alloc_d
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_d (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type_r8), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
@ -63,19 +64,16 @@
ELSE
istat = 0
ENDIF
IF (istat /= 0) THEN
IF (PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not deallocate GPU device memory",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) &
STOP "dev_mem_dealloc_d: Could not deallocate GPU device memory"
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_dealloc_d
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_d (dev_mem, n, error, stat, memory_crunch)
TYPE(dbcsr_cuda_mem_type_r8), &
INTENT(INOUT) :: dev_mem
@ -112,108 +110,134 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not reallocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_realloc_d: Could not reallocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_realloc_d
SUBROUTINE host_mem_alloc_d (host_mem, n, write_combined, all_contexts, error, stat)
REAL(kind=real_8), DIMENSION(:), &
POINTER :: host_mem
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_d (host_mem, n, error)
REAL(kind=real_8), DIMENSION(:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n
LOGICAL, INTENT(IN), OPTIONAL :: write_combined, all_contexts
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_d', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=c_int) :: all_flag, wc_flag
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
IF (PRESENT (write_combined)) THEN
wc_flag = 1
ELSE
wc_flag = 0
ENDIF
IF (PRESENT (all_contexts)) THEN
all_flag = 1
ELSE
all_flag = 0
ENDIF
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr,&
INT(MAX(1,n)*real_8_size, KIND=C_SIZE_T),&
wc_flag, all_flag)
#else
istat = -1
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate host pinned memory",&
__LINE__, error=error)
NULLIFY (host_mem)
ELSE
#if defined (__DBCSR_CUDA)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
#endif
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
CALL host_mem_alloc_raw(host_mem_c_ptr, MAX(1,n)*real_8_size, error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
END SUBROUTINE host_mem_alloc_d
SUBROUTINE host_mem_dealloc_d (host_mem, error, stat)
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_d_2D (host_mem, n1, n2, error)
REAL(kind=real_8), DIMENSION(:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*real_8_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n1),MAX(1,n2) /))
END SUBROUTINE host_mem_alloc_d_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_d_4D (host_mem, n1, n2, n3, n4, error)
REAL(kind=real_8), DIMENSION(:,:,:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2, n3, n4
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*MAX(1,n3)*MAX(1,n4)*real_8_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, &
(/ MAX(1,n1),MAX(1,n2),MAX(1,n3),MAX(1,n4) /))
END SUBROUTINE host_mem_alloc_d_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_d (host_mem, error)
REAL(kind=real_8), DIMENSION(:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_d', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
REAL(kind=real_8), POINTER :: host_first_el
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
IF (SIZE (host_mem) .GT. 0) THEN
host_first_el => host_mem(1)
host_mem_c_ptr = C_LOC (host_first_el)
ENDIF
IF (SIZE (host_mem) .GT. 0) THEN
istat = cuda_host_mem_dealloc_cu (host_mem_c_ptr)
ENDIF
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_d: Error deallocating host pinned memory"
#else
istat = -1
STOP "host_mem_dealloc_d: DBCSR_CUDA not compiled in."
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Error deallocating host pinned memory",&
__LINE__, error=error)
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_d
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_d_2D (host_mem, error)
REAL(kind=real_8), DIMENSION(:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_d_2D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_d_2D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_d: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_d_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_d_4D (host_mem, error)
REAL(kind=real_8), DIMENSION(:,:,:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_d_4D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1,1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_d_4D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_d: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_d_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_d (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type_r8), &
INTENT(INOUT) :: dev_mem

View file

@ -3,6 +3,9 @@
! Copyright (C) 2000 - 2013 Urban Borstnik and the CP2K developers group !
!-----------------------------------------------------------------------------!
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_i (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type_i4), &
INTENT(OUT) :: dev_mem
@ -28,18 +31,16 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_alloc_i: Could not allocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_alloc_i
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_i (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type_i4), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
@ -63,19 +64,16 @@
ELSE
istat = 0
ENDIF
IF (istat /= 0) THEN
IF (PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not deallocate GPU device memory",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) &
STOP "dev_mem_dealloc_i: Could not deallocate GPU device memory"
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_dealloc_i
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_i (dev_mem, n, error, stat, memory_crunch)
TYPE(dbcsr_cuda_mem_type_i4), &
INTENT(INOUT) :: dev_mem
@ -112,108 +110,134 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not reallocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_realloc_i: Could not reallocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_realloc_i
SUBROUTINE host_mem_alloc_i (host_mem, n, write_combined, all_contexts, error, stat)
INTEGER(KIND=int_4), DIMENSION(:), &
POINTER :: host_mem
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_i (host_mem, n, error)
INTEGER(KIND=int_4), DIMENSION(:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n
LOGICAL, INTENT(IN), OPTIONAL :: write_combined, all_contexts
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_i', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=c_int) :: all_flag, wc_flag
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
IF (PRESENT (write_combined)) THEN
wc_flag = 1
ELSE
wc_flag = 0
ENDIF
IF (PRESENT (all_contexts)) THEN
all_flag = 1
ELSE
all_flag = 0
ENDIF
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr,&
INT(MAX(1,n)*int_4_size, KIND=C_SIZE_T),&
wc_flag, all_flag)
#else
istat = -1
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate host pinned memory",&
__LINE__, error=error)
NULLIFY (host_mem)
ELSE
#if defined (__DBCSR_CUDA)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
#endif
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
CALL host_mem_alloc_raw(host_mem_c_ptr, MAX(1,n)*int_4_size, error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
END SUBROUTINE host_mem_alloc_i
SUBROUTINE host_mem_dealloc_i (host_mem, error, stat)
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_i_2D (host_mem, n1, n2, error)
INTEGER(KIND=int_4), DIMENSION(:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*int_4_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n1),MAX(1,n2) /))
END SUBROUTINE host_mem_alloc_i_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_i_4D (host_mem, n1, n2, n3, n4, error)
INTEGER(KIND=int_4), DIMENSION(:,:,:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2, n3, n4
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*MAX(1,n3)*MAX(1,n4)*int_4_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, &
(/ MAX(1,n1),MAX(1,n2),MAX(1,n3),MAX(1,n4) /))
END SUBROUTINE host_mem_alloc_i_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_i (host_mem, error)
INTEGER(KIND=int_4), DIMENSION(:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_i', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=int_4), POINTER :: host_first_el
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
IF (SIZE (host_mem) .GT. 0) THEN
host_first_el => host_mem(1)
host_mem_c_ptr = C_LOC (host_first_el)
ENDIF
IF (SIZE (host_mem) .GT. 0) THEN
istat = cuda_host_mem_dealloc_cu (host_mem_c_ptr)
ENDIF
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_i: Error deallocating host pinned memory"
#else
istat = -1
STOP "host_mem_dealloc_i: DBCSR_CUDA not compiled in."
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Error deallocating host pinned memory",&
__LINE__, error=error)
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_i
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_i_2D (host_mem, error)
INTEGER(KIND=int_4), DIMENSION(:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_i_2D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_i_2D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_i: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_i_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_i_4D (host_mem, error)
INTEGER(KIND=int_4), DIMENSION(:,:,:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_i_4D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1,1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_i_4D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_i: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_i_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_i (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type_i4), &
INTENT(INOUT) :: dev_mem

View file

@ -3,6 +3,9 @@
! Copyright (C) 2000 - 2013 Urban Borstnik and the CP2K developers group !
!-----------------------------------------------------------------------------!
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_l (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type_i8), &
INTENT(OUT) :: dev_mem
@ -28,18 +31,16 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_alloc_l: Could not allocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_alloc_l
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_l (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type_i8), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
@ -63,19 +64,16 @@
ELSE
istat = 0
ENDIF
IF (istat /= 0) THEN
IF (PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not deallocate GPU device memory",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) &
STOP "dev_mem_dealloc_l: Could not deallocate GPU device memory"
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_dealloc_l
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_l (dev_mem, n, error, stat, memory_crunch)
TYPE(dbcsr_cuda_mem_type_i8), &
INTENT(INOUT) :: dev_mem
@ -112,108 +110,134 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not reallocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_realloc_l: Could not reallocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_realloc_l
SUBROUTINE host_mem_alloc_l (host_mem, n, write_combined, all_contexts, error, stat)
INTEGER(KIND=int_8), DIMENSION(:), &
POINTER :: host_mem
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_l (host_mem, n, error)
INTEGER(KIND=int_8), DIMENSION(:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n
LOGICAL, INTENT(IN), OPTIONAL :: write_combined, all_contexts
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_l', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=c_int) :: all_flag, wc_flag
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
IF (PRESENT (write_combined)) THEN
wc_flag = 1
ELSE
wc_flag = 0
ENDIF
IF (PRESENT (all_contexts)) THEN
all_flag = 1
ELSE
all_flag = 0
ENDIF
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr,&
INT(MAX(1,n)*int_8_size, KIND=C_SIZE_T),&
wc_flag, all_flag)
#else
istat = -1
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate host pinned memory",&
__LINE__, error=error)
NULLIFY (host_mem)
ELSE
#if defined (__DBCSR_CUDA)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
#endif
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
CALL host_mem_alloc_raw(host_mem_c_ptr, MAX(1,n)*int_8_size, error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
END SUBROUTINE host_mem_alloc_l
SUBROUTINE host_mem_dealloc_l (host_mem, error, stat)
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_l_2D (host_mem, n1, n2, error)
INTEGER(KIND=int_8), DIMENSION(:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*int_8_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n1),MAX(1,n2) /))
END SUBROUTINE host_mem_alloc_l_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_l_4D (host_mem, n1, n2, n3, n4, error)
INTEGER(KIND=int_8), DIMENSION(:,:,:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2, n3, n4
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*MAX(1,n3)*MAX(1,n4)*int_8_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, &
(/ MAX(1,n1),MAX(1,n2),MAX(1,n3),MAX(1,n4) /))
END SUBROUTINE host_mem_alloc_l_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_l (host_mem, error)
INTEGER(KIND=int_8), DIMENSION(:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_l', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=int_8), POINTER :: host_first_el
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
IF (SIZE (host_mem) .GT. 0) THEN
host_first_el => host_mem(1)
host_mem_c_ptr = C_LOC (host_first_el)
ENDIF
IF (SIZE (host_mem) .GT. 0) THEN
istat = cuda_host_mem_dealloc_cu (host_mem_c_ptr)
ENDIF
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_l: Error deallocating host pinned memory"
#else
istat = -1
STOP "host_mem_dealloc_l: DBCSR_CUDA not compiled in."
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Error deallocating host pinned memory",&
__LINE__, error=error)
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_l
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_l_2D (host_mem, error)
INTEGER(KIND=int_8), DIMENSION(:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_l_2D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_l_2D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_l: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_l_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_l_4D (host_mem, error)
INTEGER(KIND=int_8), DIMENSION(:,:,:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_l_4D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1,1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_l_4D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_l: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_l_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_l (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type_i8), &
INTENT(INOUT) :: dev_mem

View file

@ -3,6 +3,9 @@
! Copyright (C) 2000 - 2013 Urban Borstnik and the CP2K developers group !
!-----------------------------------------------------------------------------!
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_r (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type_r4), &
INTENT(OUT) :: dev_mem
@ -28,18 +31,16 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_alloc_r: Could not allocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_alloc_r
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_r (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type_r4), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
@ -63,19 +64,16 @@
ELSE
istat = 0
ENDIF
IF (istat /= 0) THEN
IF (PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not deallocate GPU device memory",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) &
STOP "dev_mem_dealloc_r: Could not deallocate GPU device memory"
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_dealloc_r
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_r (dev_mem, n, error, stat, memory_crunch)
TYPE(dbcsr_cuda_mem_type_r4), &
INTENT(INOUT) :: dev_mem
@ -112,108 +110,134 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not reallocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_realloc_r: Could not reallocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_realloc_r
SUBROUTINE host_mem_alloc_r (host_mem, n, write_combined, all_contexts, error, stat)
REAL(kind=real_4), DIMENSION(:), &
POINTER :: host_mem
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_r (host_mem, n, error)
REAL(kind=real_4), DIMENSION(:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n
LOGICAL, INTENT(IN), OPTIONAL :: write_combined, all_contexts
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_r', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=c_int) :: all_flag, wc_flag
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
IF (PRESENT (write_combined)) THEN
wc_flag = 1
ELSE
wc_flag = 0
ENDIF
IF (PRESENT (all_contexts)) THEN
all_flag = 1
ELSE
all_flag = 0
ENDIF
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr,&
INT(MAX(1,n)*real_4_size, KIND=C_SIZE_T),&
wc_flag, all_flag)
#else
istat = -1
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate host pinned memory",&
__LINE__, error=error)
NULLIFY (host_mem)
ELSE
#if defined (__DBCSR_CUDA)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
#endif
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
CALL host_mem_alloc_raw(host_mem_c_ptr, MAX(1,n)*real_4_size, error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
END SUBROUTINE host_mem_alloc_r
SUBROUTINE host_mem_dealloc_r (host_mem, error, stat)
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_r_2D (host_mem, n1, n2, error)
REAL(kind=real_4), DIMENSION(:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*real_4_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n1),MAX(1,n2) /))
END SUBROUTINE host_mem_alloc_r_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_r_4D (host_mem, n1, n2, n3, n4, error)
REAL(kind=real_4), DIMENSION(:,:,:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2, n3, n4
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*MAX(1,n3)*MAX(1,n4)*real_4_size
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, &
(/ MAX(1,n1),MAX(1,n2),MAX(1,n3),MAX(1,n4) /))
END SUBROUTINE host_mem_alloc_r_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_r (host_mem, error)
REAL(kind=real_4), DIMENSION(:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_r', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
REAL(kind=real_4), POINTER :: host_first_el
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
IF (SIZE (host_mem) .GT. 0) THEN
host_first_el => host_mem(1)
host_mem_c_ptr = C_LOC (host_first_el)
ENDIF
IF (SIZE (host_mem) .GT. 0) THEN
istat = cuda_host_mem_dealloc_cu (host_mem_c_ptr)
ENDIF
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_r: Error deallocating host pinned memory"
#else
istat = -1
STOP "host_mem_dealloc_r: DBCSR_CUDA not compiled in."
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Error deallocating host pinned memory",&
__LINE__, error=error)
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_r
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_r_2D (host_mem, error)
REAL(kind=real_4), DIMENSION(:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_r_2D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_r_2D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_r: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_r_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_r_4D (host_mem, error)
REAL(kind=real_4), DIMENSION(:,:,:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_r_4D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1,1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_r_4D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_r: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_r_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_r (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type_r4), &
INTENT(INOUT) :: dev_mem

View file

@ -3,6 +3,9 @@
! Copyright (C) 2000 - 2013 Urban Borstnik and the CP2K developers group !
!-----------------------------------------------------------------------------!
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_alloc_z (dev_mem, n, error, stat)
TYPE(dbcsr_cuda_mem_type_c8), &
INTENT(OUT) :: dev_mem
@ -28,18 +31,16 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_alloc_z: Could not allocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_alloc_z
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_dealloc_z (dev_mem, error, stat)
TYPE(dbcsr_cuda_mem_type_c8), INTENT(INOUT) :: dev_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
@ -63,19 +64,16 @@
ELSE
istat = 0
ENDIF
IF (istat /= 0) THEN
IF (PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not deallocate GPU device memory",&
__LINE__, error=error)
ENDIF
ENDIF
IF (istat /= 0) &
STOP "dev_mem_dealloc_z: Could not deallocate GPU device memory"
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_dealloc_z
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_realloc_z (dev_mem, n, error, stat, memory_crunch)
TYPE(dbcsr_cuda_mem_type_c8), &
INTENT(INOUT) :: dev_mem
@ -112,108 +110,134 @@
#endif
IF (istat /= 0) THEN
dev_mem%len = -1
IF (PRESENT(error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not reallocate GPU device memory",&
__LINE__, error=error)
ENDIF
STOP "dev_mem_realloc_z: Could not reallocate GPU device memory"
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dev_mem_realloc_z
SUBROUTINE host_mem_alloc_z (host_mem, n, write_combined, all_contexts, error, stat)
COMPLEX(kind=real_8), DIMENSION(:), &
POINTER :: host_mem
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_z (host_mem, n, error)
COMPLEX(kind=real_8), DIMENSION(:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n
LOGICAL, INTENT(IN), OPTIONAL :: write_combined, all_contexts
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_alloc_z', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
INTEGER(KIND=c_int) :: all_flag, wc_flag
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
IF (PRESENT (write_combined)) THEN
wc_flag = 1
ELSE
wc_flag = 0
ENDIF
IF (PRESENT (all_contexts)) THEN
all_flag = 1
ELSE
all_flag = 0
ENDIF
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_alloc_cu(host_mem_c_ptr,&
INT(MAX(1,n)*(2*real_8_size), KIND=C_SIZE_T),&
wc_flag, all_flag)
#else
istat = -1
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Could not allocate host pinned memory",&
__LINE__, error=error)
NULLIFY (host_mem)
ELSE
#if defined (__DBCSR_CUDA)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
#endif
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
CALL host_mem_alloc_raw(host_mem_c_ptr, MAX(1,n)*(2*real_8_size), error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n) /))
END SUBROUTINE host_mem_alloc_z
SUBROUTINE host_mem_dealloc_z (host_mem, error, stat)
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_z_2D (host_mem, n1, n2, error)
COMPLEX(kind=real_8), DIMENSION(:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*(2*real_8_size)
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, (/ MAX(1,n1),MAX(1,n2) /))
END SUBROUTINE host_mem_alloc_z_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_alloc_z_4D (host_mem, n1, n2, n3, n4, error)
COMPLEX(kind=real_8), DIMENSION(:,:,:,:), POINTER :: host_mem
INTEGER, INTENT(IN) :: n1, n2, n3, n4
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
TYPE(C_PTR) :: host_mem_c_ptr
INTEGER :: n_bytes
n_bytes = MAX(1,n1)*MAX(1,n2)*MAX(1,n3)*MAX(1,n4)*(2*real_8_size)
CALL host_mem_alloc_raw(host_mem_c_ptr,n_bytes , error)
CALL C_F_POINTER (host_mem_c_ptr, host_mem, &
(/ MAX(1,n1),MAX(1,n2),MAX(1,n3),MAX(1,n4) /))
END SUBROUTINE host_mem_alloc_z_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_z (host_mem, error)
COMPLEX(kind=real_8), DIMENSION(:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
INTEGER, INTENT(OUT), OPTIONAL :: stat
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_z', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
COMPLEX(kind=real_8), POINTER :: host_first_el
TYPE(C_PTR) :: host_mem_c_ptr
! ---------------------------------------------------------------------------
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
IF (SIZE (host_mem) .GT. 0) THEN
host_first_el => host_mem(1)
host_mem_c_ptr = C_LOC (host_first_el)
ENDIF
IF (SIZE (host_mem) .GT. 0) THEN
istat = cuda_host_mem_dealloc_cu (host_mem_c_ptr)
ENDIF
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_z: Error deallocating host pinned memory"
#else
istat = -1
STOP "host_mem_dealloc_z: DBCSR_CUDA not compiled in."
#endif
IF (istat /= 0 .AND. PRESENT (error)) THEN
CALL dbcsr_assert (istat, "EQ", 0,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Error deallocating host pinned memory",&
__LINE__, error=error)
ENDIF
IF (PRESENT (stat)) stat = istat
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_z
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_z_2D (host_mem, error)
COMPLEX(kind=real_8), DIMENSION(:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_z_2D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_z_2D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_z: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_z_2D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE host_mem_dealloc_z_4D (host_mem, error)
COMPLEX(kind=real_8), DIMENSION(:,:,:,:), &
POINTER :: host_mem
TYPE(dbcsr_error_type), INTENT(INOUT), OPTIONAL :: error
CHARACTER(len=*), PARAMETER :: routineN = 'host_mem_dealloc_z_4D', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, istat
IF (SIZE (host_mem) == 0) RETURN
IF (PRESENT (error)) CALL dbcsr_error_set (routineN, error_handle, error)
#if defined (__DBCSR_CUDA)
istat = cuda_host_mem_dealloc_cu(C_LOC(host_mem(1,1,1,1)))
IF (istat /= 0 ) &
STOP "host_mem_dealloc_z_4D: Error deallocating host pinned memory"
#else
STOP "host_mem_dealloc_z: DBCSR_CUDA not compiled in."
#endif
IF (PRESENT (error)) CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE host_mem_dealloc_z_4D
! *****************************************************************************
!> \brief
! *****************************************************************************
SUBROUTINE dev_mem_zero_z (dev_mem, first, last, error)
TYPE(dbcsr_cuda_mem_type_c8), &
INTENT(INOUT) :: dev_mem

View file

@ -32,7 +32,10 @@ MODULE dbcsr_cuda_types
PUBLIC :: dbcsr_cuda_mem_type_i4, dbcsr_cuda_mem_type_i8,&
dbcsr_cuda_mem_type_r4, dbcsr_cuda_mem_type_r8,&
dbcsr_cuda_mem_type_c4, dbcsr_cuda_mem_type_c8
PUBLIC :: dbcsr_cuda_mem_p_type
! TODO: these datastructures are pointless. They hold C_PTR anyways.
TYPE dbcsr_cuda_mem_type_i4
TYPE(C_PTR) :: ref
INTEGER :: len, req_len
@ -74,7 +77,10 @@ MODULE dbcsr_cuda_types
TYPE(dbcsr_cuda_mem_type_c8) :: d_z
INTEGER :: refcount
END TYPE dbcsr_cuda_mem_type
TYPE dbcsr_cuda_mem_p_type
TYPE(dbcsr_cuda_mem_type), POINTER :: p
END TYPE dbcsr_cuda_mem_p_type
END MODULE dbcsr_cuda_types

View file

@ -292,17 +292,13 @@ CONTAINS
CASE (dbcsr_memory_CUDA_host_pinned)
SELECT CASE (data_type)
CASE (dbcsr_type_real_4)
CALL dbcsr_cuda_host_mem_alloc (area%r_sp, sizes(1), stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_alloc (area%r_sp, sizes(1), error=error)
CASE (dbcsr_type_real_8)
CALL dbcsr_cuda_host_mem_alloc (area%r_dp, sizes(1), stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_alloc (area%r_dp, sizes(1), error=error)
CASE (dbcsr_type_complex_4)
CALL dbcsr_cuda_host_mem_alloc (area%c_sp, sizes(1), stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_alloc (area%c_sp, sizes(1), error=error)
CASE (dbcsr_type_complex_8)
CALL dbcsr_cuda_host_mem_alloc (area%c_dp, sizes(1), stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_alloc (area%c_dp, sizes(1), error=error)
CASE (dbcsr_type_real_8_2d, dbcsr_type_real_4_2d,&
dbcsr_type_complex_8_2d, dbcsr_type_complex_4_2d)
CALL dbcsr_assert (.FALSE.,&
@ -399,20 +395,16 @@ CONTAINS
CASE (dbcsr_memory_CUDA_host_pinned)
SELECT CASE (data_type)
CASE (dbcsr_type_real_4)
CALL dbcsr_cuda_host_mem_dealloc (area%r_sp, stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_dealloc (area%r_sp, error=error)
NULLIFY (area%r_sp)
CASE (dbcsr_type_real_8)
CALL dbcsr_cuda_host_mem_dealloc (area%r_dp, stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_dealloc (area%r_dp, error=error)
NULLIFY (area%r_dp)
CASE (dbcsr_type_complex_4)
CALL dbcsr_cuda_host_mem_dealloc (area%c_sp, stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_dealloc (area%c_sp, error=error)
NULLIFY (area%c_sp)
CASE (dbcsr_type_complex_8)
CALL dbcsr_cuda_host_mem_dealloc (area%c_dp, stat=stat,&
error=error)
CALL dbcsr_cuda_host_mem_dealloc (area%c_dp, error=error)
NULLIFY (area%c_dp)
CASE (dbcsr_type_real_8_2d, dbcsr_type_real_4_2d,&
dbcsr_type_complex_8_2d, dbcsr_type_complex_4_2d)

View file

@ -63,10 +63,12 @@ MODULE dbcsr_mm_cannon
dbcsr_mp_numnodes, dbcsr_mp_pgrid, dbcsr_nblkcols_local, &
dbcsr_nblkcols_total, dbcsr_nblkrows_local, dbcsr_nblkrows_total, &
dbcsr_nfullcols_local, dbcsr_nfullrows_local, dbcsr_valid_index
USE dbcsr_mm_multrec, ONLY: dbcsr_multrec_finalize,&
dbcsr_multrec_init,&
dbcsr_multrec_sparse
USE dbcsr_mm_types, ONLY: carrier_type
USE dbcsr_mm_multrec, ONLY: dbcsr_mm_multrec_finalize,&
dbcsr_mm_multrec_init,&
dbcsr_mm_multrec_lib_finalize,&
dbcsr_mm_multrec_lib_init,&
dbcsr_mm_multrec_multiply
USE dbcsr_mm_types, ONLY: multrec_carrier_type
USE dbcsr_mp_operations, ONLY: dbcsr_irecv_any,&
dbcsr_isend_any,&
dbcsr_mp_type_from_anytype
@ -91,10 +93,33 @@ MODULE dbcsr_mm_cannon
LOGICAL, PARAMETER :: debug_mod = .FALSE.
LOGICAL, PARAMETER :: careful_mod = .FALSE.
PUBLIC :: dbcsr_mult_m_e_e
PUBLIC :: dbcsr_mm_cannon_lib_init, dbcsr_mm_cannon_lib_finalize
PUBLIC :: dbcsr_mm_cannon_multiply
CONTAINS
CONTAINS
! *****************************************************************************
! \brief Initialize the library
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_cannon_lib_init(error)
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CALL dbcsr_mm_multrec_lib_init(error)
END SUBROUTINE
! *****************************************************************************
! \brief Finalize the library
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_cannon_lib_finalize(error)
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CALL dbcsr_mm_multrec_lib_finalize(error)
END SUBROUTINE
! *****************************************************************************
@ -108,7 +133,7 @@ CONTAINS
!> existing product matrix; default is no
!> \param[out] flop (optional) effective flop
! *****************************************************************************
SUBROUTINE dbcsr_mult_m_e_e (left_set, right_set, product_matrix,&
SUBROUTINE dbcsr_mm_cannon_multiply(left_set, right_set, product_matrix,&
error, retain_sparsity, &
filter_eps, flop)
TYPE(dbcsr_2d_array_type), POINTER :: left_set, right_set
@ -118,7 +143,7 @@ CONTAINS
REAL(kind=real_8), INTENT(in), OPTIONAL :: filter_eps
INTEGER(KIND=int_8), INTENT(OUT) :: flop
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mult_m_e_e', &
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_cannon_multiply', &
routineP = moduleN//':'//routineN
CHARACTER(LEN=80), PARAMETER :: &
fdata = '(A,1X,I4,"(",2(I3),"x",2(I3),")","(",I3,"x",I3,")")', fxfer = &
@ -179,7 +204,6 @@ CONTAINS
right_fill
REAL(kind=sp), ALLOCATABLE, DIMENSION(:) :: left_norms, right_norms, &
row_max_epss
TYPE(carrier_type), SAVE :: carrier
TYPE(dbcsr_2d_array_type), POINTER :: left_buffer_1, left_buffer_2, &
left_buffer_calc, left_buffer_comm, right_buffer_1, right_buffer_2, &
right_buffer_calc, right_buffer_comm
@ -199,8 +223,9 @@ CONTAINS
left_send_subtypes, &
right_recv_subtypes, &
right_send_subtypes
TYPE(multrec_carrier_type), SAVE :: multrec_carrier
!$OMP THREADPRIVATE (carrier)
!$OMP THREADPRIVATE (multrec_carrier)
! ---------------------------------------------------------------------------
@ -659,13 +684,11 @@ CONTAINS
!$omp right_data_sr, right_data_rr, left_data_sr, left_data_rr,&
!$omp right_index_sr, right_index_rr, left_index_sr, left_index_rr), &
!$omp shared (a_dev, b_dev) firstprivate (error)
CALL dbcsr_multrec_init(&
CALL dbcsr_mm_multrec_init(&
left_buffer_comm%mats(1, 1)%m,&
right_buffer_comm%mats(1, 1)%m,&
product_matrix%m,&
carrier,&
right_data_sr, right_data_rr, left_data_sr, left_data_rr,&
right_index_sr, right_index_rr, left_index_sr, left_index_rr, &
multrec_carrier,&
keep_sparsity=keep_sparsity,&
eps=filter_eps,&
row_max_epss = row_max_epss,&
@ -1042,16 +1065,15 @@ CONTAINS
!$omp reduction (+: flop_single)
t_error=error
CALL dbcsr_error_set(routineN//"_multrec_sparse", error_handler2, t_error)
CALL dbcsr_multrec_sparse(&
CALL dbcsr_mm_multrec_multiply(&
left_buffer_calc%mats(1, v_ki+1)%m,&
right_buffer_calc%mats(v_ki+1, 1)%m,&
product_matrix%m,&
carrier,&
keep_sparsity=keep_sparsity,&
multrec_carrier,&
keep_sparsity=keep_sparsity,&
flop=flop_single,&
a_norms=left_norms, b_norms=right_norms,&
eps=filter_eps,&
row_max_epss = row_max_epss, error=t_error)
error=t_error)
! once multiplication is complete, thread zero polls MPI until
! all other threads have also completed
!$omp atomic
@ -1062,7 +1084,14 @@ CONTAINS
!$ nthreads = omp_get_num_threads()
IF (use_comm_thread .AND. (ithread .EQ. 0)) THEN
DO WHILE (threads_finished .NE. nthreads)
CALL progress_comms(carrier)
CALL mp_testany(right_data_sr)
CALL mp_testany(right_data_rr)
CALL mp_testany(left_data_sr)
CALL mp_testany(left_data_rr)
CALL mp_testany(right_index_sr)
CALL mp_testany(right_index_rr)
CALL mp_testany(left_index_sr)
CALL mp_testany(left_index_rr)
!$omp flush (threads_finished)
END DO
END IF
@ -1097,14 +1126,14 @@ CONTAINS
!
!$omp parallel &
!$omp parallel &
!$omp default (none) &
!$omp private (t_error) &
!$omp shared (product_matrix, error)
!
t_error = error
! Finalize multiplication
CALL dbcsr_multrec_finalize (carrier, product_matrix%m, t_error)
CALL dbcsr_mm_multrec_finalize(multrec_carrier, product_matrix%m, t_error)
!$omp end parallel
!
IF (ALLOCATED (right_norms)) THEN
@ -1181,7 +1210,7 @@ CONTAINS
i = real_8_size * 2
END SELECT
CALL dbcsr_error_stop(error_handler, error)
END SUBROUTINE dbcsr_mult_m_e_e
END SUBROUTINE dbcsr_mm_cannon_multiply
! ******************************************************************************
@ -1726,24 +1755,6 @@ CONTAINS
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE calculate_norms
! *****************************************************************************
!> \brief Call in MPI to progrss any outstanding communications
! *****************************************************************************
SUBROUTINE progress_comms(carrier)
TYPE(carrier_type), INTENT(inout) :: carrier
CALL mp_testany(carrier%right_data_sr)
CALL mp_testany(carrier%right_data_rr)
CALL mp_testany(carrier%left_data_sr)
CALL mp_testany(carrier%left_data_rr)
CALL mp_testany(carrier%right_index_sr)
CALL mp_testany(carrier%right_index_rr)
CALL mp_testany(carrier%left_index_sr)
CALL mp_testany(carrier%left_index_rr)
END SUBROUTINE
#include "dbcsr_mm_cannon_d.F"
#include "dbcsr_mm_cannon_z.F"
#include "dbcsr_mm_cannon_s.F"

File diff suppressed because it is too large Load diff

View file

@ -5,378 +5,604 @@
! *****************************************************************************
!> \brief Fourth layer of the dbcsr matrix-matrix multiplication.
!> It hides the differences between performing calculations on the
!> GPU or the CPU.
!> It hides the differences between performing calculations on the
!> GPU or the CPU.
!>
!> \author Urban Borstnik
!>
!>
!> <b>Modification history:</b>
!> - 2010-02-23 Moved from dbcsr_operations
!> - 2011-11 Moved parameter-stack processing routines to
!> dbcsr_mm_methods.
!> - 2013-01 reorganized code (Ole Schuett)
!> - 2013-01 extensive refactoring (Ole Schuett)
! *****************************************************************************
MODULE dbcsr_mm_driver
USE dbcsr_block_operations, ONLY: block_add
USE dbcsr_config, ONLY: mm_driver,&
USE dbcsr_block_operations, ONLY: block_add,&
dbcsr_data_clear
USE dbcsr_config, ONLY: dbcsr_get_conf_nstacks,&
mm_driver,&
mm_driver_cuda,&
mm_stack_size,&
mm_thread_workshare
mm_stack_size
USE dbcsr_cuda_device, ONLY: dbcsr_cuda_create_streams,&
dbcsr_cuda_destroy_streams
dbcsr_cuda_destroy_streams,&
dbcsr_cuda_device_sync,&
dbcsr_cuda_init
USE dbcsr_cuda_memory, ONLY: dbcsr_cuda_dev_mem_alloc,&
dbcsr_cuda_dev_mem_hold,&
dbcsr_cuda_dev_mem_dealloc,&
dbcsr_cuda_dev_mem_new,&
dbcsr_cuda_dev_mem_realloc,&
dbcsr_cuda_dev_mem_release,&
dbcsr_cuda_dev_mem_zero
dbcsr_cuda_dev_mem_zero,&
dbcsr_cuda_host_mem_alloc,&
dbcsr_cuda_host_mem_dealloc
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
USE dbcsr_cuda_types, ONLY: dbcsr_cuda_mem_type
USE dbcsr_cuda_types, ONLY: dbcsr_cuda_mem_p_type,&
dbcsr_cuda_mem_type
USE dbcsr_data_methods, ONLY: dbcsr_data_ensure_size,&
dbcsr_data_get_size,&
dbcsr_data_get_type,&
dbcsr_data_init,&
dbcsr_data_new,&
dbcsr_data_release
dbcsr_data_release,&
dbcsr_data_set_size_referenced
USE dbcsr_error_handling
USE dbcsr_kinds, ONLY: int_4,&
USE dbcsr_kinds, ONLY: dp,&
int_8,&
real_8
USE dbcsr_mm_types, ONLY: carrier_type
USE dbcsr_pq_methods, ONLY: dbcsr_ps_set_get_group_p,&
dbcsr_ps_target_add_data_cuda,&
dbcsr_psg_add_data_cuda_ab,&
dbcsr_psg_rm_data_cuda_ab
USE dbcsr_pq_types, ONLY: dbcsr_ps_group_type,&
dbcsr_ps_state_empty,&
dbcsr_ps_target_obj,&
dbcsr_ps_target_type,&
dbcsr_ps_width
USE dbcsr_mm_stack, ONLY: process_mm_stack_cpu,&
process_mm_stack_cuda
USE dbcsr_mm_types, ONLY: dbcsr_ps_state_empty,&
dbcsr_ps_state_working,&
dbcsr_ps_width,&
driver_carrier_type,&
stack_descriptor_type
USE dbcsr_types, ONLY: dbcsr_data_obj,&
dbcsr_type,&
dbcsr_type_int_4
!$ USE OMP_LIB
IMPLICIT NONE
PRIVATE
CHARACTER(len=*), PARAMETER, PRIVATE :: moduleN = 'dbcsr_mm_driver'
PUBLIC :: dbcsr_mm_driver_outer_init, dbcsr_mm_driver_outer_finalize
PUBLIC :: dbcsr_mm_driver_inner_init
PUBLIC :: dbcsr_mm_driver_lib_init, dbcsr_mm_driver_lib_finalize
PUBLIC :: dbcsr_mm_driver_outer_init, dbcsr_mm_driver_outer_finalize
PUBLIC :: dbcsr_mm_driver_inner_init, dbcsr_mm_driver_inner_finalize
PUBLIC :: dbcsr_mm_driver_process_stack
! ===== Global Cuda Memory =====
! Allocating memory for cuda (on the card and host-pinned) is slow.
! Therefore, the memory is allocated once and stored in global variables.
! Variable cuda_mem_in_use should prevent concurred use of this global memory.
LOGICAL cuda_mem_in_use
!$OMP THREADPRIVATE(cuda_mem_in_use)
TYPE(dbcsr_cuda_mem_type), POINTER :: cuda_dev_c_data, cuda_dev_c_locks
!$OMP THREADPRIVATE(cuda_dev_c_data, cuda_dev_c_locks)
TYPE(dbcsr_cuda_mem_p_type), DIMENSION(:), POINTER :: cuda_dev_stacklist_data
!$OMP THREADPRIVATE(cuda_dev_stacklist_data)
INTEGER, DIMENSION(:,:,:,:), POINTER :: cuda_host_stackgrid_data
!$OMP THREADPRIVATE(cuda_host_stackgrid_data)
INTEGER, DIMENSION(:,:), POINTER :: cuda_host_stackgrid_state
!$OMP THREADPRIVATE(cuda_host_stackgrid_state)
!These are shared between threads
TYPE(dbcsr_cuda_mem_type), POINTER, SAVE :: cuda_dev_a_data => Null()
TYPE(dbcsr_cuda_mem_type), POINTER, SAVE :: cuda_dev_b_data => Null()
TYPE(dbcsr_cuda_mem_type), POINTER, SAVE :: cuda_dev_state_empty => Null()
!INTEGER , SAVE :: num_cpu_stacks = 0
!INTEGER , SAVE :: num_gpu_stacks = 0
!INTEGER(kind=int_8), SAVE :: gpu_flop = 0
!INTEGER(kind=int_8), SAVE :: cpu_flop = 0
CONTAINS
! *****************************************************************************
!> \brief Moved out of internal_operations
!> This loads the data for matrizes A and B onto the card.
!> \author Ole Schuett
! \brief Initialize the library
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_driver_inner_init(left, right, carrier, error)
TYPE(dbcsr_type), INTENT(IN) :: left, right
TYPE(carrier_type), INTENT(INOUT) :: carrier
SUBROUTINE dbcsr_mm_driver_lib_init(error)
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_driver_inner_init', &
routineP = moduleN//':'//routineN
INTEGER :: i, ithread, n_stack_buffers, &
nstacks, nthreads
INTEGER, DIMENSION(1), TARGET :: tmp_arr
INTEGER, DIMENSION(3) :: nxstacks
INTEGER, DIMENSION(:), POINTER :: tmp_ptr
IF (mm_driver .EQ. mm_driver_cuda) THEN
!$OMP MASTER
CALL dbcsr_cuda_cp_host_to_dev(&
left%data_area,&
carrier%param_sets%groups(1,1)%master%s%left_data_cuda,&
async = .TRUE.,&
error=error)
CALL dbcsr_cuda_cp_host_to_dev(&
right%data_area,&
carrier%param_sets%groups(1,1)%master%s%right_data_cuda,&
async = .TRUE.,&
error=error)
!$OMP END MASTER
nthreads = 1; ithread = 0
!$ nthreads = OMP_GET_NUM_THREADS () ; ithread = OMP_GET_THREAD_NUM ()
IF (mm_driver == mm_driver_cuda) THEN
CALL dbcsr_get_conf_nstacks(nxstacks, n_stack_buffers, error)
nstacks = nxstacks(1) * nxstacks(2) * nxstacks(3) + 1
!Ensure that lock_offset in dbcsr_cuda_calc.cu won't overflow (it's a int4).
! TODO: Check this code, before you start to used locking again.
IF(4*mm_stack_size > 2**17)& ! The "4" due to stack_mm_mnk_sq5_d
STOP "dbcsr_mm_driver_lib_init: mm_stack_size is too large"
IF(nthreads*nstacks*n_stack_buffers > 2**15)&
STOP "dbcsr_mm_driver_lib_init: nthreads*nstacks*n_stack_buffers is too large"
!TODO: do we want to support multiple cards? Search: ma_set_gpu_affinity
CALL dbcsr_cuda_init(error=error)
! 1. Things that are shared between threads.
!-----------------------------------------------------------------------
!$OMP BARRIER
!$OMP MASTER
CALL dbcsr_cuda_create_streams(nthreads*nstacks*n_stack_buffers, error)
! Constant dbcsr_ps_state_empty needed on card to signal completed transfers.
tmp_arr = dbcsr_ps_state_empty
tmp_ptr => tmp_arr
CALL dbcsr_cuda_dev_mem_new(cuda_dev_state_empty, dbcsr_type_int_4, error=error)
CALL dbcsr_cuda_dev_mem_alloc(cuda_dev_state_empty, 1, error=error)
CALL dbcsr_cuda_cp_host_to_dev(host_mem=tmp_ptr, dev_mem=cuda_dev_state_empty%d_i, n=1,&
async=.FALSE.,error=error)
!$OMP END MASTER
!$OMP BARRIER
! 2. Things that are local to each thread
!------------------------------------------------------------------------
! Allocate the parameter stack on the card
ALLOCATE(cuda_dev_stacklist_data(nstacks*n_stack_buffers))
DO i=1, nstacks*n_stack_buffers
NULLIFY(cuda_dev_stacklist_data(i)%p)
CALL dbcsr_cuda_dev_mem_new(cuda_dev_stacklist_data(i)%p, dbcsr_type_int_4, error)
CALL dbcsr_cuda_dev_mem_alloc(cuda_dev_stacklist_data(i)%p,&
dbcsr_ps_width*mm_stack_size, error=error)
END DO
! Allocate the parameter stack on the host with pinned memory
CALL dbcsr_cuda_host_mem_alloc(cuda_host_stackgrid_data, &
dbcsr_ps_width, mm_stack_size, nstacks, n_stack_buffers, error=error)
CALL dbcsr_cuda_host_mem_alloc(cuda_host_stackgrid_state, &
nstacks, n_stack_buffers, error=error)
! Size of the C-Blocks unkown at this point, allocation is done late.
NULLIFY(cuda_dev_c_data, cuda_dev_c_locks)
!Setup safety guard against overlapping multiplications.
cuda_mem_in_use = .FALSE.
ENDIF
END SUBROUTINE dbcsr_mm_driver_inner_init
END SUBROUTINE dbcsr_mm_driver_lib_init
! *****************************************************************************
!> \brief Moved out of dbcsr_multrec_init (actually csr_multiply_outer_init)
!> This allocates the memory for the matrizes A and B on the card and
!> and leaves pointers to them in every stack.
! \brief Finalize the library
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_driver_lib_finalize(error)
USE ISO_C_BINDING
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
INTEGER :: i
IF (mm_driver == mm_driver_cuda) THEN
! Releasing host pinned memory of stackgrid on the host
CALL dbcsr_cuda_host_mem_dealloc(cuda_host_stackgrid_data, error)
CALL dbcsr_cuda_host_mem_dealloc(cuda_host_stackgrid_state, error)
! releasing stuff on the card if necessary
IF(ASSOCIATED(cuda_dev_c_data))&
CALL dbcsr_cuda_dev_mem_release(cuda_dev_c_data, error=error)
IF(ASSOCIATED(cuda_dev_c_locks))&
CALL dbcsr_cuda_dev_mem_release(cuda_dev_c_locks, error=error)
DO i=1, SIZE(cuda_dev_stacklist_data)
CALL dbcsr_cuda_dev_mem_release(cuda_dev_stacklist_data(i)%p, error=error)
END DO
DEALLOCATE(cuda_dev_stacklist_data)
!$OMP BARRIER
!$OMP MASTER
IF(ASSOCIATED(cuda_dev_a_data))&
CALL dbcsr_cuda_dev_mem_release(cuda_dev_a_data, error=error)
IF(ASSOCIATED(cuda_dev_b_data))&
CALL dbcsr_cuda_dev_mem_release(cuda_dev_b_data, error=error)
CALL dbcsr_cuda_dev_mem_release(cuda_dev_state_empty, error=error)
CALL dbcsr_cuda_destroy_streams(error=error)
!How much memory is still allocated on the card?
!istat = dbcsr_cuda_dev_mem_info(mem_free, mem_avail)
!WRITE (*,*) "after outer_finalize cuda mem: ",mem_free, mem_avail, istat
!$OMP END MASTER
!$OMP BARRIER
END IF
! !$OMP MASTER
! WRITE (*,*) "num_stacks gpu:", num_gpu_stacks, " cpu: ",num_cpu_stacks, " ratio: ",REAL(num_gpu_stacks)/REAL(num_cpu_stacks)
! WRITE (*,*) "flop gpu:", gpu_flop, " cpu: ",cpu_flop, " ratio: ",REAL(gpu_flop)/REAL(cpu_flop)
! !$OMP END MASTER
END SUBROUTINE dbcsr_mm_driver_lib_finalize
! *****************************************************************************
!> \brief Initializes a multiplication cycle for new set of C-blocks.
!> \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_driver_outer_init(left, right, product, product_target, carrier, fill_guess, error)
SUBROUTINE dbcsr_mm_driver_outer_init(left, right, product, drv_carrier,&
fill_guess, nstacks, n_stack_buffers, stackgrid_data, stackgrid_state, error)
TYPE(dbcsr_type), INTENT(IN) :: left, right
TYPE(dbcsr_type), INTENT(INOUT) :: product
TYPE(dbcsr_ps_target_obj), INTENT(INOUT) :: product_target
TYPE(carrier_type), INTENT(INOUT) :: carrier
TYPE(driver_carrier_type), INTENT(INOUT) :: drv_carrier
REAL(KIND=real_8) :: fill_guess
INTEGER, INTENT(IN) :: nstacks, n_stack_buffers
INTEGER, DIMENSION(:, :, :, :), POINTER :: stackgrid_data
INTEGER, DIMENSION(:, :), POINTER :: stackgrid_state
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_driver_outer_init', &
routineP = moduleN//':'//routineN
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_driver_outer_init'
INTEGER :: data_type, nthreads, &
ps_buffer, ps_memreg
TYPE(dbcsr_cuda_mem_type), POINTER, SAVE :: a_dev, b_dev
TYPE(dbcsr_ps_group_type), POINTER :: ps_group
INTEGER :: a_data_size, allocated_datasize, b_data_size, c_data_size, &
c_locks_size, data_type, error_handler, ithread, nthreads, used_datasize
nthreads = 1
!$ nthreads = OMP_GET_NUM_THREADS ()
! first driver init part ---------------------------------------------------
CALL dbcsr_error_set(routineN, error_handler, error)
IF (mm_driver .EQ. mm_driver_cuda) THEN
!$OMP MASTER
!
! Setup threads<->streams
CALL dbcsr_cuda_create_streams (nthreads, error)
!
data_type = dbcsr_data_get_type (left%data_area)
CALL dbcsr_cuda_dev_mem_new (a_dev, data_type, error=error)
CALL dbcsr_cuda_dev_mem_new (b_dev, data_type, error=error)
CALL dbcsr_cuda_dev_mem_alloc (a_dev,&
dbcsr_data_get_size (left%data_area), error=error)
CALL dbcsr_cuda_dev_mem_alloc (b_dev,&
dbcsr_data_get_size (right%data_area), error=error)
!$OMP END MASTER
CALL init_card_c (product, product_target%t,fill_guess, error)
!$OMP BARRIER
IF (.FALSE.) THEN
! Left in for mental accounting
!$OMP CRITICAL
CALL dbcsr_cuda_dev_mem_hold (a_dev, error=error)
CALL dbcsr_cuda_dev_mem_hold (b_dev, error=error)
!$OMP END CRITICAL
ENDIF
nthreads = 1; ithread = 0
!$ nthreads = OMP_GET_NUM_THREADS () ; ithread = OMP_GET_THREAD_NUM ()
data_type = dbcsr_data_get_type(left%data_area)
drv_carrier%product_wm => product%wms(ithread+1) !TODO: why do we need this ref?
used_datasize = drv_carrier%product_wm%datasize
allocated_datasize = dbcsr_data_get_size(drv_carrier%product_wm%data_area)
! The product's data_area could already contain some data.
! ( see: keep_product_data in dbcsr_operations.F )
! But this data might not occupy all the allocated memory in the data_area.
! Since, we don't want to keep track of unitialized memory we just zero it now.
CALL dbcsr_data_clear(drv_carrier%product_wm%data_area, lb=used_datasize+1, ub=allocated_datasize)
IF (mm_driver /= mm_driver_cuda) THEN
ALLOCATE(stackgrid_data(dbcsr_ps_width, mm_stack_size, nstacks, n_stack_buffers))
ALLOCATE(stackgrid_state(nstacks, n_stack_buffers))
ELSE
!$OMP MASTER
NULLIFY (a_dev, b_dev)
!$OMP END MASTER
!$OMP BARRIER
ENDIF
! end of first part of driver init -----------------------------------------
! second part of driver init -----------------------------------------------
IF (mm_driver .EQ. mm_driver_cuda) THEN
DO ps_memreg = 1, carrier%param_sets%nmemregs
DO ps_buffer = 1, carrier%param_sets%nbuffers
ps_group => dbcsr_ps_set_get_group_p (carrier%param_sets,&
ps_buffer, ps_memreg, error=error)
!$OMP CRITICAL (crit_data_card)
CALL dbcsr_psg_add_data_cuda_ab (ps_group,&
a_dev, b_dev, error=error)
!$OMP END CRITICAL (crit_data_card)
ENDDO
ENDDO
IF (.FALSE.) THEN
! Left in for mental accounting
!$OMP CRITICAL
CALL dbcsr_cuda_dev_mem_release (a_dev, error=error)
CALL dbcsr_cuda_dev_mem_release (b_dev, error=error)
!$OMP END CRITICAL
ENDIF
! 0. Making shure there is only one multiplication going on at a time.
!------------------------------------------------------------------------
IF(.NOT. ASSOCIATED(cuda_dev_state_empty)) &
STOP "dbcsr_mm_driver_outer_init: dbcsr not initialized"
IF(cuda_mem_in_use) STOP "dbcsr_mm_driver: found overlapping multiplications"
cuda_mem_in_use = .TRUE.
! 1. Things that are shared between threads.
!------------------------------------------------------------------------
!$OMP BARRIER
!$OMP MASTER
CALL dbcsr_cuda_dev_mem_release (a_dev, error=error)
CALL dbcsr_cuda_dev_mem_release (b_dev, error=error)
! (re)Allocate Memory for A and B blocks on the Card
a_data_size = dbcsr_data_get_size(left%data_area)
b_data_size = dbcsr_data_get_size(right%data_area)
CALL cuda_prepare_mem(cuda_dev_a_data, data_type, a_data_size, error=error)
CALL cuda_prepare_mem(cuda_dev_b_data, data_type, b_data_size, error=error)
!$OMP END MASTER
!$OMP BARRIER
! 2. Things that are local to each thread
!------------------------------------------------------------------------
IF(SIZE(cuda_host_stackgrid_data, 1) /= dbcsr_ps_width .OR. &
SIZE(cuda_host_stackgrid_data, 2) /= mm_stack_size .OR. &
SIZE(cuda_host_stackgrid_data, 3) /= nstacks .OR. &
SIZE(cuda_host_stackgrid_data, 4) /= n_stack_buffers ) &
STOP "dbcsr_mm_driver_outer_init: cuda_host_stackgrid_data wrong dims"
IF(SIZE(cuda_host_stackgrid_state, 1) /= nstacks .OR. &
SIZE(cuda_host_stackgrid_state, 2) /= n_stack_buffers ) &
STOP "dbcsr_mm_driver_outer_init: cuda_host_stackgrid_state wrong dims"
stackgrid_data => cuda_host_stackgrid_data
stackgrid_state => cuda_host_stackgrid_state
! (re)Allocate Memory for C blocks on the Card and zero it.
c_data_size = dbcsr_data_get_size(product%wms(ithread+1)%data_area)
CALL cuda_prepare_mem(cuda_dev_c_data, data_type, c_data_size, zero=.TRUE., error=error)
! (re)Allocate and zero C locks on GPU.
c_locks_size = INT(product%nblkcols_local,KIND=int_8)* &
INT(product%nblkrows_local,KIND=int_8) * fill_guess + 1
CALL cuda_prepare_mem(cuda_dev_c_locks, dbcsr_type_int_4, c_locks_size, zero=.TRUE., error=error)
ENDIF
CALL dbcsr_error_stop(error_handler, error)
END SUBROUTINE dbcsr_mm_driver_outer_init
! *****************************************************************************
!> \brief Moved out of dbcsr_multrec_finalize (actually csr_multiply_outer_finalize)
!> Copies the result matrix C back from the card,
!> deallocates the memory on the card, and destroys the streams.
! \brief Helper routine used by dbcsr_mm_driver_outer_init().
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE cuda_prepare_mem(dev_mem, data_type, data_size, zero, error)
TYPE(dbcsr_cuda_mem_type), POINTER :: dev_mem
INTEGER, INTENT(IN) :: data_type, data_size
LOGICAL, INTENT(IN), OPTIONAL :: zero
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
IF (.NOT. ASSOCIATED(dev_mem)) THEN
!WRITE (*,*) "cuda_prepare_mem: allocting for the first time"
CALL dbcsr_cuda_dev_mem_new(dev_mem, data_type, error=error)
CALL dbcsr_cuda_dev_mem_alloc(dev_mem, data_size, error=error)
END IF
IF (dev_mem%data_type /= data_type) THEN
!WRITE (*,*) "cuda_prepare_mem: changing datatype"
CALL dbcsr_cuda_dev_mem_release(dev_mem, error=error)
CALL dbcsr_cuda_dev_mem_new(dev_mem, data_type, error=error)
CALL dbcsr_cuda_dev_mem_alloc(dev_mem, data_size, error=error)
END IF
IF (dbcsr_cuda_dev_mem_get_alloc(dev_mem) < data_size) THEN
!WRITE (*,*) "cuda_prepare_mem: growing dev_mem"
CALL dbcsr_cuda_dev_mem_dealloc(dev_mem, error=error)
CALL dbcsr_cuda_dev_mem_alloc(dev_mem, data_size, error=error)
END IF
IF(PRESENT(zero)) THEN
IF(zero) CALL dbcsr_cuda_dev_mem_zero(dev_mem, first=1, &
last=dbcsr_cuda_dev_mem_get_alloc(dev_mem), error=error)
END IF
END SUBROUTINE
! *****************************************************************************
!> \brief Initializes a multiplication cycle for a new set of A- and B-blocks.
!> \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_driver_outer_finalize(carrier, product, error)
TYPE(carrier_type), INTENT(INOUT) :: carrier
TYPE(dbcsr_type), INTENT(inout) :: product
SUBROUTINE dbcsr_mm_driver_inner_init(left, right, drv_carrier, error)
TYPE(dbcsr_type), INTENT(IN) :: left, right
TYPE(driver_carrier_type), INTENT(INOUT) :: drv_carrier
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
INTEGER :: data_size, data_type, &
ithread, ps_buffer, ps_memreg
TYPE(dbcsr_data_obj) :: tmp_data
TYPE(dbcsr_ps_group_type), POINTER :: ps_group
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_driver_inner_init'
IF (mm_driver .EQ. mm_driver_cuda) THEN
! this part is from dbcsr_internal_operations ---------------------------
ithread = 0
!$ ithread = omp_get_thread_num()
! Transfer data from GPU to host.
CALL dbcsr_data_init (tmp_data)
CALL dbcsr_data_new (tmp_data,&
data_type = dbcsr_data_get_type(product%wms(ithread+1)%data_area),&
data_size = carrier%datasize)
CALL dbcsr_cuda_cp_dev_to_host(&
carrier%param_sets%groups(1,1)%master%s%t%t%product_data_cuda,&
!product%wms(1)%data_area,&
tmp_data,&
error=error)
CALL dbcsr_data_ensure_size(product%wms(ithread+1)%data_area,&
carrier%datasize, error=error)
CALL block_add (product%wms(ithread+1)%data_area, tmp_data,&
len=carrier%datasize, error=error)
CALL dbcsr_data_release (tmp_data)
! end of part from dbcsr_internal_operations ----------------------------
! this part is from dbcsr_multrec_finalize ------------------------------
DO ps_memreg = 1, carrier%param_sets%nmemregs
DO ps_buffer = 1, carrier%param_sets%nbuffers
ps_group => dbcsr_ps_set_get_group_p (carrier%param_sets,&
ps_buffer, ps_memreg, error=error)
!$OMP CRITICAL (crit_data_card)
CALL dbcsr_psg_rm_data_cuda_ab (ps_group, error=error)
!$OMP END CRITICAL (crit_data_card)
ENDDO
ENDDO
ENDIF
INTEGER :: error_handler
!$OMP BARRIER
!$OMP MASTER
IF (mm_driver .EQ. mm_driver_cuda) THEN
CALL dbcsr_cuda_destroy_streams (error=error)
ENDIF
!$OMP END MASTER
END SUBROUTINE dbcsr_mm_driver_outer_finalize
! *****************************************************************************
!> \brief Perform allocations and setup needed for multiplication on an
!> accelerator card.
! *****************************************************************************
SUBROUTINE init_card_c (host_matrix, product_target, fill_guess, error)
TYPE(dbcsr_type), INTENT(IN) :: host_matrix
TYPE(dbcsr_ps_target_type), &
INTENT(INOUT) :: product_target
REAL(KIND=real_8) :: fill_guess
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CALL dbcsr_error_set(routineN, error_handler, error)
CHARACTER(len=*), PARAMETER :: routineN = 'init_card_c', &
routineP = moduleN//':'//routineN
drv_carrier%left_data_area = left%data_area
drv_carrier%right_data_area = right%data_area
INTEGER :: clocks_size, data_size, &
data_type, error_handle, &
ithread
INTEGER, POINTER :: state_tmp
INTEGER, TARGET :: state_tmp_tgt
LOGICAL :: do_master_work
TYPE(dbcsr_cuda_mem_type), POINTER :: c_dev, c_locks_dev, &
params_dev_l
TYPE(dbcsr_cuda_mem_type), POINTER, SAVE :: params_dev, stack_state_dev
! ---------------------------------------------------------------------------
CALL dbcsr_error_set(routineN, error_handle, error)
! Only the master thread does common work in when threads share work.
do_master_work = .NOT. mm_thread_workshare
do_master_work = .TRUE.
!$OMP MASTER
do_master_work = .TRUE.
!$OMP END MASTER
ithread = 0
!$ ithread = OMP_GET_THREAD_NUM ()
data_type = dbcsr_data_get_type (host_matrix%wms(ithread+1)%data_area)
data_size = dbcsr_data_get_size (host_matrix%wms(ithread+1)%data_area)
!
! Allocate space for product data on the card.
CALL dbcsr_cuda_dev_mem_new (c_dev, data_type, error)
CALL dbcsr_cuda_dev_mem_alloc (c_dev, data_size, error=error)
!
! To finish allocations before the zeroing is launched.
!
!$OMP BARRIER
!
! Zero C data on card. This assumes that data will be
! summed with the host data after the multiplication is done.
CALL dbcsr_cuda_dev_mem_zero(c_dev, first=1,&
last=data_size, error=error)
!
! Attach C data on card to the product target descriptor
CALL dbcsr_ps_target_add_data_cuda (product_target, c_dev, error)
CALL dbcsr_cuda_dev_mem_release (c_dev, error=error)
!
! Allocate C locks on GPU. Locks should be cleared as soon as
! allocated.
CALL dbcsr_cuda_dev_mem_new (c_locks_dev,&
dbcsr_type_int_4, error)
clocks_size = INT(host_matrix%nblkcols_local,KIND=int_8)* &
INT(host_matrix%nblkrows_local,KIND=int_8)* &
fill_guess
clocks_size = MAX(1,INT(clocks_size, KIND=int_4))
CALL dbcsr_cuda_dev_mem_alloc (c_locks_dev,INT(clocks_size,KIND=int_4), error=error)
CALL dbcsr_cuda_dev_mem_zero (c_locks_dev,&
first=1,last=dbcsr_cuda_dev_mem_get_alloc(c_locks_dev),&
error=error)
product_target%c_locks_dev => c_locks_dev
!
! Have a source from which to update states and setup the on-card
! parameter stacks.
!
!$OMP MASTER
CALL dbcsr_cuda_dev_mem_new (stack_state_dev, dbcsr_type_int_4,&
error=error)
CALL dbcsr_cuda_dev_mem_alloc (stack_state_dev, 1, error=error)
state_tmp_tgt = dbcsr_ps_state_empty
state_tmp => state_tmp_tgt
CALL dbcsr_cuda_cp_host_to_dev (state_tmp, stack_state_dev%d_i, error=error)
!$OMP END MASTER
IF (do_master_work) THEN
! Allocate the on-card parameter stack.
CALL dbcsr_cuda_dev_mem_new (params_dev_l, dbcsr_type_int_4, error)
CALL dbcsr_cuda_dev_mem_alloc (params_dev_l, mm_stack_size*dbcsr_ps_width,&
error=error)
IF (mm_driver == mm_driver_cuda) THEN
!$OMP BARRIER
!$OMP MASTER
params_dev => params_dev_l
!TODO: could be async if we use events to make the other streams wait
! for a and b to be uploaded.
CALL dbcsr_cuda_cp_host_to_dev(host_mem=left%data_area,&
dev_mem=cuda_dev_a_data, async = .FALSE., error=error)
CALL dbcsr_cuda_cp_host_to_dev(host_mem=right%data_area,&
dev_mem=cuda_dev_b_data, async = .FALSE., error=error)
!$OMP END MASTER
!$OMP BARRIER
ENDIF
!$OMP BARRIER
!$OMP CRITICAL (crit_data_card)
!
IF (.FALSE.) THEN
! Left in for mental accounting
CALL dbcsr_cuda_dev_mem_hold (stack_state_dev, error=error)
CALL dbcsr_error_stop(error_handler, error)
END SUBROUTINE dbcsr_mm_driver_inner_init
! *****************************************************************************
!> \brief Finalizes a multiplication cycle for a set of A- and B-blocks.
!> \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_driver_inner_finalize(drv_carrier, error)
TYPE(driver_carrier_type), INTENT(INOUT) :: drv_carrier
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
IF (mm_driver == mm_driver_cuda) THEN
!TODO: waiting for ALL streams - should only wait for those of this thread
!$OMP BARRIER
CALL dbcsr_cuda_device_sync(error)
!$OMP BARRIER
END IF
END SUBROUTINE dbcsr_mm_driver_inner_finalize
! *****************************************************************************
!> \brief Finalizes a multiplication cycle for a set of C-blocks.
!> \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_driver_outer_finalize(drv_carrier, product,stackgrid_data, stackgrid_state, error)
USE ISO_C_BINDING
TYPE(driver_carrier_type), INTENT(INOUT) :: drv_carrier
TYPE(dbcsr_type), INTENT(inout) :: product
INTEGER, DIMENSION(:, :, :, :), POINTER :: stackgrid_data
INTEGER, DIMENSION(:, :), POINTER :: stackgrid_state
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
INTEGER :: data_type, datasize, ithread
TYPE(dbcsr_data_obj) :: tmp_data
ithread = 0
!$ ithread = omp_get_thread_num()
IF (mm_driver /= mm_driver_cuda) THEN
DEALLOCATE(stackgrid_data)
DEALLOCATE(stackgrid_state)
ELSE
! Host-pinned memory not dealloced, will be reused by next multiplication
NULLIFY(stackgrid_data, stackgrid_state)
! Transfer C-data from GPU to host and adding it to host's result
datasize = product%wms(ithread+1)%datasize
data_type = dbcsr_data_get_type(product%wms(ithread+1)%data_area)
CALL dbcsr_data_init(tmp_data)
CALL dbcsr_data_new(tmp_data, data_type=data_type, data_size=datasize)
CALL dbcsr_cuda_cp_dev_to_host(dev_mem=cuda_dev_c_data,&
host_mem=tmp_data, error=error)
CALL block_add(product%wms(ithread+1)%data_area, tmp_data,&
len=datasize, error=error)
CALL dbcsr_data_release(tmp_data)
cuda_mem_in_use = .FALSE.
ENDIF
! all threads share the master's stack_state_dev
product_target%stack_state_dev => stack_state_dev
IF (do_master_work) THEN
product_target%params_dev => params_dev_l
ELSE
product_target%params_dev => params_dev
ENDIF
CALL dbcsr_cuda_dev_mem_hold (product_target%params_dev, error=error)
IF (.FALSE.) THEN
! Left in for mental accounting
CALL dbcsr_cuda_dev_mem_release (stack_state_dev, error=error)
ENDIF
!
!$OMP END CRITICAL (crit_data_card)
!$OMP BARRIER
IF (do_master_work) THEN
!$OMP CRITICAL (crit_data_card)
CALL dbcsr_cuda_dev_mem_release (params_dev_l, error=error)
!$OMP END CRITICAL (crit_data_card)
ENDIF
!
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE init_card_c
END SUBROUTINE dbcsr_mm_driver_outer_finalize
! *****************************************************************************
! \brief Processes a given stack.
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_driver_process_stack(drv_carrier, driver, stack_data,&
stack_fillcount, stack_state, stack_descr, stack_id, error)
TYPE(driver_carrier_type), INTENT(INOUT) :: drv_carrier
INTEGER, INTENT(IN) :: driver
INTEGER, DIMENSION(:, :), POINTER :: stack_data
INTEGER, POINTER :: stack_fillcount, stack_state
TYPE(stack_descriptor_type), POINTER :: stack_descr
INTEGER, INTENT(IN) :: stack_id
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
REAL, PARAMETER :: resize_factor = 1.618034
INTEGER :: allocated_datasize, nblks, &
new_c_size, new_locks_size, &
old_c_size, stacked_datasize
!INTEGER(kind=int_8) :: flop
!TODO: make difference clear between driver, mm_driver and mm_host_driver.
! collect statistics
stack_descr%total_entries = stack_descr%total_entries + stack_fillcount
stacked_datasize = drv_carrier%product_wm%datasize
allocated_datasize = dbcsr_data_get_size(drv_carrier%product_wm%data_area)
! tell the data_area how much of its memory is actually beeing used.
CALL dbcsr_data_set_size_referenced(drv_carrier%product_wm%data_area, stacked_datasize)
! Resizing the product data (on host and card) ---------------------------
IF (stacked_datasize > allocated_datasize) THEN
! Resize the target data area if the stack references a target
! matrix data element outside of its current range. When
! appropriate a lock on the data area is taken.
!WRITE (*,*) "dbcsr_mm_driver_process_stack: We need to resize c data"
!IF (dbg) &
! WRITE(*,*)routineN//" Resizing to", LOG(REAL(maxs))/LOG(10.0)
CALL dbcsr_data_ensure_size(drv_carrier%product_wm%data_area,&
stacked_datasize, factor=resize_factor, zero_pad=.TRUE., error=error)
IF(stacked_datasize > dbcsr_data_get_size(drv_carrier%product_wm%data_area))&
STOP "dbcsr_mm_driver_process_stack: resizing failed"
! Cuda on-device resizing
IF (mm_driver == mm_driver_cuda) THEN
! All kernels in the stream must be finished before
! reallocation is performed.
!TODO: waiting for ALL streams - should only wait for those of this thread
CALL dbcsr_cuda_device_sync(error)
! dbcsr_data_ensure_size oversizes the array - we need to know how much.
old_c_size = dbcsr_cuda_dev_mem_get_alloc(cuda_dev_c_data)
new_c_size = dbcsr_data_get_size(drv_carrier%product_wm%data_area)
CALL dbcsr_cuda_dev_mem_realloc(cuda_dev_c_data, new_c_size, error=error)
CALL dbcsr_cuda_dev_mem_zero(cuda_dev_c_data,&
first=old_c_size, last=new_c_size, error=error)
END IF
END IF
! Resizing the product locks (only on the card) --------------------------
IF (mm_driver == mm_driver_cuda) THEN
! Resize locks, which are equal to the new block count.
nblks = drv_carrier%product_wm%lastblk
IF (nblks > dbcsr_cuda_dev_mem_get_alloc(cuda_dev_c_locks)) THEN
!WRITE (*,*) "dbcsr_mm_driver_process_stack: We need to resize c_locks_cuda"
! All kernels in the stream must be finished before
! reallocation is performed.
!TODO: waiting for ALL streams - should only wait for those of this thread
CALL dbcsr_cuda_device_sync(error)
new_locks_size = INT(REAL(nblks*4,kind=dp)*resize_factor)
CALL dbcsr_cuda_dev_mem_realloc(cuda_dev_c_locks, new_locks_size, error=error)
CALL dbcsr_cuda_dev_mem_zero(cuda_dev_c_locks,&
first=1, last=new_locks_size, error=error)
END IF
END IF
!!From here on there is no boundary checking due to assumed-SIZE-arguments.
!!This is usefull to check stack parameters, BUT it works only for kind=dp
!DO sp = 1, stack_fillcount
! IF(stack_data(p_a_first,sp) > SIZE(drv_carrier%left_data_area%d%r_dp)) &
! STOP "left data out of range"
! IF(stack_data(p_b_first,sp) > SIZE(drv_carrier%right_data_area%d%r_dp)) &
! STOP "right data out of range"
! IF(stack_data(p_c_first,sp) > SIZE(drv_carrier%product_wm%data_area%d%r_dp)) THEN
! WRITE (*,*) "blub: ",stack_data(p_c_first,sp) , SIZE(drv_carrier%product_wm%data_area%d%r_dp), &
! dbcsr_data_get_size(drv_carrier%product_wm%data_area), stacked_datasize
! STOP "product data out of range"
! END IF
!END DO
! Submitting the stack for processing -------------------------------------
stack_state = dbcsr_ps_state_working
! flop = 2*stack_fillcount*stack_descr%max_m*stack_descr%max_n*stack_descr%max_k
IF (driver == mm_driver_cuda) THEN
! !$OMP critical
! num_gpu_stacks = num_gpu_stacks + 1
! gpu_flop = gpu_flop + flop
! !$OMP end critical
CALL process_mm_stack_cuda(params=stack_data,&
stack_size=stack_fillcount,&
data_a_dev=cuda_dev_a_data,&
data_b_dev=cuda_dev_b_data,&
data_c_dev=cuda_dev_c_data,&
c_locks=cuda_dev_c_locks,&
params_dev=cuda_dev_stacklist_data(stack_id)%p,&
m=stack_descr%m,&
n=stack_descr%n,&
k=stack_descr%k,&
max_m=stack_descr%max_m,&
max_n=stack_descr%max_n,&
max_k=stack_descr%max_k,&
defined_mnk=stack_descr%defined_mnk,&
state=stack_state,&
stack_state_dev=cuda_dev_state_empty,&
stack_id=stack_id,&
error=error)
ELSE
! !$OMP critical
! num_cpu_stacks = num_cpu_stacks + 1
! cpu_flop = cpu_flop + flop
! !$OMP end critical
CALL process_mm_stack_cpu(params=stack_data,&
driver=driver,&
stack_size=stack_fillcount, &
left_data_area=drv_carrier%left_data_area,&
right_data_area=drv_carrier%right_data_area,&
product_data_area=drv_carrier%product_wm%data_area,&
state=stack_state,&
m=stack_descr%m,&
n=stack_descr%n,&
k=stack_descr%k,&
max_m=stack_descr%max_m,&
max_n=stack_descr%max_n,&
max_k=stack_descr%max_k,&
defined_mnk=stack_descr%defined_mnk,&
stack_id=stack_id,&
error=error)
END IF
stack_fillcount = 0 !otherwise the stackcolumn gets processed again.
END SUBROUTINE dbcsr_mm_driver_process_stack
END MODULE dbcsr_mm_driver

View file

@ -13,14 +13,13 @@
!> - 2010-02-23 Moved from dbcsr_operations
!> - 2011-11 Moved parameter-stack processing routines to
!> dbcsr_mm_methods.
!> - 2013-01 reorganized code (Ole Schuett)
!> - 2013-01 extensive refactoring (Ole Schuett)
! *****************************************************************************
MODULE dbcsr_mm_multrec
USE array_types, ONLY: array_data,&
array_equality
USE dbcsr_config, ONLY: mm_thread_workshare
USE dbcsr_error_handling
USE dbcsr_kinds, ONLY: int_8,&
real_8,&
@ -31,29 +30,53 @@ MODULE dbcsr_mm_multrec
dbcsr_distribution_local_rows,&
dbcsr_distribution_row_dist,&
dbcsr_distribution_thread_dist
USE dbcsr_mm_csr, ONLY: csr_multiply_inner_finalize,&
csr_multiply_inner_init,&
csr_multiply_outer_finalize,&
csr_multiply_outer_init,&
csr_multiply_unwrap
USE dbcsr_mm_types, ONLY: carrier_type
USE dbcsr_mm_csr, ONLY: dbcsr_mm_csr_inner_finalize,&
dbcsr_mm_csr_inner_init,&
dbcsr_mm_csr_lib_finalize,&
dbcsr_mm_csr_lib_init,&
dbcsr_mm_csr_multiply,&
dbcsr_mm_csr_outer_finalize,&
dbcsr_mm_csr_outer_init
USE dbcsr_mm_types, ONLY: multrec_carrier_type
USE dbcsr_ptr_util, ONLY: ensure_array_size
USE dbcsr_types, ONLY: dbcsr_type
!$ USE OMP_LIB
IMPLICIT NONE
PRIVATE
CHARACTER(len=*), PARAMETER, PRIVATE :: moduleN = 'dbcsr_mm_multrec'
LOGICAL, PARAMETER :: careful_mod = .FALSE.
PUBLIC :: dbcsr_multrec_init
PUBLIC :: dbcsr_multrec_sparse
PUBLIC :: dbcsr_multrec_finalize
CONTAINS
PUBLIC :: dbcsr_mm_multrec_lib_init, dbcsr_mm_multrec_lib_finalize
PUBLIC :: dbcsr_mm_multrec_init, dbcsr_mm_multrec_finalize
PUBLIC :: dbcsr_mm_multrec_multiply
CONTAINS
! *****************************************************************************
! \brief Initialize the library
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_multrec_lib_init(error)
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CALL dbcsr_mm_csr_lib_init(error)
END SUBROUTINE
! *****************************************************************************
! \brief Finalize the library
! \author Ole Schuett
! *****************************************************************************
SUBROUTINE dbcsr_mm_multrec_lib_finalize(error)
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CALL dbcsr_mm_csr_lib_finalize(error)
END SUBROUTINE
! *****************************************************************************
!> \brief Sets up recursive multiplication
@ -65,23 +88,18 @@ CONTAINS
!> existing product matrix, default is no
!> \param[in] eps (optional) on-the-fly filtering epsilon
! *****************************************************************************
SUBROUTINE dbcsr_multrec_init(left, right, product, carrier,&
right_data_sr, right_data_rr, left_data_sr, left_data_rr, &
right_index_sr, right_index_rr, left_index_sr, left_index_rr, &
SUBROUTINE dbcsr_mm_multrec_init(left, right, product, multrec_carrier,&
keep_sparsity, eps, row_max_epss, fill_guess, error)
TYPE(dbcsr_type), INTENT(IN) :: left, right
TYPE(dbcsr_type), INTENT(INOUT) :: product
TYPE(carrier_type), INTENT(out) :: carrier
INTEGER, DIMENSION(:, :), POINTER :: right_data_sr, right_data_rr, &
left_data_sr, left_data_rr, right_index_sr, right_index_rr, &
left_index_sr, left_index_rr
TYPE(multrec_carrier_type), INTENT(out) :: multrec_carrier
LOGICAL, INTENT(IN) :: keep_sparsity
REAL(kind=real_8), INTENT(in), OPTIONAL :: eps
REAL(kind=sp), DIMENSION(:), INTENT(IN) :: row_max_epss
REAL(KIND=real_8) :: fill_guess
TYPE(dbcsr_error_type), INTENT(inout) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_multrec_init', &
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_multrec_init', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: dbg = .FALSE.
@ -90,35 +108,30 @@ CONTAINS
error_handler, i, ithread
INTEGER, DIMENSION(:), POINTER :: c_local_cols, c_local_rows, &
product_thread_dist
LOGICAL :: local_cols, local_indexing, &
local_rows
! ---------------------------------------------------------------------------
CALL dbcsr_error_set(routineN, error_handler, error)
ithread = 0
!$ ithread = OMP_GET_THREAD_NUM ()
carrier%id = ithread
!
! Ensures that the index is correctly defined.
CALL dbcsr_assert (left%list_indexing,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Must use list indexing for this routine.", __LINE__, error=error)
!
CALL dbcsr_assert ("NOT", left%bcsc,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Wrong routine for BCSC matrices.", __LINE__, error=error)
CALL dbcsr_assert ("NOT", right%bcsc,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Wrong routine for BCSC matrices.", __LINE__, error=error)
local_cols = right%local_indexing
local_rows = left%local_indexing
CALL dbcsr_assert (local_cols, "EQV", local_rows,&
CALL dbcsr_assert (right%local_indexing, "EQV", left%local_indexing,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Local index useage must be consistent.", __LINE__, error=error)
local_indexing = local_rows
IF (local_cols) THEN
IF (right%local_indexing) THEN
CALL dbcsr_assert (left%local_indexing,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"Wrong left format for local_cols.", __LINE__, error=error)
@ -134,28 +147,18 @@ CONTAINS
"Wrong right format for not local_cols.", __LINE__, error=error)
ENDIF
!
! Fill carrier data structure.
carrier%local_indexing = local_indexing
carrier%keep_sparsity = keep_sparsity
carrier%c_has_symmetry = product%symmetry
carrier%use_eps = PRESENT (eps)
carrier%my_wm = product%wms(ithread+1)
carrier%lastblk = product%wms(ithread+1)%lastblk
carrier%original_lastblk = carrier%lastblk
carrier%datasize = product%wms(ithread+1)%datasize
carrier%flop = INT(0, int_8)
carrier%right_data_sr => right_data_sr
carrier%right_data_rr => right_data_rr
carrier%right_index_sr => right_index_sr
carrier%right_index_rr => right_index_rr
carrier%left_data_sr => left_data_sr
carrier%left_data_rr => left_data_rr
carrier%left_index_sr => left_index_sr
carrier%left_index_rr => left_index_rr
! Fill multrec_carrier data structure.
multrec_carrier%local_indexing = left%local_indexing
multrec_carrier%keep_sparsity = keep_sparsity
multrec_carrier%c_has_symmetry = product%symmetry
multrec_carrier%use_eps = PRESENT (eps)
multrec_carrier%original_lastblk = product%wms(ithread+1)%lastblk
multrec_carrier%flop = INT(0, int_8)
IF (PRESENT (eps)) THEN
carrier%eps = eps
multrec_carrier%eps = eps
ELSE
carrier%eps = 0.0_real_8
multrec_carrier%eps = 0.0_real_8
ENDIF
!
!
@ -171,17 +174,17 @@ CONTAINS
c_local_rows => array_data (product%local_rows)
c_nblkcols_local = product%nblkcols_local
c_local_cols => array_data (product%local_cols)
IF (local_indexing) THEN
carrier%c_local_rows => c_local_rows
carrier%c_local_cols => c_local_cols
IF (multrec_carrier%local_indexing) THEN
multrec_carrier%c_local_rows => c_local_rows
multrec_carrier%c_local_cols => c_local_cols
ELSE
ALLOCATE (carrier%c_local_rows (product%nblkrows_total))
ALLOCATE (carrier%c_local_cols (product%nblkcols_total))
ALLOCATE (multrec_carrier%c_local_rows (product%nblkrows_total))
ALLOCATE (multrec_carrier%c_local_cols (product%nblkcols_total))
FORALL (i = 1 : product%nblkrows_total)
carrier%c_local_rows(i) = i
multrec_carrier%c_local_rows(i) = i
END FORALL
FORALL (i = 1 : product%nblkcols_total)
carrier%c_local_cols(i) = i
multrec_carrier%c_local_cols(i) = i
END FORALL
ENDIF
IF (dbg) WRITE(*,*)"setting up for product", product%name
@ -213,54 +216,57 @@ CONTAINS
ENDIF
!
! And the k epsilons
IF (local_indexing) THEN
ALLOCATE (carrier%row_max_epss(c_nblkrows_local))
IF (multrec_carrier%local_indexing) THEN
ALLOCATE (multrec_carrier%row_max_epss(c_nblkrows_local))
ELSE
ALLOCATE (carrier%row_max_epss(product%nblkrows_total))
ALLOCATE (multrec_carrier%row_max_epss(product%nblkrows_total))
ENDIF
IF (carrier%use_eps) THEN
IF (local_indexing) THEN
IF (multrec_carrier%use_eps) THEN
IF (multrec_carrier%local_indexing) THEN
CALL local_filter_sp(row_max_epss, c_nblkrows_local, c_local_rows,&
carrier%row_max_epss)
multrec_carrier%row_max_epss)
ELSE
IF (careful_mod) &
CALL dbcsr_assert (SIZE(row_max_epss) .EQ. SIZE(carrier%row_max_epss),&
CALL dbcsr_assert (SIZE(row_max_epss) .EQ. SIZE(multrec_carrier%row_max_epss),&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"max epss local/global mismatch.", __LINE__, error=error)
carrier%row_max_epss(:) = row_max_epss(:)
multrec_carrier%row_max_epss(:) = row_max_epss(:)
ENDIF
ELSE
carrier%row_max_epss(:) = -HUGE(0.0_sp)
multrec_carrier%row_max_epss(:) = -HUGE(0.0_sp)
ENDIF
!
IF (local_indexing) THEN
ALLOCATE(carrier%m_sizes(c_nblkrows_local))
IF (multrec_carrier%local_indexing) THEN
ALLOCATE(multrec_carrier%m_sizes(c_nblkrows_local))
CALL local_filter(array_data (product%row_blk_size), SIZE(c_local_rows),&
c_local_rows, carrier%m_sizes)
ALLOCATE(carrier%n_sizes(c_nblkcols_local))
c_local_rows, multrec_carrier%m_sizes)
ALLOCATE(multrec_carrier%n_sizes(c_nblkcols_local))
CALL local_filter(array_data (product%col_blk_size), SIZE(c_local_cols),&
c_local_cols, carrier%n_sizes)
c_local_cols, multrec_carrier%n_sizes)
ELSE
ALLOCATE(carrier%m_sizes(product%nblkrows_total))
carrier%m_sizes(:) = array_data (product%row_blk_size)
ALLOCATE(carrier%n_sizes(product%nblkcols_total))
carrier%n_sizes(:) = array_data (product%col_blk_size)
ALLOCATE(multrec_carrier%m_sizes(product%nblkrows_total))
multrec_carrier%m_sizes(:) = array_data (product%row_blk_size)
ALLOCATE(multrec_carrier%n_sizes(product%nblkcols_total))
multrec_carrier%n_sizes(:) = array_data (product%col_blk_size)
ENDIF
!
NULLIFY (carrier%k_locals)
NULLIFY (carrier%k_sizes)
IF (.NOT. local_indexing) THEN
ALLOCATE (carrier%k_locals(right%nblkrows_total))
NULLIFY (multrec_carrier%k_locals)
NULLIFY (multrec_carrier%k_sizes)
IF (.NOT. multrec_carrier%local_indexing) THEN
ALLOCATE (multrec_carrier%k_locals(right%nblkrows_total))
FORALL (i = 1:right%nblkrows_total)
carrier%k_locals(i) = i
multrec_carrier%k_locals(i) = i
END FORALL
carrier%k_sizes => array_data (right%row_blk_size)
multrec_carrier%k_sizes => array_data (right%row_blk_size)
ENDIF
CALL csr_multiply_outer_init(left, right, product, carrier, fill_guess, error=error)
CALL dbcsr_mm_csr_outer_init(csr_carrier=multrec_carrier%csr_carrier, &
left=left, right=right, product=product,fill_guess=fill_guess,&
m_sizes=multrec_carrier%m_sizes, n_sizes=multrec_carrier%n_sizes,&
k_sizes=multrec_carrier%k_sizes, error=error)
CALL dbcsr_error_stop(error_handler, error)
END SUBROUTINE dbcsr_multrec_init
END SUBROUTINE dbcsr_mm_multrec_init
! *****************************************************************************
@ -282,121 +288,85 @@ CONTAINS
!> \param[out] flop (optional) number of effective double-precision
!> floating point operations performed
! *****************************************************************************
SUBROUTINE dbcsr_multrec_sparse(left, right, product, carrier, flop,&
keep_sparsity, a_norms, b_norms, eps, row_max_epss, error)
SUBROUTINE dbcsr_mm_multrec_multiply(left, right, product, multrec_carrier, flop,&
keep_sparsity, a_norms, b_norms, error)
TYPE(dbcsr_type), INTENT(IN) :: left, right
TYPE(dbcsr_type), INTENT(INOUT) :: product
TYPE(carrier_type), INTENT(inout) :: carrier
TYPE(multrec_carrier_type), &
INTENT(inout) :: multrec_carrier
INTEGER(KIND=int_8), INTENT(OUT) :: flop
LOGICAL, INTENT(IN) :: keep_sparsity
REAL(kind=sp), DIMENSION(:), &
INTENT(in), TARGET :: a_norms, b_norms
REAL(kind=real_8), INTENT(in), OPTIONAL :: eps
REAL(kind=sp), DIMENSION(:), &
INTENT(in), TARGET :: row_max_epss
TYPE(dbcsr_error_type), INTENT(inout) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_multrec_sparse', &
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_multrec_multiply', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: dbg = .FALSE.
INTEGER :: ithread, nthreads, t_a_f, &
t_a_l, t_b_f, t_b_l
INTEGER, DIMENSION(:), POINTER :: k_locals
LOGICAL :: local_cols, local_indexing, &
local_rows
! ---------------------------------------------------------------------------
!
ithread = 0 ; nthreads = 1
!$ ithread = OMP_GET_THREAD_NUM () ; nthreads = OMP_GET_NUM_THREADS ()
carrier%flop = 0
CALL csr_multiply_inner_init(left, right, carrier, error=error)
local_cols = right%local_indexing
local_rows = left%local_indexing
local_indexing = local_rows
multrec_carrier%flop = 0
CALL dbcsr_mm_csr_inner_init(left=left, right=right, &
csr_carrier=multrec_carrier%csr_carrier, error=error)
! Find out the local A columns / B rows and sizes
! The right%local_rows is setup by the communication engine.
IF (local_indexing) THEN
IF (multrec_carrier%local_indexing) THEN
k_locals => array_data (right%local_rows)
carrier%k_locals => k_locals
CALL ensure_array_size (carrier%k_sizes, ub=SIZE(k_locals), error=error)
multrec_carrier%k_locals => k_locals
CALL ensure_array_size (multrec_carrier%k_sizes, ub=SIZE(k_locals), error=error)
CALL local_filter(array_data(right%row_blk_size), SIZE(k_locals),&
k_locals, carrier%k_sizes)
k_locals, multrec_carrier%k_sizes)
ELSE
k_locals => carrier%k_locals
k_locals => multrec_carrier%k_locals
ENDIF
! Setup the block norms
carrier%a_norms => a_norms
carrier%b_norms => b_norms
!
IF (mm_thread_workshare) THEN
!$OMP BARRIER
ENDIF
!
multrec_carrier%a_norms => a_norms
multrec_carrier%b_norms => b_norms
! Start local multiplication
IF (.TRUE.) THEN
t_a_f = 1
t_a_l = left%nblks
t_b_f = 1
t_b_l = right%nblks
!$ ithread = OMP_GET_THREAD_NUM()
!$ t_a_f = left%thr_c(ithread+1)+1
!$ t_a_l = left%thr_c(ithread+2)
IF (left%local_indexing) THEN
CALL sparse_multrec(&
1, left%nblkrows_local,&
1, right%nblkcols_local,&
1, SIZE(k_locals),&
t_a_f, t_a_l, left%coo_l,&
t_b_f, t_b_l, right%coo_l,&
carrier, error, 0)
ELSE
CALL sparse_multrec(&
1, left%nblkrows_total,&
1, right%nblkcols_total,&
1, SIZE(k_locals),&
t_a_f, t_a_l, left%coo_l,&
t_b_f, t_b_l, right%coo_l,&
carrier, error, 0)
ENDIF
t_a_f = 1
t_a_l = left%nblks
t_b_f = 1
t_b_l = right%nblks
!$ ithread = OMP_GET_THREAD_NUM()
!$ t_a_f = left%thr_c(ithread+1)+1
!$ t_a_l = left%thr_c(ithread+2)
IF (multrec_carrier%local_indexing) THEN
CALL sparse_multrec(&
1, left%nblkrows_local,&
1, right%nblkcols_local,&
1, SIZE(k_locals),&
t_a_f, t_a_l, left%coo_l,&
t_b_f, t_b_l, right%coo_l,&
product, multrec_carrier, error, 0)
ELSE
! For debugging. This avoids the recursion but it might not
! work anymore.
IF (left%local_indexing) THEN
CALL csr_multiply_unwrap(&
1, left%nblkrows_local,&
1, right%nblkrows_local,&
1, SIZE(k_locals),&
1, left%nblks, left%row_p,&
1, right%nblks, right%row_p,&
carrier, error)
ELSE
CALL csr_multiply_unwrap(&
1, left%nblkrows_total,&
1, right%nblkrows_total,&
1, SIZE(k_locals),&
1, left%nblks, left%row_p,&
1, right%nblks, right%row_p,&
carrier, error)
ENDIF
CALL sparse_multrec(&
1, left%nblkrows_total,&
1, right%nblkcols_total,&
1, SIZE(k_locals),&
t_a_f, t_a_l, left%coo_l,&
t_b_f, t_b_l, right%coo_l,&
product, multrec_carrier, error, 0)
ENDIF
CALL dbcsr_mm_csr_inner_finalize(csr_carrier=multrec_carrier%csr_carrier, error=error)
flop = multrec_carrier%flop
!
CALL csr_multiply_inner_finalize(carrier, error=error)
flop = carrier%flop
!
IF (ASSOCIATED (carrier%k_sizes) .AND. carrier%local_indexing) &
DEALLOCATE (carrier%k_sizes)
!
product%wms(ithread+1)%lastblk = carrier%lastblk
product%wms(ithread+1)%datasize = carrier%datasize
!
END SUBROUTINE dbcsr_multrec_sparse
IF (ASSOCIATED (multrec_carrier%k_sizes) .AND. multrec_carrier%local_indexing) &
DEALLOCATE (multrec_carrier%k_sizes)
END SUBROUTINE dbcsr_mm_multrec_multiply
! *****************************************************************************
!> \brief Sets up recursive multiplication
@ -408,12 +378,13 @@ CONTAINS
!> existing product matrix, default is no
!> \param[in] eps (optional) on-the-fly filtering epsilon
! *****************************************************************************
SUBROUTINE dbcsr_multrec_finalize(carrier, product, error)
TYPE(carrier_type), INTENT(inout) :: carrier
SUBROUTINE dbcsr_mm_multrec_finalize(multrec_carrier, product, error)
TYPE(multrec_carrier_type), &
INTENT(inout) :: multrec_carrier
TYPE(dbcsr_type), INTENT(inout) :: product
TYPE(dbcsr_error_type), INTENT(inout) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_multrec_finalize', &
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_mm_multrec_finalize', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: dbg = .FALSE.
@ -422,64 +393,52 @@ CONTAINS
! ---------------------------------------------------------------------------
CALL dbcsr_error_set(routineN, error_handler, error)
CALL csr_multiply_outer_finalize(carrier, product, error=error)
CALL dbcsr_mm_csr_outer_finalize(csr_carrier=multrec_carrier%csr_carrier,&
product=product, error=error)
ithread = 0
!$ ithread = OMP_GET_THREAD_NUM()
! Release the carrier
DEALLOCATE (carrier%m_sizes)
DEALLOCATE (carrier%n_sizes)
DEALLOCATE (carrier%row_max_epss)
IF (carrier%local_indexing) THEN
CALL remap_local2global(carrier%my_wm%row_i, carrier%my_wm%col_i,&
carrier%c_local_rows, carrier%c_local_cols,&
carrier%original_lastblk+1, carrier%lastblk)
ENDIF
!
DEALLOCATE (carrier%m_size_maps)
DEALLOCATE (carrier%n_size_maps)
DEALLOCATE (carrier%k_size_maps)
DEALLOCATE (carrier%stack_map)
IF (.NOT. carrier%local_indexing) THEN
DEALLOCATE (carrier%c_local_rows)
DEALLOCATE (carrier%c_local_cols)
DEALLOCATE (carrier%k_locals)
NULLIFY (carrier%k_sizes)
ENDIF
!
! Reinstate WM
carrier%my_wm%lastblk = carrier%lastblk
carrier%my_wm%datasize = carrier%datasize
product%wms(ithread+1) = carrier%my_wm
NULLIFY(carrier%right_data_sr)
NULLIFY(carrier%right_data_rr)
NULLIFY(carrier%right_index_sr)
NULLIFY(carrier%right_index_rr)
NULLIFY(carrier%left_data_sr)
NULLIFY(carrier%left_data_rr)
NULLIFY(carrier%left_index_sr)
NULLIFY(carrier%left_index_rr)
! Release the carrier
DEALLOCATE (multrec_carrier%m_sizes)
DEALLOCATE (multrec_carrier%n_sizes)
DEALLOCATE (multrec_carrier%row_max_epss)
IF (multrec_carrier%local_indexing) THEN
CALL remap_local2global(product%wms(ithread+1)%row_i, &
product%wms(ithread+1)%col_i, &
multrec_carrier%c_local_rows, multrec_carrier%c_local_cols,&
multrec_carrier%original_lastblk+1, product%wms(ithread+1)%lastblk)
ENDIF
!
IF (.NOT. multrec_carrier%local_indexing) THEN
DEALLOCATE (multrec_carrier%c_local_rows)
DEALLOCATE (multrec_carrier%c_local_cols)
DEALLOCATE (multrec_carrier%k_locals)
NULLIFY (multrec_carrier%k_sizes)
ENDIF
CALL dbcsr_error_stop(error_handler, error)
END SUBROUTINE dbcsr_multrec_finalize
END SUBROUTINE dbcsr_mm_multrec_finalize
! *****************************************************************************
!> \brief Performs recursive multiplication
!> \author Joost VandeVondele
! *****************************************************************************
RECURSIVE SUBROUTINE sparse_multrec(mi, mf, ni, nf, ki, kf,&
ai, af, a_index, bi, bf, b_index, &
carrier, error, d)
product, carrier, error, d)
INTEGER, INTENT(IN) :: mi, mf, ni, nf, ki, kf, ai, af
INTEGER, DIMENSION(3, 1:af), INTENT(IN) :: a_index
INTEGER, INTENT(IN) :: bi, bf
INTEGER, DIMENSION(3, 1:bf), INTENT(IN) :: b_index
TYPE(carrier_type), INTENT(INOUT) :: carrier
TYPE(dbcsr_type), INTENT(INOUT) :: product
TYPE(multrec_carrier_type), &
INTENT(INOUT) :: carrier
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
INTEGER, INTENT(IN) :: d
@ -502,13 +461,23 @@ CONTAINS
RETURN
ENDIF
ENDIF
IF (af-ai+1 <= norec .AND. bf-bi+1 <= norec) THEN
CALL csr_multiply_unwrap(&
mi, mf, ni, nf, ki, kf,&
ai, af, a_index,&
bi, bf, b_index,&
carrier, error)
RETURN
IF (af-ai+1 .GT. 0 .AND. bf-bi+1 .GT. 0) &
CALL dbcsr_mm_csr_multiply(mi=mi, mf=mf,ni=ni, nf=nf, ki=ki, kf=kf,&
ai=ai, af=af, &
bi=bi, bf=bf,&
m_sizes=carrier%m_sizes, n_sizes=carrier%n_sizes, k_sizes=carrier%k_sizes,&
c_local_rows=carrier%c_local_rows, c_local_cols=carrier%c_local_cols,&
c_has_symmetry=carrier%c_has_symmetry, keep_sparsity=carrier%keep_sparsity,&
use_eps=carrier%use_eps, row_max_epss=carrier%row_max_epss,&
flop=carrier%flop,&
a_index=a_index, b_index=b_index,&
a_norms=carrier%a_norms, b_norms=carrier%b_norms,&
csr_carrier=carrier%csr_carrier, &
product=product,&
error=error)
RETURN
ENDIF
M = mf-mi + 1
@ -525,9 +494,9 @@ CONTAINS
s1=M/2
acut = find_cut_row(ai,af,a_index,mi+s1-1)
CALL sparse_multrec(mi,mi+s1-1, ni,nf, ki,kf,&
ai,acut-1,a_index, bi,bf,b_index, carrier, error,d+1)
ai,acut-1,a_index, bi,bf,b_index, product, carrier, error,d+1)
CALL sparse_multrec(mi+s1,mf, ni,nf, ki,kf,&
acut,af,a_index, bi,bf,b_index, carrier, error,d+1)
acut,af,a_index, bi,bf,b_index, product, carrier, error,d+1)
CASE(2)
s1=K/2
acut = find_cut_col(ai,af,a_index,ki+s1-1)
@ -541,9 +510,9 @@ CONTAINS
WRITE(*,'(3(I7))')b_index
ENDIF
CALL sparse_multrec(mi,mf, ni,nf, ki,ki+s1-1,&
ai,acut-1,a_index, bi,bcut-1,b_index, carrier, error,d+1)
ai,acut-1,a_index, bi,bcut-1,b_index, product, carrier, error,d+1)
CALL sparse_multrec(mi,mf, ni,nf, ki+s1,kf,&
acut,af,a_index, bcut,bf,b_index, carrier, error,d+1)
acut,af,a_index, bcut,bf,b_index, product, carrier, error,d+1)
CASE(3)
s1=N/2
bcut = find_cut_col(bi,bf,b_index,ni+s1-1)
@ -552,15 +521,15 @@ CONTAINS
WRITE(*,'(3(I7))')b_index
ENDIF
CALL sparse_multrec(mi,mf, ni,ni+s1-1, ki,kf,&
ai,af,a_index, bi,bcut-1,b_index, carrier, error,d+1)
ai,af,a_index, bi,bcut-1,b_index, product, carrier, error,d+1)
CALL sparse_multrec(mi,mf, ni+s1,nf, ki,kf,&
ai,af,a_index, bcut,bf,b_index, carrier, error,d+1)
ai,af,a_index, bcut,bf,b_index, product, carrier, error,d+1)
END SELECT
END SUBROUTINE sparse_multrec
! ******************************************************************************
!> \brief
!> \brief
!> \author JV
! ******************************************************************************
PURE FUNCTION find_cut_row(ai,af,a,val) RESULT(res)
@ -601,13 +570,13 @@ CONTAINS
! IF (a(i)%r>val) EXIT
!ENDDO
!res=i
END FUNCTION find_cut_row
END FUNCTION find_cut_row
! ******************************************************************************
!> \brief
!> \brief
!> \author JV
! ******************************************************************************
! ******************************************************************************
PURE FUNCTION find_cut_col(ai,af,a,val) RESULT(res)
INTEGER, INTENT(IN) :: ai, af
INTEGER, DIMENSION(3, 1:af), INTENT(IN) :: a
@ -648,11 +617,11 @@ CONTAINS
!res=i
END FUNCTION find_cut_col
! ******************************************************************************
!> \brief Packs a globally-indexed array into a locally-indexed array.
! ******************************************************************************
! ******************************************************************************
PURE SUBROUTINE remap_local2global(row_i, col_i, local_rows, local_cols,&
first, last)
INTEGER, INTENT(in) :: last, first
@ -727,7 +696,7 @@ CONTAINS
local_data(l) = full_data(local_elements(l))
END FORALL
END SUBROUTINE local_filter_sp
END MODULE dbcsr_mm_multrec

View file

@ -13,43 +13,24 @@
! - 2011-09-26 Split dbcsr_internal_operations
! *****************************************************************************
MODULE dbcsr_mm_stack
USE dbcsr_block_operations, ONLY: dbcsr_data_clear
USE dbcsr_config, ONLY: &
driver_is_async, mm_async, mm_driver, mm_driver_blas, mm_driver_cuda, &
mm_driver_matmul, mm_driver_plasma, mm_driver_smm, mm_thread_workshare
USE dbcsr_cuda_device, ONLY: dbcsr_cuda_device_sync,&
dbcsr_cuda_stream_sync
USE dbcsr_cuda_memory, ONLY: dbcsr_cuda_dev_mem_alloc,&
dbcsr_cuda_dev_mem_dealloc,&
dbcsr_cuda_dev_mem_realloc,&
dbcsr_cuda_dev_mem_zero
USE dbcsr_cuda_methods, ONLY: dbcsr_cuda_dev_mem_get_alloc
USE dbcsr_config, ONLY: driver_is_async,&
mm_driver_blas,&
mm_driver_matmul,&
mm_driver_plasma,&
mm_driver_smm
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_ensure_size,&
dbcsr_data_get_size,&
dbcsr_data_set_size_referenced
USE dbcsr_data_methods, ONLY: dbcsr_data_get_size
USE dbcsr_error_handling
USE dbcsr_kinds, ONLY: dp,&
real_4,&
real_8,&
sp
USE dbcsr_pq_methods, ONLY: &
dbcsr_pq_add_stack, dbcsr_pq_flush_level_chg, dbcsr_pq_get_any_stack, &
dbcsr_pq_get_far_stack, dbcsr_pq_get_own_stack, &
dbcsr_pq_handoff_level_chg, dbcsr_ps_set_get_group_p, &
dbcsr_ps_set_get_n_working, dbcsr_ps_target_lock_main, &
dbcsr_ps_target_lock_regions, dbcsr_ps_target_new_regions, &
dbcsr_ps_target_unlock_main, dbcsr_ps_target_unlock_regions, &
dbcsr_psg_get_state, dbcsr_psg_set_state
USE dbcsr_pq_types, ONLY: &
dbcsr_pq_type, dbcsr_ps_group_type, dbcsr_ps_obj, dbcsr_ps_set_type, &
dbcsr_ps_state_empty, dbcsr_ps_state_queued, dbcsr_ps_state_working, &
dbcsr_ps_target_obj, dbcsr_ps_target_type, dbcsr_ps_type, &
dbcsr_ps_width, max_regions, p_a_first, p_b_first, p_c_first, p_k, &
p_m, p_n
USE dbcsr_mm_types, ONLY: &
dbcsr_ps_state_empty, dbcsr_ps_state_working, dbcsr_ps_width, &
p_a_first, p_b_first, p_c_first, p_k, p_m, p_n
USE dbcsr_toollib, ONLY: sort
USE dbcsr_types, ONLY: dbcsr_data_obj,&
dbcsr_type_complex_4,&
@ -76,548 +57,13 @@ MODULE dbcsr_mm_stack
LOGICAL, PARAMETER, PUBLIC :: show_levels = .FALSE.
LOGICAL, PARAMETER :: verbose_acc = .FALSE.
PUBLIC :: enqueue_ps_group
PUBLIC :: process_queue_mine,&
process_queue_others,&
process_queue_preempt,&
get_stack_or_process_queue
PUBLIC :: process_mm_stack_cpu, process_mm_stack_cuda
LOGICAL, PARAMETER :: debug_mod = .FALSE.
LOGICAL, PARAMETER :: careful_mod = .FALSE.
!> \var max_stack_block_size The maximal block size to be specially
!> treated.
INTEGER, PARAMETER :: max_stack_block_size = HUGE (INT (0))
!$ INTEGER(KIND=omp_lock_kind), SAVE :: accel_lock
CONTAINS
! *****************************************************************************
!> \brief Puts stacks ready to be processed into a queue
!>
!> All non-empty stacks in the param_group are added to the queue.
!> \par Non-worksharing
!> When mm_thread_workshare is FALSE the stacks are instead processessed
!> immediately, bypassing the queue.
!> \param[in,out] queue My queue, to which parameter stacks are added.
!> \param[in] param_group Parameter stacks to add to queue
!> \param[in,out] error error
!> \note The queue and param_groups parameters could be INTENT(INOUT)
!> when OpenMP is not used.
! *****************************************************************************
SUBROUTINE enqueue_ps_group (queue, param_group, error)
TYPE(dbcsr_pq_type), POINTER :: queue
TYPE(dbcsr_ps_group_type), POINTER :: param_group
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'enqueue_ps_group', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: careful = careful_mod, &
dbg = debug_mod
INTEGER :: error_handle, i
IF (careful_mod) &
CALL dbcsr_error_set(routineN, error_handle, error)
CALL dbcsr_psg_set_state (param_group, dbcsr_ps_state_queued, error)
param_group%master%s%driver = mm_driver
!$ if (dbg) write(*,*)" lckng targt", OMP_GET_THREAD_NUM(),&
!$ param_group%master%s%t%t%owner
CALL process_ps_target_low (param_group%master%s%t,&
param_group%master%s%t%t%product_data_area,&
param_group%master%s%t%t%zero_first,&
param_group%master%s%t%t%zero_last,&
param_group%master%s%t%t%last_c_blk,&
param_group%master%s%driver,&
param_group%master%s%t%t%product_data_cuda,&
param_group%master%s%t%t%c_locks_dev,&
error)
DO i = SIZE(param_group%stacks), 1, -1
param_group%stacks(i)%s%driver = mm_driver
IF (param_group%stacks(i)%s%stack_p .EQ. 0) THEN
param_group%stacks(i)%s%state = dbcsr_ps_state_empty
IF (dbg) WRITE(*,*)routineN//" not adding empty stack"
ELSE
IF (dbg) WRITE(*,*)routineN//" adding stack len", param_group%stacks(i)%s%stack_p
IF (mm_thread_workshare .OR. mm_async) THEN
CALL dbcsr_pq_add_stack (queue, param_group%stacks(i), error=error)
ELSE
! If no worksharing is active, just process the stack immediately.
param_group%stacks(i)%s%driver = mm_driver
CALL process_ps_stack (param_group%stacks(i)%s, error=error)
ENDIF
ENDIF
ENDDO
IF (careful_mod) &
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE enqueue_ps_group
! *****************************************************************************
!> \brief Processes stacks from my queue
!>
!> \param[in,out] queue My queue
!> \param[in] driver Driver to use for processing queue
!> \param[in,out] error error
!> \note The queue parameter could be INTENT(INOUT) when OpenMP is not used.
! *****************************************************************************
SUBROUTINE process_queue_mine (queue, driver, error)
TYPE(dbcsr_pq_type), POINTER :: queue
INTEGER, INTENT(IN) :: driver
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'process_queue_mine', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: careful = careful_mod, &
dbg = debug_mod
INTEGER :: cnt, error_handle, max_handoff
LOGICAL :: easy, found
TYPE(dbcsr_ps_obj) :: stack
IF (careful_mod) &
CALL dbcsr_error_set(routineN, error_handle, error)
easy = driver_is_async (driver)
cnt = 0
IF (easy) THEN
max_handoff = queue%handoff - queue%nworking
ENDIF
DO
CALL dbcsr_pq_get_own_stack (queue, stack, found,&
easy = easy, error=error)
IF (dbg) THEN
!$omp master
WRITE(*,*)"Found my stack"
!$omp end master
ENDIF
IF (found) THEN
stack%s%driver = driver
CALL process_ps_stack(stack%s, error=error)
cnt = cnt + 1
IF (easy .AND. cnt .GT. max_handoff) EXIT
ELSE
EXIT
ENDIF
ENDDO
!
IF (careful_mod) &
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE process_queue_mine
! *****************************************************************************
!> \brief Processes stacks from other threads
!>
!> \param[in,out] queue My queue
!> \param[in] forever Processes until all other threads declare
!> they are done.
!> \param[in] driver Driver to use for processing queue
!> \param[in,out] error error
!> \note The queue parameter could be INTENT(INOUT) when OpenMP is not used.
! *****************************************************************************
SUBROUTINE process_queue_others (queue, forever, driver, error)
TYPE(dbcsr_pq_type), POINTER :: queue
LOGICAL, INTENT(IN) :: forever
INTEGER, INTENT(IN) :: driver
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'process_queue_others', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: careful = careful_mod, &
dbg = debug_mod
INTEGER :: cnt, error_handle, nt, &
old_thread_ptr, t, thread_ptr
LOGICAL :: easy, found, keep_cycling
TYPE(dbcsr_ps_obj) :: stack
IF (careful_mod) &
CALL dbcsr_error_set(routineN, error_handle, error)
!
!
nt = 1
!$nt = omp_get_num_threads()
easy = driver_is_async (driver)
cnt = 1
old_thread_ptr = 0
thread_ptr = 1
found = .TRUE.
keep_cycling = found .OR.&
(forever .AND. queue%all_queues%n_working .GT. 0)
DO WHILE (keep_cycling)
CALL dbcsr_pq_get_any_stack (queue, stack, found,&
easy = easy,&
anyways = easy,&
thread_ptr = thread_ptr, error=error)
IF (thread_ptr .NE. old_thread_ptr) cnt = 1
old_thread_ptr = thread_ptr
IF (found) THEN
IF (dbg) THEN
!$omp master
WRITE(*,*)"Found other stack owned by", stack%s%t%t%owner
!$omp end master
ENDIF
stack%s%driver = driver
CALL process_ps_stack(stack%s, error=error)
cnt = cnt + 1
t = 0
!$ t = omp_get_thread_num ()
t = MOD(thread_ptr + t, nt)+1
IF (easy .AND.&
cnt .GT. queue%all_queues%queues(t)%pq%handoff&
- queue%all_queues%queues(t)%pq%nworking) THEN
thread_ptr = thread_ptr + 1
WRITE(*,*)'new thread', thread_ptr
ENDIF
ENDIF
!$OMP FLUSH
keep_cycling = found .OR.&
(forever .AND. queue%all_queues%n_working .GT. 0)
ENDDO
!
IF (careful_mod) &
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE process_queue_others
! *****************************************************************************
!> \brief Processes stacks while the accelerator processor is busy.
!>
!> Processing is not done if number of stacks being processed is less
!> than the queue's flush_level.
!> \see get_stack_or_process_queue
!> \see csr_multiply
!> \param[in,out] queue My queue
!> \param[in,out] param_sets Parameter group sets
!> \param[in] driver Driver to use for processing queue
!> \param[in,out] error error
!> \note The queue parameter could be INTENT(INOUT) when OpenMP is not used.
! *****************************************************************************
SUBROUTINE process_queue_preempt (queue, param_set, driver, error)
TYPE(dbcsr_pq_type), POINTER :: queue
TYPE(dbcsr_ps_set_type), INTENT(INOUT) :: param_set
INTEGER, INTENT(IN) :: driver
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'process_queue_preempt', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: careful = careful_mod, &
dbg = debug_mod
INTEGER :: error_handle, iam, n
LOGICAL :: found
TYPE(dbcsr_ps_obj) :: stack
IF (careful_mod) &
CALL dbcsr_error_set(routineN, error_handle, error)
!
! Fetch & process only if no one else is done (i.e., has spare
! time to process my stacks).
found = .TRUE.
n = dbcsr_ps_set_get_n_working (param_set, error)
!$ if (.FALSE. .AND. dbg) &
!$ write(*,*)routineN, omp_get_thread_num(), n, queue%flush_level
DO WHILE (n .GT. queue%flush_level .AND. found)
CALL dbcsr_pq_get_far_stack (queue, stack, found, error=error)
IF (found) THEN
IF (dbg) THEN
iam = 0
!$ iam = omp_get_thread_num ()
WRITE(*,*)"Preempt: Found own stack", iam
ENDIF
stack%s%driver = driver
CALL process_ps_stack(stack%s, error=error)
ENDIF
IF (dbg) THEN
n = dbcsr_ps_set_get_n_working (param_set, error)
IF (n .EQ. 0) WRITE(*,*)"oops"
ENDIF
ENDDO
!
IF (careful_mod) &
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE process_queue_preempt
! *****************************************************************************
!> \brief Gets a new stack group prepared for writing.
!> \par Dealing with busy stack groups.
!> The queue will be processed while the stack group is busy (it is
!> queued or in processing).
!>
!> If the stack_group is not immediately available the queue's
!> flush_level is lowered. It thus becomes more favorable to have the
!> host CPU process stacks.
!> \see process_queue_preempt
!> \see csr_multiply
!> \param[in,out] param_sets The parameter group sets
!> \param[in,out] queue My queue
!> \param[out] stack_group Retrieved stack group
!> \param[in] driver Driver to use for processing queue
!> \param[in,out] error error
!> \note The queue parameter could be INTENT(INOUT) when OpenMP is not used.
! *****************************************************************************
SUBROUTINE get_stack_or_process_queue (param_sets, queue, stack_group,&
driver, error)
TYPE(dbcsr_ps_set_type), INTENT(INOUT) :: param_sets
TYPE(dbcsr_pq_type), POINTER :: queue
TYPE(dbcsr_ps_group_type), POINTER :: stack_group
INTEGER, INTENT(IN) :: driver
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
LOGICAL, PARAMETER :: dbg = debug_mod
INTEGER :: cnt, n, stack_state
LOGICAL :: found
TYPE(dbcsr_ps_obj) :: stack
stack_group => dbcsr_ps_set_get_group_p (param_sets, wait=.FALSE.,&
error=error)
IF (mm_thread_workshare) THEN
!$OMP FLUSH
ENDIF
! If the parameter stack we get is still busy then
! just do more crunching.
cnt = 1
stack_state = dbcsr_psg_get_state (stack_group, error)
IF (stack_state .GE. dbcsr_ps_state_queued .AND. mm_async) THEN
! The parameter stack is still busy. The others processors are
! too slow so we lower the threshold for local computation.
!
!$ if (dbg) write(*,*)"still in queue:", queue%nstacks, omp_get_thread_num()
n = dbcsr_ps_set_get_n_working (param_sets, error)
CALL dbcsr_pq_flush_level_chg (queue, -(n/4),&
param_sets)
CALL dbcsr_pq_handoff_level_chg (queue, -MAX((param_sets%group_size*3)/2,1),&
param_sets)
ENDIF
DO WHILE (stack_state .GE. dbcsr_ps_state_queued)
! Just process the queue until the parameter stack is clear.
CALL dbcsr_pq_get_far_stack (queue, stack, found, error=error)
IF (found) THEN
stack%s%driver = driver
CALL process_ps_stack(stack%s, error=error)
ENDIF
IF (mm_thread_workshare) THEN
!$OMP FLUSH
ENDIF
cnt = cnt + 1
stack_state = dbcsr_psg_get_state (stack_group, error)
ENDDO
END SUBROUTINE get_stack_or_process_queue
! *****************************************************************************
!> \brief Prepares the target for stack processing.
!> \note Resizes data area, zeros new data, resets device locks, ...
! *****************************************************************************
SUBROUTINE process_ps_target_low (target_desc,&
product_data_area,&
zero_first, zero_last, nblks, driver, card_data, c_locks_dev, error)
TYPE(dbcsr_ps_target_obj), INTENT(INOUT) :: target_desc
TYPE(dbcsr_data_obj), INTENT(INOUT) :: product_data_area
INTEGER, INTENT(INOUT) :: zero_first, zero_last
INTEGER, INTENT(IN) :: nblks, driver
TYPE(dbcsr_cuda_mem_type), POINTER :: card_data, c_locks_dev
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'process_ps_target_low', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: careful = careful_mod, &
dbg = debug_mod, &
info = .FALSE.
REAL, PARAMETER :: resize_factor = 1.618034
INTEGER :: c_size, istat, maxs, tmp_i
LOGICAL :: stream_is_synced
! ---------------------------------------------------------------------------
!
! Resize target data area if necessary.
! Here we want the actual allocation size.
maxs = dbcsr_data_get_size(product_data_area)
IF (zero_last .GT. maxs) THEN
! Resize the target data area if the stack references a target
! matrix data element outside of its current range. When
! appropriate a lock on the data area is taken.
IF (dbg) WRITE(*,*)"Target too small"
maxs = zero_last
IF (mm_thread_workshare) THEN
! Lock the target and wait for others to unlock their
! regions.
CALL dbcsr_ps_target_lock_main (target_desc%t,&
flush=.TRUE.,&
error=error)
ENDIF
IF (dbg) &
WRITE(*,*)routineN//" Resizing to", LOG(REAL(maxs))/LOG(10.0)
CALL dbcsr_data_ensure_size (product_data_area,&
maxs, factor=resize_factor, error=error)
IF (mm_thread_workshare) THEN
! Potentially resize locks. The number of region locks is
! adjusted to the allocated size of the new data area (to
! prevent changing them whenever the last addressed block is
! changed).
maxs = dbcsr_data_get_size(product_data_area)
CALL dbcsr_ps_target_new_regions (target_desc%t,&
maxs, error=error)
CALL dbcsr_ps_target_unlock_main (target_desc%t, error)
ENDIF
ENDIF
IF (dbg) WRITE(*,*)"Setting referenced size"
CALL dbcsr_data_set_size_referenced (product_data_area, zero_last)
IF (dbg) WRITE(*,*)"Done"
!
! Zero new blocks
IF (zero_last .GE. zero_first) THEN
IF (dbg) WRITE(*,*)"Zeroing"
CALL dbcsr_data_clear (product_data_area, lb=zero_first, ub=zero_last)
IF (dbg) WRITE(*,*)"Done"
ENDIF
!
! Cuda on-device resizing
cuda_process_target: IF (driver .EQ. mm_driver_cuda) THEN
stream_is_synced = .FALSE.
IF (dbg) WRITE(*,*)"Getting allocated size..."
!$ IF (dbg) WRITE(*,*)"thread, target", omp_get_thread_num(),&
!$ target_desc%t%owner
c_size = dbcsr_cuda_dev_mem_get_alloc(card_data)
IF (dbg) WRITE(*,*)"done", c_size, "vs", zero_last
! Resize & zero product data if too big.
IF (zero_last .GT. c_size) THEN
!
! All kernels in the stream must be finished before
! reallocation is performed.
CALL dbcsr_cuda_stream_sync (target_desc%t%owner+1, error=error)
stream_is_synced = .TRUE.
IF (dbg .OR. info) WRITE(*,*)routineN//" reallocating c_dev",&
c_size, zero_last
tmp_i = MAX(zero_last, INT(REAL(zero_last)*resize_factor) )
IF (verbose_acc) WRITE(*,*)routineN//" reallocating c_dev",&
tmp_i, c_size
!--- !$OMP CRITICAL (crit_cuda)
!$ IF (mm_thread_workshare) CALL OMP_SET_LOCK (accel_lock)
CALL dbcsr_cuda_dev_mem_realloc (card_data, tmp_i, stat=istat)
IF (istat /= 0) THEN
! If the resize, including pre-oversizing failed, try to
! resize it to the actual number of locks used.
IF (verbose_acc) WRITE(*,*)routineN//" Running out of memory"
CALL dbcsr_cuda_dev_mem_realloc (card_data, zero_last, stat=istat)
ENDIF
c_size = dbcsr_cuda_dev_mem_get_alloc(card_data)
IF (dbg) WRITE(*,*)routineN//" zeroing c_dev"
CALL dbcsr_cuda_dev_mem_zero(card_data,&
first=zero_first, last=c_size, error=error)
!$ IF (mm_thread_workshare) CALL OMP_UNSET_LOCK (accel_lock)
!--- !$OMP END CRITICAL (crit_cuda)
ELSE IF (dbg) THEN
WRITE(*,*)"not reallocating c_dev"
ENDIF
!
! Resize locks, which are equal to the new block count.
IF (dbg) WRITE(*,*)routineN//" lock size..."
IF (dbg) WRITE(*,*)routineN, dbcsr_cuda_dev_mem_get_alloc(c_locks_dev)
IF (nblks .GT. dbcsr_cuda_dev_mem_get_alloc(c_locks_dev)) THEN
!
! All kernels in the stream must be finished before
! reallocation is performed.
IF (.NOT. stream_is_synced) THEN
CALL dbcsr_cuda_stream_sync (target_desc%t%owner+1, error=error)
ENDIF
stream_is_synced = .TRUE.
IF (dbg) WRITE(*,*)routineN//" reallocing locks"
maxs = dbcsr_cuda_dev_mem_get_alloc(c_locks_dev)
IF (dbg) WRITE(*,*)routineN//" lock size", maxs
IF (dbg) WRITE(*,*)routineN//" reallocing locks"
IF (verbose_acc .OR. info) WRITE(*,*)routineN//" reallocating locks", maxs
IF (dbg) WRITE(*,*)routineN//" deallocating..."
!--- !$OMP CRITICAL (crit_cuda)
!$ IF (mm_thread_workshare) CALL OMP_SET_LOCK (accel_lock)
CALL dbcsr_cuda_dev_mem_dealloc(c_locks_dev, error=error)
tmp_i = INT(REAL(nblks*4,kind=dp)*default_resize_factor)
!WRITE(*,*)routineN//" reallocating locks", tmp_i, maxs
IF (dbg) WRITE(*,*)routineN//" allocating..."
CALL dbcsr_cuda_dev_mem_alloc(c_locks_dev, tmp_i, stat=istat)
IF (istat /= 0) THEN
IF (verbose_acc) WRITE(*,*)routineN//" Trying smaller allocation"
tmp_i = nblks
IF (dbg) WRITE(*,*)routineN//" failed, trying smaller..."
CALL dbcsr_cuda_dev_mem_alloc(c_locks_dev, tmp_i, error=error)
ENDIF
IF (dbg) WRITE(*,*)routineN//" zeroing..."
CALL dbcsr_cuda_dev_mem_zero(c_locks_dev,&
first=1, last=dbcsr_cuda_dev_mem_get_alloc(c_locks_dev),&
error=error)
!$ IF (mm_thread_workshare) CALL OMP_UNSET_LOCK (accel_lock)
!--- !$OMP END CRITICAL (crit_cuda)
!write(*,*)routineN//" done zeroing"
ENDIF
ENDIF cuda_process_target
IF (dbg) WRITE(*,*)routineN//" done"
zero_first = zero_last + 1
END SUBROUTINE process_ps_target_low
! *****************************************************************************
!> \brief Unpacks stack variables to call lower-level stack processing
!> drivers.
!*****************************************************************************
SUBROUTINE process_ps_stack(param_stack, error)
TYPE(dbcsr_ps_type), INTENT(INOUT) :: param_stack
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'process_ps_stack', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: careful = careful_mod
INTEGER :: error_handle, stack_size
IF (careful_mod) &
CALL dbcsr_error_set(routineN, error_handle, error)
IF (careful_mod) THEN
IF (param_stack%driver .EQ. mm_driver_cuda) THEN
CALL dbcsr_assert (param_stack%has_cuda_ab_data,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"A or B matrix data not present in stack.",&
__LINE__, error=error)
ELSE
CALL dbcsr_assert (param_stack%has_ab_data,&
dbcsr_fatal_level, dbcsr_wrong_args_error, routineN,&
"A or B matrix data not present in stack.",&
__LINE__, error=error)
ENDIF
ENDIF
!if (dbg) &
! write(*,*)routineN, OMP_GET_THREAD_NUM(), param_stack%stack_p,&
! param_stack%id, param_stack%sid, param_stack%t%t%owner
!$ IF (.FALSE.) &
!$ WRITE(*,*)"process_ps_stack stack size is", param_stack%stack_p,&
!$ omp_get_thread_num(), param_stack%t%t%owner, param_stack%driver,&
!$ param_stack%sid, param_stack%id
! This has to be reset before the stack may be marked as empty.
stack_size = param_stack%stack_p
param_stack%stack_p = 0
CALL process_mm_stack(param_stack%parameters,&
param_stack%driver,&
stack_size,&
param_stack%left_data_area, param_stack%right_data_area,&
param_stack%t%t%product_data_area,&
param_stack%t%t%product_data_cuda, param_stack%t%t%has_cuda_c_data,&
param_stack%left_data_cuda, param_stack%right_data_cuda,&
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,&
param_stack%t%t, &
error=error)
IF (careful_mod) &
CALL dbcsr_error_stop(error_handle, error)
END SUBROUTINE process_ps_stack
! *****************************************************************************
!> \brief Calls the various drivers that process the stack.
@ -632,15 +78,11 @@ CONTAINS
!> \param[in] lastblk Number of blocks in product
!> \param[in,out] product_data_area Data for results
! *****************************************************************************
SUBROUTINE process_mm_stack(params, driver,&
SUBROUTINE process_mm_stack_cpu(params, driver,&
stack_size, &
left_data_area, right_data_area, product_data_area,&
product_data_card, has_product_data_card,&
a_dev, b_dev,&
state, stack_state_dev,&
m, n, k, max_m, max_n, max_k, defined_mnk,&
c_locks_dev, params_dev, &
target_desc, &
state, m, n, k, max_m, max_n, max_k, defined_mnk,&
stack_id,&
error)
INTEGER, INTENT(IN) :: stack_size, driver
INTEGER, DIMENSION(1:dbcsr_ps_width, &
@ -648,30 +90,21 @@ CONTAINS
TYPE(dbcsr_data_obj), INTENT(IN) :: left_data_area, &
right_data_area
TYPE(dbcsr_data_obj), INTENT(INOUT) :: product_data_area
TYPE(dbcsr_cuda_mem_type), POINTER :: product_data_card
LOGICAL, INTENT(IN) :: has_product_data_card
TYPE(dbcsr_cuda_mem_type), POINTER :: a_dev, b_dev
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_ps_target_type), POINTER :: target_desc
INTEGER, INTENT(IN) :: stack_id
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'process_mm_stack', &
CHARACTER(len=*), PARAMETER :: routineN = 'process_mm_stack_cpu', &
routineP = moduleN//':'//routineN
LOGICAL, PARAMETER :: careful = careful_mod, &
dbg = .FALSE.
REAL, PARAMETER :: resize_factor = 1.618034
INTEGER :: nregions, sp
LOGICAL, DIMENSION(max_regions) :: which_locks
INTEGER :: sp
REAL(KIND=dp) :: rnd
!stack_size), INTENT(IN) :: params
! ---------------------------------------------------------------------------
state = dbcsr_ps_state_working
IF (dbg) THEN
CALL RANDOM_NUMBER (rnd)
@ -680,7 +113,7 @@ CONTAINS
CALL print_gemm_parameters(params(:,1:stack_size))
ENDIF
ENDIF
!
! Verify stack consistency. Only the upper bound is verified.
IF (careful) THEN
DO sp = 1, stack_size
@ -701,21 +134,7 @@ CONTAINS
"C data out of bounds.", __LINE__, error=error)
ENDDO
ENDIF
!
! Locks the target (data) area if there could be an access
! conflict there.
!
!$ IF (mm_thread_workshare .AND. .NOT. driver_is_async (driver)) THEN
!$ CALL dbcsr_ps_target_lock_regions (target_desc,&
!$ stack_size, params,&
!$ target_desc%bit_shift,&
!$ which_locks, nregions,&
!$ error=error)
!$ ENDIF
IF (mm_thread_workshare) THEN
!$OMP FLUSH
ENDIF
SELECT CASE (driver)
CASE (mm_driver_matmul)
SELECT CASE (product_data_area%d%data_type)
@ -821,43 +240,15 @@ CONTAINS
CALL dbcsr_assert (.FALSE., dbcsr_fatal_level, dbcsr_caller_error,&
routineN, "Invalid data type",__LINE__,error)
END SELECT
CASE (mm_driver_cuda)
IF (.NOT. has_product_data_card) &
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"No C data on card is specified.",&
__LINE__, error=error)
CALL cuda_process_mm_stack (params,&
stack_size,&
a_dev, b_dev, product_data_card,&
c_locks_dev,&
params_dev,&
m, n, k, max_m, max_n, max_k, defined_mnk,&
state, stack_state_dev,&
target_desc%owner,&
error=error)
CASE default
CALL dbcsr_assert (.FALSE., dbcsr_fatal_level, dbcsr_caller_error,&
routineN, "Invalid multiplication driver",__LINE__,error)
END SELECT
!
! These writes have to be ordered. Specifically the state should not
! be set before the stack has been used and the results written.
IF (mm_thread_workshare) THEN
!$OMP FLUSH
ENDIF
IF (.NOT. driver_is_async (driver)) THEN
state = dbcsr_ps_state_empty
ENDIF
IF (mm_thread_workshare) THEN
!$OMP FLUSH
ENDIF
!$ IF (mm_thread_workshare .AND. .NOT. driver_is_async (driver)) THEN
!$ CALL dbcsr_ps_target_unlock_regions (target_desc,&
!$ which_locks, nregions,&
!$ error=error)
!$ ENDIF
END SUBROUTINE process_mm_stack
END SUBROUTINE process_mm_stack_cpu
! *****************************************************************************
@ -870,13 +261,14 @@ CONTAINS
!> \param[in,out] c_data Product data
!> \param[in,out] error error
! *****************************************************************************
SUBROUTINE cuda_process_mm_stack(params,&
SUBROUTINE process_mm_stack_cuda(params,&
stack_size,&
data_a_dev, data_b_dev, data_c_dev,&
c_locks,&
params_dev,&
m, n, k, max_m, max_n, max_k, defined_mnk,&
state, stack_state_dev, owner, &
state, stack_state_dev,&
stack_id,&
error)
INTEGER, INTENT(IN) :: stack_size
INTEGER, &
@ -889,46 +281,45 @@ CONTAINS
LOGICAL, INTENT(IN) :: defined_mnk
INTEGER, POINTER :: state
TYPE(dbcsr_cuda_mem_type), INTENT(IN) :: stack_state_dev
INTEGER, INTENT(IN) :: owner
INTEGER, INTENT(IN) :: stack_id
TYPE(dbcsr_error_type), INTENT(INOUT) :: error
CHARACTER(len=*), PARAMETER :: routineN = 'cuda_process_mm_stack', &
CHARACTER(len=*), PARAMETER :: routineN = 'process_mm_stack_cuda', &
routineP = moduleN//':'//routineN
INTEGER, PARAMETER :: flops_to_sort = 2*13*13*13
INTEGER :: error_handle, error_handle2, &
i, k_max, m_max, n_max, stream
i, ithread, k_max, m_max, &
n_max, nthreads, stream
INTEGER, DIMENSION(:), POINTER :: params_p
INTEGER, DIMENSION(stack_size) :: c_sort, c_sort_ind
INTEGER, &
DIMENSION(dbcsr_ps_width, stack_size) :: params_sort, params_tmp
! sort for matrices 13x13 and larger, this parameter needs tuning
! ---------------------------------------------------------------------------
nthreads = 1; ithread = 0
!$ nthreads = OMP_GET_NUM_THREADS () ; ithread = OMP_GET_THREAD_NUM ()
IF (stack_size .EQ. 0) THEN
!$OMP FLUSH
state = dbcsr_ps_state_empty
!$OMP FLUSH
WRITE(*,*)routineN//" Stack size is 0"
!$ write(*,*)"from thread", omp_get_thread_num()
RETURN
ENDIF
stream = owner+1
!$ IF (.FALSE.) &
!$ WRITE(*,*)routineN//" Stack size is", stack_size, omp_get_thread_num()
! IF (careful_mod) &
CALL dbcsr_error_set (routineN, error_handle, error)
! ithread is 0-based, stack_id is 1-based, stream should be 1-based
stream = ithread + nthreads*(stack_id-1) + 1
CALL dbcsr_error_set (routineN, error_handle, error)
params_p => params
IF (m .GT. 0) THEN ; m_max = m ; ELSE ; m_max = -max_m ; ENDIF
IF (n .GT. 0) THEN ; n_max = n ; ELSE ; n_max = -max_n ; ENDIF
IF (k .GT. 0) THEN ; k_max = k ; ELSE ; k_max = -max_k ; ENDIF
!========================================
!===========================================================================
! sort the stack. Since this costs CPU time, only a good idea if the CPUs
! are not too busy, or GPU gain is very large
CALL dbcsr_error_set (routineN//"_sort", error_handle2, error)
IF (2*m*n*k>=flops_to_sort .AND. stack_size .GT. 1) THEN
!TODO: now we are allways sorting
!IF (2*m*n*k>=flops_to_sort .AND. stack_size .GT. 1) THEN
params_tmp = RESHAPE(params, (/dbcsr_ps_width, stack_size/))
! sort by the C-blocks
c_sort = params_tmp(6,:)
@ -937,34 +328,31 @@ CONTAINS
params_sort(:, i)=params_tmp(:, c_sort_ind(i))
END DO
params = RESHAPE(params_sort, (/dbcsr_ps_width*stack_size/))
ENDIF
!ENDIF
CALL dbcsr_error_stop (error_handle2, error)
!========================================
!===========================================================================
CALL dbcsr_cuda_cp_host_to_dev (params_p,&
params_dev%d_i, dbcsr_ps_width*stack_size,&
async=.TRUE., stream=stream,&
error=error)
! Schedule the resetting of the state of this stack as soon as the
! parameter upload finishes. N.B. This must be the last change to
! the stack: its stack_p must be set to 0 and its data must be
! used <em>before</em> the state is set to dbcsr_ps_state_empty.
CALL dbcsr_cuda_cp_dev_to_host (stack_state_dev%d_i, state,&
async=.TRUE., stream=stream, error=error)
!--- !$OMP CRITICAL (crit_cuda)
!$ IF (mm_thread_workshare) CALL OMP_SET_LOCK (accel_lock)
!TODO: dbcsr_ps_width is a constant, shouldn't need to pass it as argument
CALL dbcsr_cuda_do_mm_stack (params_dev%d_i, stack_size, dbcsr_ps_width,&
data_a_dev, data_b_dev, data_c_dev,&
c_locks%d_i, ABS(m_max), ABS(n_max), ABS(k_max), defined_mnk,&
stream=stream, error=error)
!$ IF (mm_thread_workshare) CALL OMP_UNSET_LOCK (accel_lock)
!--- !$OMP END CRITICAL (crit_cuda)
IF (.FALSE.) CALL dbcsr_cuda_device_sync (error=error)
! IF (careful_mod) &
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE cuda_process_mm_stack
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE process_mm_stack_cuda

View file

@ -18,9 +18,8 @@ MODULE dbcsr_mm_types
int_8,&
real_8,&
sp
USE dbcsr_pq_types, ONLY: dbcsr_pq_type,&
dbcsr_ps_set_type
USE dbcsr_types, ONLY: dbcsr_work_type
USE dbcsr_types, ONLY: dbcsr_data_obj,&
dbcsr_work_type
IMPLICIT NONE
@ -28,7 +27,23 @@ MODULE dbcsr_mm_types
CHARACTER(len=*), PARAMETER, PRIVATE :: moduleN = 'dbcsr_mm_types'
PUBLIC :: carrier_type, hash_table_type, ele_type
PUBLIC :: multrec_carrier_type, csr_carrier_type
PUBLIC :: driver_carrier_type, stack_descriptor_type
PUBLIC :: hash_table_type, ele_type
INTEGER, PARAMETER, PUBLIC :: dbcsr_ps_state_empty = 42424242
INTEGER, PARAMETER, PUBLIC :: dbcsr_ps_state_working = 1234567890
! Definitions for the members of a parameter stack.
INTEGER, PARAMETER, PUBLIC :: dbcsr_ps_width = 7
INTEGER, PARAMETER, PUBLIC :: p_m = 1
INTEGER, PARAMETER, PUBLIC :: p_n = 2
INTEGER, PARAMETER, PUBLIC :: p_k = 3
INTEGER, PARAMETER, PUBLIC :: p_a_first = 4
INTEGER, PARAMETER, PUBLIC :: p_b_first = 5
INTEGER, PARAMETER, PUBLIC :: p_c_first = 6
INTEGER, PARAMETER, PUBLIC :: p_c_blk = 7
! *****************************************************************************
!> \brief Types needed for the hashtable.
@ -37,7 +52,7 @@ MODULE dbcsr_mm_types
INTEGER :: c=0
INTEGER :: p=0
END TYPE ele_type
TYPE hash_table_type
TYPE(ele_type), DIMENSION(:), POINTER :: table
INTEGER :: nele=0
@ -45,15 +60,58 @@ MODULE dbcsr_mm_types
INTEGER :: prime=0
END TYPE hash_table_type
! *****************************************************************************
!>
! *****************************************************************************
TYPE stack_descriptor_type
INTEGER :: m, n, k, max_m ,max_n, max_k
INTEGER :: total_entries
LOGICAL :: defined_mnk
END TYPE stack_descriptor_type
! *****************************************************************************
!>
! *****************************************************************************
TYPE driver_carrier_type
! for cpu based drivers
TYPE(dbcsr_data_obj) :: left_data_area, right_data_area
TYPE(dbcsr_work_type), POINTER :: product_wm
! for the cuda driver
END TYPE driver_carrier_type
! *****************************************************************************
!>
! *****************************************************************************
TYPE csr_carrier_type
TYPE(hash_table_type), DIMENSION(:), POINTER :: c_hashes
INTEGER :: nm_stacks, nn_stacks, nk_stacks
INTEGER(KIND=int_4), DIMENSION(:), POINTER :: m_size_maps,&
n_size_maps,&
k_size_maps
INTEGER :: max_m, max_n, max_k
INTEGER :: m_size_maps_size,&
n_size_maps_size,&
k_size_maps_size
INTEGER(KIND=int_1), DIMENSION(:,:,:), POINTER :: stack_map
INTEGER, DIMENSION(:,:,:,:), POINTER :: stackgrid_data
INTEGER, DIMENSION(:,:), POINTER :: stackgrid_fillcount
INTEGER, DIMENSION(:,:), POINTER :: stackgrid_state
TYPE(stack_descriptor_type), DIMENSION(:), POINTER :: stackrow_descr
TYPE(driver_carrier_type) :: drv_carrier
END TYPE csr_carrier_type
! *****************************************************************************
!> \brief Used to carry data among the various calls. Each thread has
!> its own private copy.
!> \var id Unique ID of each carrier_type instance
!> \var c_has_symmetry The product matrix has symmetry
!> \var keep_sparsity Sparsity of C matrix should be kept
!> \var use_eps Use on-the-fly filtering
!> \var param_sets Set of parameter stacks
!> \var local_indexing The A and B matrix indices and C matrix work indices
!> are not global but local to process rows and columns
!> \var m_sizes Block sizes of A and C matrix rows, indexed locally
@ -92,7 +150,6 @@ MODULE dbcsr_mm_types
!> \var nk_stacks The number of most common k blocks sizes.
!> \var stack_map Map from most common (n,k,m) block sizes to a stack
!> number within a stack group.
!> \var default_stack Stack to use when stack members do not have common
!> m, n, and k sizes.
!> \var row_max_epss Maximum eps to be used for one row.
!> \var a_norms Norms of A matrix blocks.
@ -105,38 +162,18 @@ MODULE dbcsr_mm_types
!> \var original_lastblk Number of work matrix blocks before addition
!> \var flop flop count
! *****************************************************************************
TYPE carrier_type
INTEGER :: id
TYPE multrec_carrier_type
LOGICAL :: c_has_symmetry, keep_sparsity, use_eps
LOGICAL :: local_indexing
TYPE(hash_table_type), DIMENSION(:), POINTER :: c_hashes
TYPE(dbcsr_pq_type), POINTER :: queue
TYPE(dbcsr_ps_set_type) :: param_sets
INTEGER, DIMENSION(:), POINTER :: m_sizes, n_sizes, k_sizes
INTEGER, DIMENSION(:), POINTER :: c_local_rows, c_local_cols, k_locals,&
c_global_rows, c_global_cols
INTEGER :: max_m, max_n, max_k
INTEGER(KIND=int_4), DIMENSION(:), POINTER :: m_size_maps,&
n_size_maps,&
k_size_maps
INTEGER :: m_size_maps_size,&
n_size_maps_size,&
k_size_maps_size
INTEGER :: nm_stacks, nn_stacks, nk_stacks
INTEGER(KIND=int_1), DIMENSION(:,:,:), POINTER :: stack_map
INTEGER :: default_stack
REAL(KIND=sp), DIMENSION(:), POINTER :: row_max_epss, a_norms, b_norms
REAL(KIND=real_8) :: eps
TYPE(dbcsr_work_type) :: my_wm
INTEGER :: lastblk, datasize
INTEGER :: original_lastblk
INTEGER(kind=int_8) :: flop
INTEGER, DIMENSION(:,:), POINTER :: right_data_sr, right_data_rr,&
left_data_sr, left_data_rr,&
right_index_sr, right_index_rr,&
left_index_sr, left_index_rr
END TYPE carrier_type
TYPE(csr_carrier_type) :: csr_carrier
END TYPE multrec_carrier_type
END MODULE dbcsr_mm_types

View file

@ -35,16 +35,11 @@ MODULE dbcsr_operations
dbcsr_data_clear,&
get_block2d_diagonal,&
set_block2d_diagonal
USE dbcsr_config, ONLY: &
dbcsr_get_conf_mm_stacksize, dbcsr_get_conf_nstacks, dbcsr_init_conf, &
has_cuda, is_configured, mm_async, 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
USE dbcsr_cuda_types, ONLY: dbcsr_cuda_mem_type
USE dbcsr_config, ONLY: dbcsr_init_conf,&
is_configured,&
mm_async,&
mm_driver,&
mm_driver_plasma
USE dbcsr_data_methods, ONLY: &
dbcsr_data_clear_pointer, dbcsr_data_ensure_size, dbcsr_data_get_size, &
dbcsr_data_get_type, dbcsr_data_init, dbcsr_data_new, &
@ -80,7 +75,6 @@ MODULE dbcsr_operations
USE dbcsr_message_passing, ONLY: dmp_max,&
mp_allgather,&
mp_bcast,&
mp_environ,&
mp_recv,&
mp_send,&
mp_sum
@ -105,13 +99,13 @@ MODULE dbcsr_operations
dbcsr_nblkrows_total, dbcsr_nfullcols_total, dbcsr_nfullrows_total, &
dbcsr_release, dbcsr_release_locals, dbcsr_row_block_offsets, &
dbcsr_row_block_sizes, dbcsr_switch_data_area, dbcsr_valid_index
USE dbcsr_mm_cannon, ONLY: dbcsr_mult_m_e_e
USE dbcsr_mm_cannon, ONLY: dbcsr_mm_cannon_lib_finalize,&
dbcsr_mm_cannon_lib_init,&
dbcsr_mm_cannon_multiply
USE dbcsr_mp_operations, ONLY: dbcsr_recv_any,&
dbcsr_send_any
USE dbcsr_plasma_interface, ONLY: dbcsr_plasma_finalize,&
dbcsr_plasma_init
USE dbcsr_pq_methods, ONLY: dbcsr_ps_set_finalize,&
dbcsr_ps_set_init
USE dbcsr_ptr_util, ONLY: ensure_array_size,&
pointer_view
USE dbcsr_toollib, ONLY: ceil_log2,&
@ -126,15 +120,14 @@ MODULE dbcsr_operations
USE dbcsr_types, ONLY: &
dbcsr_2d_array_type, dbcsr_conjugate_transpose, dbcsr_data_obj, &
dbcsr_distribution_obj, dbcsr_filter_frobenius, &
dbcsr_imagedistribution_obj, dbcsr_iterator, &
dbcsr_memory_CUDA_host_pinned, dbcsr_memory_MPI, dbcsr_memory_default, &
dbcsr_mp_obj, dbcsr_no_transpose, dbcsr_norm_column, &
dbcsr_norm_frobenius, dbcsr_norm_gershgorin, dbcsr_norm_maxabsnorm, &
dbcsr_obj, dbcsr_repl_col, dbcsr_repl_full, dbcsr_repl_none, &
dbcsr_repl_row, dbcsr_scalar_type, dbcsr_transpose, dbcsr_type, &
dbcsr_type_antisymmetric, dbcsr_type_complex_4, dbcsr_type_complex_8, &
dbcsr_type_no_symmetry, dbcsr_type_real_4, dbcsr_type_real_8, &
dbcsr_type_symmetric
dbcsr_imagedistribution_obj, dbcsr_iterator, dbcsr_memory_MPI, &
dbcsr_memory_default, dbcsr_mp_obj, dbcsr_no_transpose, &
dbcsr_norm_column, dbcsr_norm_frobenius, dbcsr_norm_gershgorin, &
dbcsr_norm_maxabsnorm, dbcsr_obj, dbcsr_repl_col, dbcsr_repl_full, &
dbcsr_repl_none, dbcsr_repl_row, dbcsr_scalar_type, dbcsr_transpose, &
dbcsr_type, dbcsr_type_antisymmetric, dbcsr_type_complex_4, &
dbcsr_type_complex_8, dbcsr_type_no_symmetry, dbcsr_type_real_4, &
dbcsr_type_real_8, dbcsr_type_symmetric
USE dbcsr_util, ONLY: dbcsr_checksum,&
dbcsr_verify_matrix,&
find_block_of_element
@ -142,8 +135,6 @@ MODULE dbcsr_operations
dbcsr_create,&
dbcsr_finalize,&
dbcsr_work_create
USE ma_affinity, ONLY: ma_set_gpu_affinity
USE machine_architecture_types, ONLY: has_ma
!$ USE OMP_LIB
@ -158,7 +149,7 @@ MODULE dbcsr_operations
PUBLIC :: dbcsr_init_lib, dbcsr_finalize_lib
PUBLIC :: dbcsr_multiply,&
dbcsr_mult_m_e_e, dbcsr_trace, dbcsr_add_on_diag,&
dbcsr_trace, dbcsr_add_on_diag,&
dbcsr_set, dbcsr_scale, dbcsr_scale_mat, dbcsr_add, dbcsr_copy,&
dbcsr_copy_submatrix, dbcsr_copy_into_existing,&
dbcsr_get_diag, dbcsr_set_diag, &
@ -378,7 +369,6 @@ CONTAINS
old_product_col_blk_sizes, old_product_row_blk_offsets, &
old_product_row_blk_sizes
TYPE(dbcsr_2d_array_type), POINTER :: m2s_left, m2s_right
TYPE(dbcsr_cuda_mem_type) :: a_tmp
TYPE(dbcsr_distribution_obj) :: dense_product_distribution, &
old_product_distribution
TYPE(dbcsr_imagedistribution_obj) :: dense_rdist_left, &
@ -388,19 +378,6 @@ CONTAINS
matrix_left, matrix_right, matrix_tmp, product_matrix
TYPE(dbcsr_scalar_type) :: eps_any
! ---------------------------------------------------------------------------
! Pre-init card for nicer time acconting.
IF (has_cuda) THEN
IF(mm_driver .NE. mm_driver_cuda) CALL dbcsr_ps_set_finalize(error=error)
ENDIF
IF (mm_driver .EQ. mm_driver_cuda) THEN
CALL dbcsr_cuda_dev_mem_setup (a_tmp, dbcsr_type_real_8)
CALL dbcsr_cuda_dev_mem_alloc(a_tmp, 1, error=error)
CALL dbcsr_cuda_dev_mem_dealloc(a_tmp, error=error)
ENDIF
!
CALL dbcsr_error_set(routineN, error_handler, error)
!
! check parameters
@ -963,7 +940,7 @@ CONTAINS
CALL dbcsr_access_flush (product_matrix, error=error)
!
my_flop = 0
CALL dbcsr_mult_m_e_e (m2s_left, m2s_right, product_matrix,&
CALL dbcsr_mm_cannon_multiply(m2s_left, m2s_right, product_matrix,&
retain_sparsity=retain_sparsity,&
filter_eps=filter_eps, error=error,&
flop=my_flop)
@ -4478,12 +4455,12 @@ CONTAINS
CHARACTER(len=*), PARAMETER :: routineN = 'dbcsr_init_lib', &
routineP = moduleN//':'//routineN
INTEGER :: error_handle, group_size, mem_type, mynode, n_stack_buffers, &
n_stack_mem_regions, ngpus, nstacks, numnode, stack_size
INTEGER, DIMENSION(3) :: nxstacks
TYPE(dbcsr_error_type) :: t_error
INTEGER :: error_handle
!n_stack_buffers,mem_type, n_stack_mem_regions, stack_size, nstacks,&
!INTEGER, DIMENSION(3) :: nxstacks
! ---------------------------------------------------------------------------
!TODO: problem: init/finalize are called by cp2k_runs AND f77_interface
IF (is_configured) RETURN
CALL dbcsr_error_set(routineN, error_handle, error)
@ -4504,45 +4481,13 @@ CONTAINS
dbcsr_fatal_level, dbcsr_internal_error, routineN,&
"Incorrect assumption of a 64-bit integer size!",&
__LINE__, error=error)
!
CALL dbcsr_init_conf (error)
!
CALL mp_environ (numnode, mynode, group)
#if defined (__DBCSR_CUDA)
IF (has_cuda) THEN
ngpus = dbcsr_cuda_get_n_devices (error)
mem_type = dbcsr_memory_CUDA_host_pinned
CALL dbcsr_get_conf_nstacks (nxstacks, n_stack_buffers, n_stack_mem_regions,&
error)
!$omp parallel default(none) shared(error)
CALL dbcsr_mm_cannon_lib_init(error)
!$omp end parallel
nstacks = nxstacks(1) * nxstacks(2) * nxstacks(3) + 1
stack_size = dbcsr_get_conf_mm_stacksize()*n_stack_buffers*nstacks
group_size = n_stack_buffers*nstacks
!$OMP PARALLEL default (none), &
!$OMP private (t_error) &
!$OMP shared (mynode, ngpus, error, &
!$OMP group_size, n_stack_mem_regions,stack_size)
! use a critical section to allow for threads initializing static
! variables without races
!$OMP CRITICAL
t_error = error
IF (has_ma) THEN
CALL dbcsr_cuda_init (card_num=ma_set_gpu_affinity(mynode), error=error)
ELSE
CALL dbcsr_cuda_init (card_num=MOD(mynode, ngpus), error=error)
ENDIF
!$OMP END CRITICAL
!allocate host memory
CALL dbcsr_ps_set_init(group_size, n_stack_mem_regions,stack_size,error=error)
!$OMP END PARALLEL
ENDIF
#endif
!
is_configured = .TRUE.
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dbcsr_init_lib
@ -4562,12 +4507,20 @@ CONTAINS
INTEGER :: error_handle
CALL dbcsr_ps_set_finalize(error=error)
!TODO: problem: init/finalize are called by cp2k_runs AND f77_interface
IF (.NOT. is_configured) RETURN
CALL dbcsr_error_set(routineN, error_handle, error)
!$omp parallel default(none) shared(error)
CALL dbcsr_mm_cannon_lib_finalize(error)
!$omp end parallel
is_configured = .FALSE.
CALL dbcsr_error_stop (error_handle, error)
END SUBROUTINE dbcsr_finalize_lib
! *****************************************************************************
!> \brief Adds blocks to a matrix

View file

@ -302,7 +302,7 @@ CONTAINS
CASE (dbcsr_memory_MPI)
CALL mp_allocate(mem, n, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_alloc(mem, n, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_alloc(mem, n, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&
@ -344,7 +344,7 @@ CONTAINS
CASE (dbcsr_memory_MPI)
CALL mp_deallocate(mem, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_dealloc(mem, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_dealloc(mem, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&

View file

@ -200,7 +200,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_allocate(mem, n, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_alloc(mem, n, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_alloc(mem, n, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&
@ -240,7 +240,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_deallocate(mem, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_dealloc(mem, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_dealloc(mem, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&

View file

@ -200,7 +200,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_allocate(mem, n, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_alloc(mem, n, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_alloc(mem, n, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&
@ -240,7 +240,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_deallocate(mem, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_dealloc(mem, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_dealloc(mem, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&

View file

@ -200,7 +200,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_allocate(mem, n, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_alloc(mem, n, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_alloc(mem, n, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&
@ -240,7 +240,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_deallocate(mem, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_dealloc(mem, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_dealloc(mem, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&

View file

@ -200,7 +200,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_allocate(mem, n, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_alloc(mem, n, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_alloc(mem, n, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&
@ -240,7 +240,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_deallocate(mem, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_dealloc(mem, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_dealloc(mem, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&

View file

@ -200,7 +200,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_allocate(mem, n, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_alloc(mem, n, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_alloc(mem, n, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&
@ -240,7 +240,7 @@
CASE (dbcsr_memory_MPI)
CALL mp_deallocate(mem, stat=stat)
CASE (dbcsr_memory_CUDA_host_pinned)
CALL dbcsr_cuda_host_mem_dealloc(mem, stat=stat, error=error)
CALL dbcsr_cuda_host_mem_dealloc(mem, error=error)
CASE default
CALL dbcsr_assert (.FALSE.,&
dbcsr_fatal_level, dbcsr_unimplemented_error_nr, routineN,&

View file

@ -28,7 +28,7 @@ MODULE input_cp2k
dbcsr_get_conf_use_comm_thread, has_blas, has_cuda, has_plasma, &
has_smm_gemm, 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, mm_thread_workshare
mm_name_matmul, mm_name_plasma, mm_name_smm
USE dbcsr_error_handling
USE dbcsr_types, ONLY: dbcsr_type_complex_4,&
dbcsr_type_complex_8,&
@ -802,7 +802,7 @@ CONTAINS
CHARACTER(len=64) :: a_blas, a_cuda, a_plasma, &
a_smm
INTEGER :: n_buffers, n_mem_regions
INTEGER :: n_buffers
INTEGER, DIMENSION(3) :: nstacks
LOGICAL :: failure
TYPE(dbcsr_error_type) :: dbcsr_error
@ -904,7 +904,7 @@ CONTAINS
CALL section_add_keyword(section,keyword,error=error)
CALL keyword_release(keyword,error=error)
!
CALL dbcsr_get_conf_nstacks (nstacks, n_buffers, n_mem_regions, dbcsr_error)
CALL dbcsr_get_conf_nstacks (nstacks, n_buffers, dbcsr_error)
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).",&
@ -919,36 +919,7 @@ CONTAINS
default_i_val=n_buffers,error=error)
CALL section_add_keyword(section,keyword,error=error)
CALL keyword_release(keyword,error=error)
CALL keyword_create(keyword, name="n_stack_memory_regions",&
description="Number of memory regions to use for stacks" &
// " (e.g., number of GPU cards per MPI process)",&
usage="n_stack_memory_regions 2",&
default_i_val=n_mem_regions,error=error)
CALL section_add_keyword(section,keyword,error=error)
CALL keyword_release(keyword,error=error)
!
CALL keyword_create(keyword, name="mm_thread_workshare",&
description="Share matrix multiplication work between multiple threads",&
usage="mm_thread_workshare T",&
variants=(/"MM_WORKSHARE"/),&
default_l_val=mm_thread_workshare,error=error)
CALL section_add_keyword(section,keyword,error=error)
CALL keyword_release(keyword,error=error)
!
CALL keyword_create(keyword, name="detailed_timing",&
description="Output multiplication timing (development)",&
usage="detailed_timing T",&
default_l_val=.FALSE.,error=error)
CALL section_add_keyword(section,keyword,error=error)
CALL keyword_release(keyword,error=error)
!
CALL keyword_create(keyword, name="kernel_timing",&
description="Time multiplication kernel (development)",&
usage="kernel_timing T",&
default_l_val=.FALSE.,error=error)
CALL section_add_keyword(section,keyword,error=error)
CALL keyword_release(keyword,error=error)
!
CALL keyword_create(keyword, name="use_comm_thread",&
description="During multiplication, use a thread to periodically poll" &
// " MPI to progress outstanding message completions. This is" &

View file

@ -8,7 +8,6 @@
&DBCSR
n_size_mnk_stacks 2
n_stack_buffers 4
n_stack_memory_regions 5
&END DBCSR
&END GLOBAL
&TEST