21#if defined(HAVE_OPENCL) && defined(HAVE_CUDA)
22#error "Cannot compile with OpenCL and Cuda support at the same time"
25#if defined(HAVE_OPENCL) || defined(HAVE_CUDA)
34#if defined(HAVE_CLBLAS) || defined(HAVE_CLBLAST)
43 use iso_c_binding,
only: c_size_t
44 use,
intrinsic :: iso_fortran_env
98 integer,
public,
parameter :: &
99 ACCEL_MEM_READ_ONLY = cl_mem_read_only, &
103 integer,
public,
parameter :: &
104 ACCEL_MEM_READ_ONLY = 0, &
112 type(cl_context) :: cl_context
113#elif defined(HAVE_CUDA)
114 type(c_ptr) :: cuda_context
123 type(cl_device_id) :: cl_device
124#elif defined(HAVE_CUDA)
125 type(c_ptr) :: cuda_device
133 type(accel_context_t) :: context
134 type(accel_device_t) :: device
136 type(cl_command_queue) :: command_queue
138 type(c_ptr) :: cublas_handle
139 type(c_ptr) :: cuda_stream
140 type(c_ptr) :: module_map
141 integer :: max_workgroup_size
142 integer(int64) :: local_memory_size
143 integer(int64) :: global_memory_size
145 logical :: allow_CPU_only
146 logical :: shared_mem
149 logical :: initialize_buffers
150 character(len=32) :: debug_flag
151 integer(int64) :: max_block_dim(3)
152 integer(int64) :: max_grid_dim(3)
162 integer(c_size_t) :: size = 0
165 logical :: allocated = .false.
171 type(cl_kernel) :: kernel
174 type(c_ptr) :: cuda_kernel
175 type(c_ptr) :: cuda_module
176 type(c_ptr) :: arguments
178 integer(int64) :: cuda_shared_mem
179 logical :: initialized = .false.
180 type(accel_kernel_t),
pointer :: next
182 character(len=128) :: kernel_name
194 type(accel_kernel_t),
public,
target,
save :: kernel_copy
195 type(accel_kernel_t),
public,
target,
save :: kernel_copy_complex_to_real
196 type(accel_kernel_t),
public,
target,
save :: kernel_copy_real_to_complex
205 type(accel_kernel_t),
public,
target,
save :: kernel_phase
206 type(accel_kernel_t),
public,
target,
save :: kernel_phase_spiral
207 type(accel_kernel_t),
public,
target,
save :: dkernel_dot_matrix
216 type(accel_kernel_t),
public,
target,
save :: dzmul
217 type(accel_kernel_t),
public,
target,
save :: zzmul
293 integer,
parameter :: &
300 integer,
parameter :: &
310 integer :: buffer_alloc_count
311 integer(int64) :: allocated_mem
312 type(accel_kernel_t),
pointer :: head
313 type(alloc_cache_t) :: memcache
319 enabled =
accel%enabled
329 allow =
accel%allow_CPU_only
338 type(mpi_grp_t),
intent(inout) :: base_grp
339 type(namespace_t),
intent(in) :: namespace
341 logical :: disable, default, run_benchmark
342 integer :: idevice, iplatform
344 integer :: device_type
345 integer :: cl_status, idev
346 integer :: ndevices, ret_devices, nplatforms, iplat
347 character(len=256) :: device_name
348 type(cl_platform_id) :: platform_id
349 type(cl_program) :: prog
350 type(cl_platform_id),
allocatable :: allplatforms(:)
351 type(cl_device_id),
allocatable :: alldevices(:)
356 character(len=256) :: sys_name
362 buffer_alloc_count = 0
380 accel%enabled = .not. disable
383 if (
accel%enabled)
then
384 message(1) =
'Octopus was compiled without OpenCL or Cuda support.'
451 if (idevice<0) idevice = 0
453 idevice, base_grp%rank)
456 write(
message(1),
'(A,I5,A,I5,2A)')
"Rank ", base_grp%rank,
" uses device number ", idevice, &
457 " on ", trim(sys_name)
464 call cublas_init(
accel%cublas_handle,
accel%cuda_stream)
470 call clgetplatformids(nplatforms, cl_status)
473 safe_allocate(allplatforms(1:nplatforms))
475 call clgetplatformids(allplatforms, iplat, cl_status)
482 do iplat = 1, nplatforms
484 call clgetplatforminfo(allplatforms(iplat), cl_platform_name, device_name, cl_status)
486 if (iplatform < 0)
then
490 if (iplatform == iplat - 1)
then
498 call clgetplatforminfo(allplatforms(iplat), cl_platform_version, device_name, cl_status)
505 if (iplatform >= nplatforms .or. iplatform < 0)
then
507 if (iplatform > 0)
then
515 platform_id = allplatforms(iplatform + 1)
517 safe_deallocate_a(allplatforms)
519 call clgetdeviceids(platform_id, cl_device_type_all, ndevices, cl_status)
525 safe_allocate(alldevices(1:ndevices))
529 call clgetdeviceids(platform_id, cl_device_type_all, alldevices, ret_devices, cl_status)
531 do idev = 1, ndevices
534 call clgetdeviceinfo(alldevices(idev), cl_device_name, device_name, cl_status)
539 select case (idevice)
541 device_type = cl_device_type_gpu
543 device_type = cl_device_type_cpu
545 device_type = cl_device_type_accelerator
547 device_type = cl_device_type_default
549 device_type = cl_device_type_all
553 call clgetdeviceids(platform_id, device_type, alldevices, ret_devices, cl_status)
555 if (ret_devices < 1)
then
557 call clgetdeviceids(platform_id, cl_device_type_default, alldevices, ret_devices, cl_status)
559 if (ret_devices < 1)
then
561 call clgetdeviceids(platform_id, cl_device_type_all, alldevices, ret_devices, cl_status)
564 if (ret_devices < 1)
then
571 ndevices = ret_devices
573 if (idevice < 0)
then
574 if (base_grp%size > 1)
then
583 if (idevice >= ndevices)
then
584 call messages_write(
'Requested CL device does not exist (device = ')
592 accel%device%cl_device = alldevices(idevice + 1)
595 accel%context%cl_context = clcreatecontext(platform_id,
accel%device%cl_device, cl_status)
598 safe_deallocate_a(alldevices)
600 accel%command_queue = clcreatecommandqueue(
accel%context%cl_context,
accel%device%cl_device, &
601 cl_queue_profiling_enable, cl_status)
604 call clgetdeviceinfo(
accel%device%cl_device, cl_device_type, device_type, cl_status)
606 select case (device_type)
607 case (cl_device_type_gpu)
609 case (cl_device_type_cpu, cl_device_type_accelerator)
610 accel%shared_mem = .false.
612 accel%shared_mem = .false.
621 call clfftsetup(cl_status)
632 call clgetdeviceinfo(
accel%device%cl_device, cl_device_global_mem_size,
accel%global_memory_size, cl_status)
633 call clgetdeviceinfo(
accel%device%cl_device, cl_device_local_mem_size,
accel%local_memory_size, cl_status)
634 call clgetdeviceinfo(
accel%device%cl_device, cl_device_max_work_group_size,
accel%max_workgroup_size, cl_status)
643 accel%max_block_dim(1) = int(dim, int64)
645 accel%max_block_dim(2) = int(dim, int64)
647 accel%max_block_dim(3) = int(dim, int64)
649 accel%max_grid_dim(1) = int(dim, int64)
651 accel%max_grid_dim(2) = int(dim, int64)
653 accel%max_grid_dim(3) = int(dim, int64)
665 accel%debug_flag =
"-g"
666#elif defined(HAVE_CUDA)
667 accel%debug_flag =
"-lineinfo"
668#elif defined(HAVE_OPENCL)
669 accel%debug_flag =
"-g"
682 flags =
' -DRTYPE_DOUBLE')
684 flags =
'-DRTYPE_COMPLEX')
686 flags =
'-DRTYPE_DOUBLE')
688 flags =
'-DRTYPE_COMPLEX')
719 call parse_variable(namespace,
'AccelBenchmark', .false., run_benchmark)
723 if (run_benchmark)
then
745 if (
accel%cuda_mpi)
then
747 call messages_write(
"Warning: trying to use GPU-aware MPI, but we have not detected support in the linked MPI library.")
763#if defined (HAVE_ACCEL)
786#if defined(HAVE_OPENCL)
788 integer,
intent(inout) :: idevice
790 character(len=256) :: device_name
794 idevice = mod(base_grp%rank, ndevices)
796 call base_grp%barrier()
799 do irank = 0, base_grp%size - 1
800 if (irank == base_grp%rank)
then
801 call clgetdeviceinfo(alldevices(idevice + 1), cl_device_name, device_name, cl_status)
809 call base_grp%barrier()
818 integer(int64) :: val
823 integer :: major, minor
824 character(len=256) :: val_str
836#ifdef __HIP_PLATFORM_AMD__
846#ifdef __HIP_PLATFORM_AMD__
854 call clgetdeviceinfo(
accel%device%cl_device, cl_device_type, val, cl_status)
856 select case (int(val, int32))
857 case (cl_device_type_gpu)
859 case (cl_device_type_cpu)
861 case (cl_device_type_accelerator)
866 call clgetdeviceinfo(
accel%device%cl_device, cl_device_vendor, val_str, cl_status)
872 call clgetdeviceinfo(
accel%device%cl_device, cl_device_name, val_str, cl_status)
891 call clgetdeviceinfo(
accel%device%cl_device, cl_driver_version, val_str, cl_status)
903 call clgetdeviceinfo(
accel%device%cl_device, cl_device_max_compute_units, val, cl_status)
908 call clgetdeviceinfo(
accel%device%cl_device, cl_device_max_clock_frequency, val, cl_status)
925 call clgetdeviceinfo(
accel%device%cl_device, cl_device_max_mem_alloc_size, val, cl_status)
930 call clgetdeviceinfo(
accel%device%cl_device, cl_device_global_mem_cache_size, val, cl_status)
935 call clgetdeviceinfo(
accel%device%cl_device, cl_device_max_constant_buffer_size, val, cl_status)
972 character(len=*),
intent(in) :: platform_name
975 if (index(platform_name,
'AMD') > 0) platform_id =
cl_plat_amd
976 if (index(platform_name,
'ATI') > 0) platform_id =
cl_plat_ati
977 if (index(platform_name,
'NVIDIA') > 0) platform_id =
cl_plat_nvidia
978 if (index(platform_name,
'Intel') > 0) platform_id =
cl_plat_intel
989 integer(int64) :: hits, misses
990 real(real64) :: volume_hits, volume_misses
1000 if (.not. found)
exit
1003 call clreleasememobject(tmp%mem, ierr)
1011 call alloc_cache_end(memcache, hits, misses, volume_hits, volume_misses)
1022 if (hits + misses > 0)
then
1023 call messages_write(hits/real(hits + misses, real64)*100, fmt=
'(f6.1)', align_left = .
true.)
1029 if (volume_hits + volume_misses > 0)
then
1030 call messages_write(volume_hits/(volume_hits + volume_misses)*100, fmt=
'(f6.1)', align_left = .
true.)
1048 call clfftteardown()
1053 call cublas_end(
accel%cublas_handle)
1054 if (.not.
accel%cuda_mpi)
then
1060 call clreleasecommandqueue(
accel%command_queue, ierr)
1066 if (buffer_alloc_count /= 0)
then
1083 integer(int64),
intent(in) :: nn
1085 integer(int64) :: modnn, bsize
1094 modnn = mod(nn, bsize)
1095 if (modnn /= 0) psize = psize + bsize - modnn
1104 integer(int32),
intent(in) :: nn
1114 integer,
intent(in) :: flags
1115 type(
type_t),
intent(in) :: type
1116 integer,
intent(in) :: size
1117 logical,
optional,
intent(in) ::
set_zero
1126 integer,
intent(in) :: flags
1127 type(
type_t),
intent(in) :: type
1128 integer(int64),
intent(in) :: size
1129 logical,
optional,
intent(in) ::
set_zero
1131 integer(int64) :: fsize
1143 this%allocated = .
true.
1149 if (.not. found)
then
1151 this%mem = clcreatebuffer(
accel%context%cl_context, flags, fsize, ierr)
1159 buffer_alloc_count = buffer_alloc_count + 1
1160 allocated_mem = allocated_mem + fsize
1180 integer(int64) :: fsize
1184 if (this%size > 0)
then
1192 call clreleasememobject(this%mem, ierr)
1200 buffer_alloc_count = buffer_alloc_count - 1
1201 allocated_mem = allocated_mem + fsize
1208 this%allocated = .false.
1215 logical pure function accel_buffer_is_allocated(this) result(allocated)
1218 allocated = this%allocated
1232 call clfinish(
accel%command_queue, ierr)
1245 integer,
intent(in) :: narg
1256 call clsetkernelarg(kernel%kernel, narg, buffer%mem, ierr)
1270 integer,
intent(in) :: narg
1271 type(
type_t),
intent(in) :: type
1272 integer,
intent(in) :: size
1277 integer(int64) :: size_in_bytes
1284 if (size_in_bytes >
accel%local_memory_size)
then
1285 write(
message(1),
'(a,f12.6,a)')
"CL Error: requested local memory: ", real(size_in_bytes, real64) /1024.0,
" Kb"
1286 write(
message(2),
'(a,f12.6,a)')
" available local memory: ", real(
accel%local_memory_size, real64) /1024.0,
" Kb"
1288 else if (size_in_bytes <= 0)
then
1289 write(
message(1),
'(a,i10)')
"CL Error: invalid local memory size: ", size_in_bytes
1294 kernel%cuda_shared_mem = size_in_bytes
1298 call clsetkernelarglocal(kernel%kernel, narg, size_in_bytes, ierr)
1309 integer(int64),
intent(in) :: globalsizes(:)
1310 integer(int64),
intent(in) :: localsizes(:)
1316 integer(int64) :: gsizes(1:3)
1317 integer(int64) :: lsizes(1:3)
1325 dim = ubound(globalsizes, dim=1)
1327 assert(dim == ubound(localsizes, dim=1))
1330 if (any(globalsizes == 0))
return
1332 assert(all(localsizes > 0))
1334 assert(all(mod(globalsizes, localsizes) == 0))
1336 gsizes(1:dim) = globalsizes(1:dim)
1337 lsizes(1:dim) = localsizes(1:dim)
1340 call clenqueuendrangekernel(
accel%command_queue, kernel%kernel, gsizes(1:dim), lsizes(1:dim), ierr)
1346 if (any(lsizes(1:3) >
accel%max_block_dim(1:3)))
then
1347 message(1) =
"Maximum dimension of a block too large in kernel "//trim(kernel%kernel_name)
1348 message(2) =
"The following conditions should be fulfilled:"
1349 write(
message(3),
"(A, I8, A, I8)")
"Dim 1: ", lsizes(1),
" <= ",
accel%max_block_dim(1)
1350 write(
message(4),
"(A, I8, A, I8)")
"Dim 2: ", lsizes(2),
" <= ",
accel%max_block_dim(2)
1351 write(
message(5),
"(A, I8, A, I8)")
"Dim 3: ", lsizes(3),
" <= ",
accel%max_block_dim(3)
1352 message(6) =
"This is an internal error, please contact the developers."
1359 message(1) =
"Maximum number of threads per block too large in kernel "//trim(kernel%kernel_name)
1360 message(2) =
"The following condition should be fulfilled:"
1362 message(4) =
"This is an internal error, please contact the developers."
1366 gsizes(1:3) = gsizes(1:3)/lsizes(1:3)
1369 if (any(gsizes(1:3) >
accel%max_grid_dim(1:3)))
then
1370 message(1) =
"Maximum dimension of grid too large in kernel "//trim(kernel%kernel_name)
1371 message(2) =
"The following conditions should be fulfilled:"
1372 write(
message(3),
"(A, I8, A, I10)")
"Dim 1: ", gsizes(1),
" <= ",
accel%max_grid_dim(1)
1373 write(
message(4),
"(A, I8, A, I10)")
"Dim 2: ", gsizes(2),
" <= ",
accel%max_grid_dim(2)
1374 write(
message(5),
"(A, I8, A, I10)")
"Dim 3: ", gsizes(3),
" <= ",
accel%max_grid_dim(3)
1375 message(6) =
"This is an internal error, please contact the developers."
1379 call cuda_launch_kernel(kernel%cuda_kernel, gsizes(1), lsizes(1), kernel%cuda_shared_mem, kernel%arguments)
1381 kernel%cuda_shared_mem = 0
1390 integer,
intent(in) :: globalsizes(:)
1391 integer,
intent(in) :: localsizes(:)
1399 integer pure function accel_max_workgroup_size() result(max_workgroup_size)
1400 max_workgroup_size =
accel%max_workgroup_size
1409 integer(int64) :: workgroup_size8
1413 integer :: max_workgroup_size
1419 call clgetkernelworkgroupinfo(kernel%kernel,
accel%device%cl_device, cl_kernel_work_group_size, workgroup_size8, ierr)
1421 workgroup_size = workgroup_size8
1425 call cuda_kernel_max_threads_per_block(kernel%cuda_kernel, max_workgroup_size)
1426 if (debug%info .and. max_workgroup_size /=
accel%max_workgroup_size)
then
1427 write(message(1),
"(A, I5, A)")
"A kernel can use only less threads per block (", workgroup_size,
")", &
1428 "than available on the device (",
accel%max_workgroup_size,
")"
1429 call messages_info(1)
1433 workgroup_size = 256
1435 workgroup_size = min(workgroup_size, max_workgroup_size)
1444 type(cl_program),
intent(inout) :: prog
1445 character(len=*),
intent(in) :: filename
1446 character(len=*),
optional,
intent(in) :: flags
1448 character(len = 1000) :: string
1449 character(len = 256) :: share_string
1450 integer :: ierr, ierrlog, iunit, irec, newlen
1454 string =
'#include "'//trim(filename)//
'"'
1456 call messages_write(
"Building CL program '"//trim(filename)//
"'.")
1457 call messages_info(debug_only=.
true.)
1459 prog = clcreateprogramwithsource(
accel%context%cl_context, trim(string), ierr)
1465 string=trim(string)//
' -cl-denorms-are-zero'
1468 string=trim(string)//
' -cl-mad-enable'
1469 string=trim(string)//
' -cl-unsafe-math-optimizations'
1470 string=trim(string)//
' -cl-finite-math-only'
1471 string=trim(string)//
' -cl-fast-relaxed-math'
1473 share_string=
'-I'//trim(conf%share)//
'/opencl/'
1476 string = trim(string)//
' -DEXT_KHR_FP64'
1478 string = trim(string)//
' -DEXT_AMD_FP64'
1480 call messages_write(
'Octopus requires an OpenCL device with double-precision support.')
1481 call messages_fatal()
1485 string = trim(string)//
' -DSHARED_MEM'
1488 if (
present(flags))
then
1489 string = trim(string)//
' '//trim(flags)
1492 call messages_write(
"Debug info: compilation flags '"//trim(string), new_line = .
true.)
1493 call messages_write(
' '//trim(share_string)//
"'.")
1494 call messages_info(debug_only=.
true.)
1496 string = trim(string)//
' '//trim(share_string)
1498 call clbuildprogram(prog, trim(string), ierr)
1500 if(ierr /= cl_success)
then
1501 call clgetprogrambuildinfo(prog,
accel%device%cl_device, cl_program_build_log, string, ierrlog)
1505 newlen = scan(string, achar(010), back = .
true.) - 1
1506 if (newlen >= 0) string = string(1:newlen)
1508 if (len(trim(string)) > 0)
write(stderr,
'(a)') trim(string)
1520 type(cl_program),
intent(inout) :: prog
1526 call clreleaseprogram(prog, ierr)
1537 type(cl_kernel),
intent(inout) :: prog
1544 call clreleasekernel(prog, ierr)
1555 type(cl_kernel),
intent(inout) :: kernel
1556 type(cl_program),
intent(inout) :: prog
1557 character(len=*),
intent(in) :: name
1562 call profiling_in(
"CL_BUILD_KERNEL", exclude = .
true.)
1565 kernel = clcreatekernel(prog, name, ierr)
1569 call profiling_out(
"CL_BUILD_KERNEL")
1577 integer,
intent(in) :: ierr
1578 character(len=*),
intent(in) :: name
1580 character(len=40) :: errcode
1585 case (cl_success); errcode =
'CL_SUCCESS '
1586 case (cl_device_not_found); errcode =
'CL_DEVICE_NOT_FOUND '
1587 case (cl_device_not_available); errcode =
'CL_DEVICE_NOT_AVAILABLE '
1588 case (cl_compiler_not_available); errcode =
'CL_COMPILER_NOT_AVAILABLE '
1589 case (cl_mem_object_allocation_failure); errcode =
'CL_MEM_OBJECT_ALLOCATION_FAILURE '
1590 case (cl_out_of_resources); errcode =
'CL_OUT_OF_RESOURCES '
1591 case (cl_out_of_host_memory); errcode =
'CL_OUT_OF_HOST_MEMORY '
1592 case (cl_profiling_info_not_available); errcode =
'CL_PROFILING_INFO_NOT_AVAILABLE '
1593 case (cl_mem_copy_overlap); errcode =
'CL_MEM_COPY_OVERLAP '
1594 case (cl_image_format_mismatch); errcode =
'CL_IMAGE_FORMAT_MISMATCH '
1595 case (cl_image_format_not_supported); errcode =
'CL_IMAGE_FORMAT_NOT_SUPPORTED '
1596 case (cl_build_program_failure); errcode =
'CL_BUILD_PROGRAM_FAILURE '
1597 case (cl_map_failure); errcode =
'CL_MAP_FAILURE '
1598 case (cl_invalid_value); errcode =
'CL_INVALID_VALUE '
1599 case (cl_invalid_device_type); errcode =
'CL_INVALID_DEVICE_TYPE '
1600 case (cl_invalid_platform); errcode =
'CL_INVALID_PLATFORM '
1601 case (cl_invalid_device); errcode =
'CL_INVALID_DEVICE '
1602 case (cl_invalid_context); errcode =
'CL_INVALID_CONTEXT '
1603 case (cl_invalid_queue_properties); errcode =
'CL_INVALID_QUEUE_PROPERTIES '
1604 case (cl_invalid_command_queue); errcode =
'CL_INVALID_COMMAND_QUEUE '
1605 case (cl_invalid_host_ptr); errcode =
'CL_INVALID_HOST_PTR '
1606 case (cl_invalid_mem_object); errcode =
'CL_INVALID_MEM_OBJECT '
1607 case (cl_invalid_image_format_descriptor); errcode =
'CL_INVALID_IMAGE_FORMAT_DESCRIPTOR '
1608 case (cl_invalid_image_size); errcode =
'CL_INVALID_IMAGE_SIZE '
1609 case (cl_invalid_sampler); errcode =
'CL_INVALID_SAMPLER '
1610 case (cl_invalid_binary); errcode =
'CL_INVALID_BINARY '
1611 case (cl_invalid_build_options); errcode =
'CL_INVALID_BUILD_OPTIONS '
1612 case (cl_invalid_program); errcode =
'CL_INVALID_PROGRAM '
1613 case (cl_invalid_program_executable); errcode =
'CL_INVALID_PROGRAM_EXECUTABLE '
1614 case (cl_invalid_kernel_name); errcode =
'CL_INVALID_KERNEL_NAME '
1615 case (cl_invalid_kernel_definition); errcode =
'CL_INVALID_KERNEL_DEFINITION '
1616 case (cl_invalid_kernel); errcode =
'CL_INVALID_KERNEL '
1617 case (cl_invalid_arg_index); errcode =
'CL_INVALID_ARG_INDEX '
1618 case (cl_invalid_arg_value); errcode =
'CL_INVALID_ARG_VALUE '
1619 case (cl_invalid_arg_size); errcode =
'CL_INVALID_ARG_SIZE '
1620 case (cl_invalid_kernel_args); errcode =
'CL_INVALID_KERNEL_ARGS '
1621 case (cl_invalid_work_dimension); errcode =
'CL_INVALID_WORK_DIMENSION '
1622 case (cl_invalid_work_group_size); errcode =
'CL_INVALID_WORK_GROUP_SIZE '
1623 case (cl_invalid_work_item_size); errcode =
'CL_INVALID_WORK_ITEM_SIZE '
1624 case (cl_invalid_global_offset); errcode =
'CL_INVALID_GLOBAL_OFFSET '
1625 case (cl_invalid_event_wait_list); errcode =
'CL_INVALID_EVENT_WAIT_LIST '
1626 case (cl_invalid_event); errcode =
'CL_INVALID_EVENT '
1627 case (cl_invalid_operation); errcode =
'CL_INVALID_OPERATION '
1628 case (cl_invalid_gl_object); errcode =
'CL_INVALID_GL_OBJECT '
1629 case (cl_invalid_buffer_size); errcode =
'CL_INVALID_BUFFER_SIZE '
1630 case (cl_invalid_mip_level); errcode =
'CL_INVALID_MIP_LEVEL '
1631 case (cl_invalid_global_work_size); errcode =
'CL_INVALID_GLOBAL_WORK_SIZE '
1632 case (cl_platform_not_found_khr); errcode =
'CL_PLATFORM_NOT_FOUND_KHR'
1634 write(errcode,
'(i10)') ierr
1635 errcode =
'UNKNOWN ERROR CODE ('//trim(adjustl(errcode))//
')'
1638 message(1) =
'OpenCL '//trim(name)//
' '//trim(errcode)
1639 call messages_fatal(1)
1648 integer,
intent(in) :: ierr
1649 character(len=*),
intent(in) :: name
1651 character(len=40) :: errcode
1654#if defined(HAVE_CLBLAS) || defined(HAVE_CLBLAST)
1656 case (clblassuccess); errcode =
'clblasSuccess'
1657 case (clblasinvalidvalue); errcode =
'clblasInvalidValue'
1658 case (clblasinvalidcommandqueue); errcode =
'clblasInvalidCommandQueue'
1659 case (clblasinvalidcontext); errcode =
'clblasInvalidContext'
1660 case (clblasinvalidmemobject); errcode =
'clblasInvalidMemObject'
1661 case (clblasinvaliddevice); errcode =
'clblasInvalidDevice'
1662 case (clblasinvalideventwaitlist); errcode =
'clblasInvalidEventWaitList'
1663 case (clblasoutofresources); errcode =
'clblasOutOfResources'
1664 case (clblasoutofhostmemory); errcode =
'clblasOutOfHostMemory'
1665 case (clblasinvalidoperation); errcode =
'clblasInvalidOperation'
1666 case (clblascompilernotavailable); errcode =
'clblasCompilerNotAvailable'
1667 case (clblasbuildprogramfailure); errcode =
'clblasBuildProgramFailure'
1668 case (clblasnotimplemented); errcode =
'clblasNotImplemented'
1669 case (clblasnotinitialized); errcode =
'clblasNotInitialized'
1670 case (clblasinvalidmata); errcode =
'clblasInvalidMatA'
1671 case (clblasinvalidmatb); errcode =
'clblasInvalidMatB'
1672 case (clblasinvalidmatc); errcode =
'clblasInvalidMatC'
1673 case (clblasinvalidvecx); errcode =
'clblasInvalidVecX'
1674 case (clblasinvalidvecy); errcode =
'clblasInvalidVecY'
1675 case (clblasinvaliddim); errcode =
'clblasInvalidDim'
1676 case (clblasinvalidleaddima); errcode =
'clblasInvalidLeadDimA'
1677 case (clblasinvalidleaddimb); errcode =
'clblasInvalidLeadDimB'
1678 case (clblasinvalidleaddimc); errcode =
'clblasInvalidLeadDimC'
1679 case (clblasinvalidincx); errcode =
'clblasInvalidIncX'
1680 case (clblasinvalidincy); errcode =
'clblasInvalidIncY'
1681 case (clblasinsufficientmemmata); errcode =
'clblasInsufficientMemMatA'
1682 case (clblasinsufficientmemmatb); errcode =
'clblasInsufficientMemMatB'
1683 case (clblasinsufficientmemmatc); errcode =
'clblasInsufficientMemMatC'
1684 case (clblasinsufficientmemvecx); errcode =
'clblasInsufficientMemVecX'
1685 case (clblasinsufficientmemvecy); errcode =
'clblasInsufficientMemVecY'
1687 case (clblastinsufficientmemorytemp); errcode =
'clblastInsufficientMemoryTemp'
1688 case (clblastinvalidbatchcount); errcode =
'clblastInvalidBatchCount'
1689 case (clblastinvalidoverridekernel); errcode =
'clblastInvalidOverrideKernel'
1690 case (clblastmissingoverrideparameter); errcode =
'clblastMissingOverrideParameter'
1691 case (clblastinvalidlocalmemusage); errcode =
'clblastInvalidLocalMemUsage'
1692 case (clblastnohalfprecision); errcode =
'clblastNoHalfPrecision'
1693 case (clblastnodoubleprecision); errcode =
'clblastNoDoublePrecision'
1694 case (clblastinvalidvectorscalar); errcode =
'clblastInvalidVectorScalar'
1695 case (clblastinsufficientmemoryscalar); errcode =
'clblastInsufficientMemoryScalar'
1696 case (clblastdatabaseerror); errcode =
'clblastDatabaseError'
1697 case (clblastunknownerror); errcode =
'clblastUnknownError'
1698 case (clblastunexpectederror); errcode =
'clblastUnexpectedError'
1702 write(errcode,
'(i10)') ierr
1703 errcode =
'UNKNOWN ERROR CODE ('//trim(adjustl(errcode))//
')'
1707 message(1) =
'Error in calling clblas routine '//trim(name)//
' : '//trim(errcode)
1708 call messages_fatal(1)
1715 integer,
intent(in) :: ierr
1716 character(len=*),
intent(in) :: name
1718 character(len=40) :: errcode
1723 case (clfft_invalid_global_work_size); errcode =
'CLFFT_INVALID_GLOBAL_WORK_SIZE'
1724 case (clfft_invalid_mip_level); errcode =
'CLFFT_INVALID_MIP_LEVEL'
1725 case (clfft_invalid_buffer_size); errcode =
'CLFFT_INVALID_BUFFER_SIZE'
1726 case (clfft_invalid_gl_object); errcode =
'CLFFT_INVALID_GL_OBJECT'
1727 case (clfft_invalid_operation); errcode =
'CLFFT_INVALID_OPERATION'
1728 case (clfft_invalid_event); errcode =
'CLFFT_INVALID_EVENT'
1729 case (clfft_invalid_event_wait_list); errcode =
'CLFFT_INVALID_EVENT_WAIT_LIST'
1730 case (clfft_invalid_global_offset); errcode =
'CLFFT_INVALID_GLOBAL_OFFSET'
1731 case (clfft_invalid_work_item_size); errcode =
'CLFFT_INVALID_WORK_ITEM_SIZE'
1732 case (clfft_invalid_work_group_size); errcode =
'CLFFT_INVALID_WORK_GROUP_SIZE'
1733 case (clfft_invalid_work_dimension); errcode =
'CLFFT_INVALID_WORK_DIMENSION'
1734 case (clfft_invalid_kernel_args); errcode =
'CLFFT_INVALID_KERNEL_ARGS'
1735 case (clfft_invalid_arg_size); errcode =
'CLFFT_INVALID_ARG_SIZE'
1736 case (clfft_invalid_arg_value); errcode =
'CLFFT_INVALID_ARG_VALUE'
1737 case (clfft_invalid_arg_index); errcode =
'CLFFT_INVALID_ARG_INDEX'
1738 case (clfft_invalid_kernel); errcode =
'CLFFT_INVALID_KERNEL'
1739 case (clfft_invalid_kernel_definition); errcode =
'CLFFT_INVALID_KERNEL_DEFINITION'
1740 case (clfft_invalid_kernel_name); errcode =
'CLFFT_INVALID_KERNEL_NAME'
1741 case (clfft_invalid_program_executable); errcode =
'CLFFT_INVALID_PROGRAM_EXECUTABLE'
1742 case (clfft_invalid_program); errcode =
'CLFFT_INVALID_PROGRAM'
1743 case (clfft_invalid_build_options); errcode =
'CLFFT_INVALID_BUILD_OPTIONS'
1744 case (clfft_invalid_binary); errcode =
'CLFFT_INVALID_BINARY'
1745 case (clfft_invalid_sampler); errcode =
'CLFFT_INVALID_SAMPLER'
1746 case (clfft_invalid_image_size); errcode =
'CLFFT_INVALID_IMAGE_SIZE'
1747 case (clfft_invalid_image_format_descriptor); errcode =
'CLFFT_INVALID_IMAGE_FORMAT_DESCRIPTOR'
1748 case (clfft_invalid_mem_object); errcode =
'CLFFT_INVALID_MEM_OBJECT'
1749 case (clfft_invalid_host_ptr); errcode =
'CLFFT_INVALID_HOST_PTR'
1750 case (clfft_invalid_command_queue); errcode =
'CLFFT_INVALID_COMMAND_QUEUE'
1751 case (clfft_invalid_queue_properties); errcode =
'CLFFT_INVALID_QUEUE_PROPERTIES'
1752 case (clfft_invalid_context); errcode =
'CLFFT_INVALID_CONTEXT'
1753 case (clfft_invalid_device); errcode =
'CLFFT_INVALID_DEVICE'
1754 case (clfft_invalid_platform); errcode =
'CLFFT_INVALID_PLATFORM'
1755 case (clfft_invalid_device_type); errcode =
'CLFFT_INVALID_DEVICE_TYPE'
1756 case (clfft_invalid_value); errcode =
'CLFFT_INVALID_VALUE'
1757 case (clfft_map_failure); errcode =
'CLFFT_MAP_FAILURE'
1758 case (clfft_build_program_failure); errcode =
'CLFFT_BUILD_PROGRAM_FAILURE'
1759 case (clfft_image_format_not_supported); errcode =
'CLFFT_IMAGE_FORMAT_NOT_SUPPORTED'
1760 case (clfft_image_format_mismatch); errcode =
'CLFFT_IMAGE_FORMAT_MISMATCH'
1761 case (clfft_mem_copy_overlap); errcode =
'CLFFT_MEM_COPY_OVERLAP'
1762 case (clfft_profiling_info_not_available); errcode =
'CLFFT_PROFILING_INFO_NOT_AVAILABLE'
1763 case (clfft_out_of_host_memory); errcode =
'CLFFT_OUT_OF_HOST_MEMORY'
1764 case (clfft_out_of_resources); errcode =
'CLFFT_OUT_OF_RESOURCES'
1765 case (clfft_mem_object_allocation_failure); errcode =
'CLFFT_MEM_OBJECT_ALLOCATION_FAILURE'
1766 case (clfft_compiler_not_available); errcode =
'CLFFT_COMPILER_NOT_AVAILABLE'
1767 case (clfft_device_not_available); errcode =
'CLFFT_DEVICE_NOT_AVAILABLE'
1768 case (clfft_device_not_found); errcode =
'CLFFT_DEVICE_NOT_FOUND'
1769 case (clfft_success); errcode =
'CLFFT_SUCCESS'
1770 case (clfft_bugcheck); errcode =
'CLFFT_BUGCHECK'
1771 case (clfft_notimplemented); errcode =
'CLFFT_NOTIMPLEMENTED'
1772 case (clfft_file_not_found); errcode =
'CLFFT_FILE_NOT_FOUND'
1773 case (clfft_file_create_failure); errcode =
'CLFFT_FILE_CREATE_FAILURE'
1774 case (clfft_version_mismatch); errcode =
'CLFFT_VERSION_MISMATCH'
1775 case (clfft_invalid_plan); errcode =
'CLFFT_INVALID_PLAN'
1776 case (clfft_device_no_double); errcode =
'CLFFT_DEVICE_NO_DOUBLE'
1777 case (clfft_endstatus); errcode =
'CLFFT_ENDSTATUS'
1779 write(errcode,
'(i10)') ierr
1780 errcode =
'UNKNOWN ERROR CODE ('//trim(adjustl(errcode))//
')'
1784 message(1) =
'clfft '//trim(name)//
' '//trim(errcode)
1785 call messages_fatal(1)
1794 type(cl_device_id),
intent(inout) :: device
1795 character(len=*),
intent(in) :: extension
1798 character(len=2048) :: all_extensions
1801 call clgetdeviceinfo(device, cl_device_extensions, all_extensions,
cl_status)
1804 has = index(all_extensions, extension) /= 0
1813 type(type_t),
intent(in) :: type
1814 integer(int64),
intent(in) :: nval
1815 integer(int64),
optional,
intent(in) :: offset
1816 logical,
optional,
intent(in) :: async
1819 integer(int64) :: nval_real, offset_real
1822 if (nval == 0)
return
1827 if (
type == type_cmplx) nval_real = nval_real * 2
1828 if (
present(offset))
then
1829 offset_real = offset
1830 if (
type == type_cmplx) offset_real = offset_real * 2
1832 offset_real = 0_int64
1835 assert(nval_real > 0)
1837 if (
type == type_integer)
then
1850 if(.not. optional_default(async, .false.))
call accel_finish()
1859 type(type_t),
intent(in) :: type
1860 integer(int32),
intent(in) :: nval
1861 integer(int32),
optional,
intent(in) :: offset
1862 logical,
optional,
intent(in) :: async
1866 if (
present(offset))
then
1879 integer,
parameter :: times = 10
1881 real(real64) :: time, stime
1882 real(real64) :: read_bw, write_bw
1884 real(real64),
allocatable :: data(:)
1886 call messages_new_line()
1887 call messages_write(
'Info: Benchmarking the bandwidth between main memory and device memory')
1888 call messages_new_line()
1889 call messages_info()
1891 call messages_write(
' Buffer size Read bw Write bw')
1892 call messages_new_line()
1893 call messages_write(
' [MiB] [MiB/s] [MiB/s]')
1894 call messages_info()
1898 safe_allocate(
data(1:size))
1901 stime = loct_clock()
1906 time = (loct_clock() - stime)/real(times, real64)
1908 write_bw = real(
size, real64) *8.0_real64/time
1910 stime = loct_clock()
1916 time = (loct_clock() - stime)/real(times, real64)
1917 read_bw = real(
size, real64) *8.0_real64/time
1919 call messages_write(size*8.0_real64/1024.0_real64**2)
1920 call messages_write(write_bw/1024.0_real64**2, fmt =
'(f10.1)')
1921 call messages_write(read_bw/1024.0_real64**2, fmt =
'(f10.1)')
1922 call messages_info()
1926 safe_deallocate_a(data)
1928 size = int(size*2.0)
1930 if (
size > 50000000)
exit
1936 logical pure function accel_use_shared_mem() result(use_shared_mem)
1938 use_shared_mem =
accel%shared_mem
1950 call cuda_module_map_init(
accel%module_map)
1963 if (.not.
associated(
head))
exit
1964 next_head =>
head%next
1970 call cuda_module_map_end(
accel%module_map)
1980 character(len=*),
intent(in) :: file_name
1981 character(len=*),
intent(in) :: kernel_name
1982 character(len=*),
optional,
intent(in) :: flags
1985 type(cl_program) :: prog
1988 character(len=1000) :: all_flags
1993 call profiling_in(
"ACCEL_COMPILE", exclude = .
true.)
1996 all_flags =
'-I'//trim(conf%share)//
'/opencl/'//
" "//trim(
accel%debug_flag)
1999 all_flags = trim(all_flags)//
' -DSHARED_MEM'
2002 if (
present(flags))
then
2003 all_flags = trim(all_flags)//
' '//trim(flags)
2006 call cuda_build_program(
accel%module_map, this%cuda_module,
accel%device%cuda_device, trim(file_name), trim(all_flags))
2008 call cuda_create_kernel(this%cuda_kernel, this%cuda_module, trim(kernel_name))
2009 call cuda_alloc_arg_array(this%arguments)
2011 this%cuda_shared_mem = 0
2020 this%initialized = .
true.
2021 this%kernel_name = trim(kernel_name)
2023 call profiling_out(
"ACCEL_COMPILE")
2039 call cuda_free_arg_array(this%arguments)
2040 call cuda_release_kernel(this%cuda_kernel)
2045 call clreleasekernel(this%kernel, ierr)
2048 this%initialized = .false.
2057 character(len=*),
intent(in) :: file_name
2058 character(len=*),
intent(in) :: kernel_name
2059 character(len=*),
optional,
intent(in) :: flags
2063 if (.not. this%initialized)
then
2076 size =
accel%global_memory_size
2084 size =
accel%local_memory_size
2091 integer,
intent(in) :: dim
2097 if (dim == 1)
size = 2**30
2101 if (dim == 1)
size = 2**30
2108 integer,
intent(in) :: stream_number
2114 call cuda_set_stream(
accel%cuda_stream, stream_number)
2115 call cublas_set_stream(
accel%cublas_handle,
accel%cuda_stream)
2125 integer,
intent(inout) :: stream_number
2131 call cuda_get_stream(stream_number)
2145 call cuda_synchronize_all_streams()
2153 type(c_ptr),
intent(in) :: buffer
2154 integer(int64),
intent(in) :: offset
2155 type(c_ptr) :: buffer_offset
2159 call cuda_get_pointer_with_offset(buffer, offset, buffer_offset)
2162 buffer_offset = buffer
2168 type(c_ptr),
intent(in) :: buffer
2169 integer(int64),
intent(in) :: offset
2170 type(c_ptr) :: buffer_offset
2174 call cuda_get_pointer_with_offset(buffer, 2_int64*offset, buffer_offset)
2177 buffer_offset = buffer
2183 type(c_ptr),
intent(in) :: buffer
2187 call cuda_clean_pointer(buffer)
2196 integer(int64),
intent(in) :: size
2197 integer(int64),
intent(out) :: grid_size
2198 integer(int64),
intent(out) :: thread_block_size
2201#ifdef __HIP_PLATFORM_AMD__
2204 thread_block_size =
size
2206 grid_size =
size *
accel%warp_size
2207 thread_block_size =
accel%warp_size
2214#include "accel_inc.F90"
2217#include "complex.F90"
2218#include "accel_inc.F90"
2221#include "integer.F90"
2222#include "accel_inc.F90"
2225#include "integer8.F90"
2226#include "accel_inc.F90"
subroutine select_device(idevice)
subroutine laccel_get_device_pointer_3l(host_pointer, device_pointer, dimensions)
integer, parameter opencl_accelerator
type(accel_kernel_t), target, save, public kernel_density_real
subroutine zaccel_get_device_pointer_2l(host_pointer, device_pointer, dimensions)
integer, parameter opencl_default
type(accel_kernel_t), target, save, public kernel_vpsi_complex
type(accel_kernel_t), target, save, public dkernel_batch_axpy
subroutine, public accel_clean_pointer(buffer)
subroutine accel_kernel_global_end()
subroutine, public accel_get_unfolded_size(size, grid_size, thread_block_size)
Get unfolded size: some kernels (e.g. projectors) unfold the array across warps as an optimization....
pure logical function, public accel_allow_cpu_only()
subroutine iaccel_read_buffer_3(this, size, data, offset, async)
subroutine daccel_get_device_pointer_1(host_pointer, device_pointer, dimensions)
subroutine accel_create_buffer_8(this, flags, type, size, set_zero)
subroutine laccel_read_buffer_2(this, size, data, offset, async)
logical pure function, public accel_use_shared_mem()
subroutine laccel_read_buffer_3_int32(this, size, data, offset, async)
subroutine zaccel_get_device_pointer_1(host_pointer, device_pointer, dimensions)
type(accel_kernel_t), target, save, public zkernel_dot_matrix
subroutine daccel_write_buffer_0_int32(this, size, data, offset, async)
type(accel_kernel_t), target, save, public zpack
type(accel_kernel_t), target, save, public dkernel_batch_dotp
subroutine laccel_write_buffer_2(this, size, data, offset, async)
subroutine zaccel_write_buffer_3_int32(this, size, data, offset, async)
type(accel_kernel_t), target, save, public kernel_vpsi_spinors
subroutine laccel_get_device_pointer_1(host_pointer, device_pointer, dimensions)
subroutine daccel_write_buffer_0(this, size, data, offset, async)
subroutine zaccel_write_buffer_single(this, data, async)
type(accel_kernel_t), target, save, public kernel_ghost_reorder
subroutine iaccel_get_device_pointer_3l(host_pointer, device_pointer, dimensions)
subroutine laccel_get_device_pointer_2(host_pointer, device_pointer, dimensions)
type(accel_kernel_t), target, save, public zkernel_batch_axpy
integer, parameter cl_plat_nvidia
subroutine, public accel_kernel_start_call(this, file_name, kernel_name, flags)
subroutine iaccel_write_buffer_2_int32(this, size, data, offset, async)
subroutine daccel_read_buffer_3(this, size, data, offset, async)
integer, parameter cl_plat_ati
subroutine, public accel_get_stream(stream_number)
integer(int64) pure function, public accel_global_memory_size()
subroutine daccel_read_buffer_1_int32(this, size, data, offset, async)
subroutine iaccel_read_buffer_1(this, size, data, offset, async)
subroutine daccel_write_buffer_3_int32(this, size, data, offset, async)
subroutine zaccel_write_buffer_2(this, size, data, offset, async)
subroutine zaccel_write_buffer_1_int32(this, size, data, offset, async)
type(accel_kernel_t), target, save, public zkernel_dot_matrix_spinors
type(accel_kernel_t), target, save, public zkernel_ax_function_py
subroutine zaccel_set_kernel_arg_data(kernel, narg, data)
subroutine daccel_get_device_pointer_3l(host_pointer, device_pointer, dimensions)
subroutine laccel_read_buffer_0_int32(this, size, data, offset, async)
subroutine accel_set_kernel_arg_local(kernel, narg, type, size)
subroutine daccel_get_device_pointer_2l(host_pointer, device_pointer, dimensions)
subroutine zaccel_write_buffer_0_int32(this, size, data, offset, async)
integer(int64) function accel_padded_size_i8(nn)
subroutine iaccel_read_buffer_2_int32(this, size, data, offset, async)
type(accel_kernel_t), target, save set_zero_int
subroutine iaccel_get_device_pointer_1l(host_pointer, device_pointer, dimensions)
subroutine, public accel_finish()
subroutine opencl_check_bandwidth()
subroutine accel_kernel_global_init()
subroutine daccel_write_buffer_1(this, size, data, offset, async)
type(accel_kernel_t), target, save, public set_one
type(accel_kernel_t), target, save, public kernel_daxpy
subroutine laccel_read_buffer_0(this, size, data, offset, async)
subroutine opencl_release_program(prog)
type(accel_kernel_t), target, save set_zero
subroutine zaccel_read_buffer_3(this, size, data, offset, async)
subroutine iaccel_get_device_pointer_1(host_pointer, device_pointer, dimensions)
subroutine laccel_write_buffer_3(this, size, data, offset, async)
integer, parameter cl_plat_invalid
type(accel_kernel_t), target, save, public kernel_density_spinors
subroutine laccel_set_kernel_arg_data(kernel, narg, data)
subroutine daccel_read_buffer_2_int32(this, size, data, offset, async)
subroutine laccel_write_buffer_single(this, data, async)
subroutine daccel_write_buffer_2(this, size, data, offset, async)
subroutine daccel_get_device_pointer_2(host_pointer, device_pointer, dimensions)
subroutine accel_set_buffer_to_zero_i8(buffer, type, nval, offset, async)
subroutine zaccel_get_device_pointer_1l(host_pointer, device_pointer, dimensions)
logical pure function, public accel_buffer_is_allocated(this)
integer, parameter, public accel_mem_read_write
subroutine zaccel_read_buffer_0_int32(this, size, data, offset, async)
subroutine, public clfft_print_error(ierr, name)
subroutine accel_kernel_end(this)
subroutine iaccel_write_buffer_0_int32(this, size, data, offset, async)
integer, parameter opencl_gpu
type(accel_kernel_t), target, save, public dkernel_ax_function_py
subroutine iaccel_read_buffer_2(this, size, data, offset, async)
subroutine opencl_release_kernel(prog)
subroutine zaccel_get_device_pointer_2(host_pointer, device_pointer, dimensions)
subroutine zaccel_write_buffer_0(this, size, data, offset, async)
subroutine zaccel_read_buffer_1_int32(this, size, data, offset, async)
type(c_ptr) function, public daccel_get_pointer_with_offset(buffer, offset)
subroutine iaccel_write_buffer_single(this, data, async)
subroutine iaccel_get_device_pointer_2(host_pointer, device_pointer, dimensions)
integer pure function, public accel_max_size_per_dim(dim)
type(accel_kernel_t), target, save, public dpack
subroutine iaccel_write_buffer_3_int32(this, size, data, offset, async)
subroutine laccel_write_buffer_1(this, size, data, offset, async)
subroutine daccel_read_buffer_1(this, size, data, offset, async)
subroutine, public accel_release_buffer(this)
subroutine laccel_get_device_pointer_2l(host_pointer, device_pointer, dimensions)
subroutine iaccel_get_device_pointer_3(host_pointer, device_pointer, dimensions)
subroutine zaccel_write_buffer_1(this, size, data, offset, async)
subroutine daccel_set_kernel_arg_data(kernel, narg, data)
subroutine zaccel_read_buffer_2_int32(this, size, data, offset, async)
subroutine accel_kernel_run_8(kernel, globalsizes, localsizes)
subroutine laccel_write_buffer_1_int32(this, size, data, offset, async)
subroutine laccel_read_buffer_3(this, size, data, offset, async)
type(accel_kernel_t), target, save, public kernel_vpsi_spinors_complex
subroutine opencl_build_program(prog, filename, flags)
subroutine laccel_write_buffer_2_int32(this, size, data, offset, async)
subroutine, public accel_kernel_build(this, file_name, kernel_name, flags)
subroutine iaccel_write_buffer_2(this, size, data, offset, async)
subroutine, public accel_init(base_grp, namespace)
subroutine, public accel_end(namespace)
subroutine iaccel_read_buffer_0(this, size, data, offset, async)
subroutine opencl_create_kernel(kernel, prog, name)
subroutine, public accel_synchronize_all_streams()
subroutine iaccel_read_buffer_1_int32(this, size, data, offset, async)
subroutine, public accel_set_stream(stream_number)
type(accel_kernel_t), target, save, public zunpack
subroutine laccel_get_device_pointer_1l(host_pointer, device_pointer, dimensions)
subroutine zaccel_read_buffer_3_int32(this, size, data, offset, async)
subroutine laccel_read_buffer_2_int32(this, size, data, offset, async)
type(accel_kernel_t), target, save, public kernel_zaxpy
subroutine laccel_read_buffer_1_int32(this, size, data, offset, async)
integer, parameter cl_plat_amd
subroutine zaccel_read_buffer_2(this, size, data, offset, async)
integer(int32) function accel_padded_size_i4(nn)
subroutine accel_set_buffer_to_zero_i4(buffer, type, nval, offset, async)
subroutine daccel_write_buffer_2_int32(this, size, data, offset, async)
subroutine iaccel_write_buffer_1_int32(this, size, data, offset, async)
type(accel_kernel_t), target, save, public zkernel_batch_dotp
subroutine iaccel_get_device_pointer_2l(host_pointer, device_pointer, dimensions)
subroutine daccel_write_buffer_1_int32(this, size, data, offset, async)
subroutine iaccel_write_buffer_0(this, size, data, offset, async)
pure logical function, public accel_is_enabled()
subroutine zaccel_read_buffer_0(this, size, data, offset, async)
integer, parameter cl_plat_intel
integer, parameter, public accel_mem_write_only
subroutine daccel_read_buffer_3_int32(this, size, data, offset, async)
subroutine daccel_read_buffer_0_int32(this, size, data, offset, async)
subroutine laccel_write_buffer_0_int32(this, size, data, offset, async)
subroutine zaccel_read_buffer_1(this, size, data, offset, async)
type(accel_kernel_t), target, save, public kernel_vpsi
logical function f90_cl_device_has_extension(device, extension)
subroutine opencl_print_error(ierr, name)
subroutine daccel_get_device_pointer_1l(host_pointer, device_pointer, dimensions)
subroutine accel_kernel_run_4(kernel, globalsizes, localsizes)
subroutine zaccel_write_buffer_2_int32(this, size, data, offset, async)
subroutine iaccel_write_buffer_1(this, size, data, offset, async)
subroutine laccel_write_buffer_0(this, size, data, offset, async)
type(accel_kernel_t), target, save, public kernel_density_complex
subroutine zaccel_get_device_pointer_3(host_pointer, device_pointer, dimensions)
type(c_ptr) function, public zaccel_get_pointer_with_offset(buffer, offset)
subroutine daccel_write_buffer_single(this, data, async)
subroutine daccel_read_buffer_0(this, size, data, offset, async)
integer function, public accel_kernel_workgroup_size(kernel)
subroutine laccel_read_buffer_1(this, size, data, offset, async)
integer, parameter opencl_cpu
subroutine zaccel_write_buffer_3(this, size, data, offset, async)
integer function get_platform_id(platform_name)
subroutine, public clblas_print_error(ierr, name)
type(accel_t), public accel
subroutine laccel_get_device_pointer_3(host_pointer, device_pointer, dimensions)
subroutine accel_create_buffer_4(this, flags, type, size, set_zero)
subroutine laccel_write_buffer_3_int32(this, size, data, offset, async)
subroutine daccel_read_buffer_2(this, size, data, offset, async)
subroutine iaccel_set_kernel_arg_data(kernel, narg, data)
subroutine daccel_write_buffer_3(this, size, data, offset, async)
integer, public cl_status
subroutine zaccel_get_device_pointer_3l(host_pointer, device_pointer, dimensions)
subroutine iaccel_read_buffer_0_int32(this, size, data, offset, async)
type(accel_kernel_t), target, save, public dunpack
subroutine iaccel_write_buffer_3(this, size, data, offset, async)
integer(int64) pure function, public accel_local_memory_size()
subroutine accel_set_kernel_arg_buffer(kernel, narg, buffer)
integer pure function, public accel_max_workgroup_size()
subroutine daccel_get_device_pointer_3(host_pointer, device_pointer, dimensions)
type(accel_kernel_t), pointer head
subroutine iaccel_read_buffer_3_int32(this, size, data, offset, async)
subroutine, public alloc_cache_put(alloc_cache, size, loc, put)
subroutine, public alloc_cache_get(alloc_cache, size, found, loc)
integer(int64), parameter, public alloc_cache_any_size
integer, parameter, public clblassuccess
real(real64), parameter, public m_zero
subroutine, public messages_print_with_emphasis(msg, iunit, namespace)
character(len=512), private msg
subroutine, public messages_warning(no_lines, all_nodes, namespace)
subroutine, public messages_obsolete_variable(namespace, name, rep)
subroutine, public messages_new_line()
character(len=256), dimension(max_lines), public message
to be output by fatal, warning
subroutine, public messages_fatal(no_lines, only_root_writes, namespace)
subroutine, public messages_info(no_lines, iunit, debug_only, stress, all_nodes, namespace)
logical function mpi_grp_is_root(grp)
Is the current MPI process of grpcomm, root.
subroutine, public profiling_out(label)
Increment out counter and sum up difference between entry and exit time.
subroutine, public profiling_in(label, exclude)
Increment in counter and save entry time.
integer pure function, public types_get_size(this)
This module defines the unit system, used for input and output.
type(unit_t), public unit_gigabytes
For larger amounts of data (natural code units are bytes)
type(unit_t), public unit_megabytes
For large amounts of data (natural code units are bytes)
type(unit_t), public unit_kilobytes
For small amounts of data (natural code units are bytes)