dtfft_interface_cuda Module

CUDA Driver Interfaces

CUDA Driver is loaded at runtime via dynamic loading.


Uses

  • module~~dtfft_interface_cuda~~UsesGraph module~dtfft_interface_cuda dtfft_interface_cuda iso_c_binding iso_c_binding module~dtfft_interface_cuda->iso_c_binding iso_fortran_env iso_fortran_env module~dtfft_interface_cuda->iso_fortran_env module~dtfft_errors dtfft_errors module~dtfft_interface_cuda->module~dtfft_errors module~dtfft_parameters dtfft_parameters module~dtfft_interface_cuda->module~dtfft_parameters module~dtfft_utils dtfft_utils module~dtfft_interface_cuda->module~dtfft_utils module~dtfft_errors->iso_fortran_env module~dtfft_parameters->iso_c_binding module~dtfft_parameters->iso_fortran_env mpi_f08 mpi_f08 module~dtfft_parameters->mpi_f08 module~dtfft_utils->iso_c_binding module~dtfft_utils->iso_fortran_env module~dtfft_utils->module~dtfft_errors module~dtfft_utils->module~dtfft_parameters module~dtfft_utils->mpi_f08

Used by

  • module~~dtfft_interface_cuda~~UsedByGraph module~dtfft_interface_cuda dtfft_interface_cuda module~dtfft_kernel_device dtfft_kernel_device module~dtfft_kernel_device->module~dtfft_interface_cuda module~dtfft_nvrtc_block_optimizer dtfft_nvrtc_block_optimizer module~dtfft_kernel_device->module~dtfft_nvrtc_block_optimizer module~dtfft_nvrtc_module_cache dtfft_nvrtc_module_cache module~dtfft_kernel_device->module~dtfft_nvrtc_module_cache module~dtfft_nvrtc_block_optimizer->module~dtfft_interface_cuda module~dtfft_nvrtc_module dtfft_nvrtc_module module~dtfft_nvrtc_module->module~dtfft_interface_cuda module~dtfft_nvrtc_module->module~dtfft_nvrtc_block_optimizer module~dtfft_nvrtc_module_cache->module~dtfft_interface_cuda module~dtfft_nvrtc_module_cache->module~dtfft_nvrtc_block_optimizer module~dtfft_nvrtc_module_cache->module~dtfft_nvrtc_module module~dtfft_transpose_plan dtfft_transpose_plan module~dtfft_transpose_plan->module~dtfft_interface_cuda module~dtfft_transpose_plan->module~dtfft_kernel_device module~dtfft_transpose_handle_generic dtfft_transpose_handle_generic module~dtfft_transpose_plan->module~dtfft_transpose_handle_generic module~dtfft_plan dtfft_plan module~dtfft_plan->module~dtfft_transpose_plan module~dtfft_transpose_handle_generic->module~dtfft_kernel_device module~dtfft dtfft module~dtfft->module~dtfft_plan module~dtfft_api dtfft_api module~dtfft_api->module~dtfft_plan

Variables

Type Visibility Attributes Name Initial
integer(kind=int32), public, parameter :: MAX_KERNEL_ARGS = 9
procedure(cuModuleLoadData_interface), public, pointer :: cuModuleLoadData

Fortran pointer to the cuModuleLoadData function
procedure(cuModuleUnload_interface), public, pointer :: cuModuleUnload

Fortran pointer to the cuModuleUnload function
procedure(cuModuleGetFunction_interface), public, pointer :: cuModuleGetFunction

Fortran pointer to the cuModuleGetFunction function
logical, private, save :: is_loaded = .false.

Flag indicating whether the library is loaded
type(c_ptr), private, save :: libcuda

Handle to the loaded library
type(c_funptr), private, save :: cuFunctions(4)

Array of pointers to the CUDA functions
procedure(cuLaunchKernel_interface), private, pointer :: cuLaunchKernel_

Fortran pointer to the cuLaunchKernel function

Abstract Interfaces

abstract interface

  • private function cuModuleLoadData_interface(mod, image) result(cuResult)

    Load a module’s data with options.

    Takes a pointer image and loads the corresponding module module into the current context. The image may be a cubin or fatbin as output by nvcc, or a NULL-terminated PTX, either as output by nvcc or hand-written.

    Arguments

    Type IntentOptional Attributes Name
    type(CUmodule) :: mod

    Returned module
    type(c_ptr), value :: image

    Module data to load

    Return Value integer(kind=c_int)

    Driver result code

abstract interface

  • private function cuModuleUnload_interface(hmod) result(cuResult)

    Unloads a module.

    Unloads a module hmod from the current context. Attempting to unload a module which was obtained from the Library Management API such as cuLibraryGetModule will return CUDA_ERROR_NOT_PERMITTED.

    Arguments

    Type IntentOptional Attributes Name
    type(CUmodule), value :: hmod

    Module to unload

    Return Value integer(kind=c_int)

    Driver result code

abstract interface

  • private function cuModuleGetFunction_interface(hfunc, hmod, name) result(cuResult)

    Returns a function handle.

    Returns in hfunc the handle of the function of name name located in module hmod. If no function of that name exists, cuModuleGetFunction returns CUDA_ERROR_NOT_FOUND.

    Arguments

    Type IntentOptional Attributes Name
    type(CUfunction) :: hfunc

    Returns a function handle.
    type(CUmodule), value :: hmod

    Module to retrieve function from
    type(c_ptr), value :: name

    Name of function to retrieve

    Return Value integer(kind=c_int)

    Driver result code

abstract interface

  • private function cuLaunchKernel_interface(func, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, stream, kernelParams, extra) result(cuResult)

    Launches a CUDA function CUfunction.

    Arguments

    Type IntentOptional Attributes Name
    type(CUfunction), value :: func

    CUDA function to launch
    integer(kind=c_int), value :: gridDimX

    Grid dimensions in X
    integer(kind=c_int), value :: gridDimY

    Grid dimensions in Y
    integer(kind=c_int), value :: gridDimZ

    Grid dimensions in Z
    integer(kind=c_int), value :: blockDimX

    Block dimensions in X
    integer(kind=c_int), value :: blockDimY

    Block dimensions in Y
    integer(kind=c_int), value :: blockDimZ

    Block dimensions in Z
    integer(kind=c_int), value :: sharedMemBytes

    Dynamic shared memory size
    type(dtfft_stream_t), value :: stream

    Stream identifier
    type(c_ptr) :: kernelParams(*)

    Array of pointers to kernel parameters
    type(c_ptr) :: extra

    Dynamic shared-memory size per thread block in bytes

    Return Value integer(kind=c_int)

    Driver result code


Derived Types

type, public, bind(C) ::  dim3

Dimension specification type

Components

Type Visibility Attributes Name Initial
integer(kind=c_int), public :: x
integer(kind=c_int), public :: y
integer(kind=c_int), public :: z

type, public, bind(C) ::  CUmodule

CUDA module

Components

Type Visibility Attributes Name Initial
type(c_ptr), public :: ptr

Actual pointer

type, public, bind(C) ::  CUfunction

CUDA function

Components

Type Visibility Attributes Name Initial
type(c_ptr), public :: ptr

Actual pointer

Functions

public function load_cuda() result(error_code)

Loads the CUDA Driver library and needed symbols

Arguments

None

Return Value integer(kind=int32)

Error code

public function cuLaunchKernel(func, in, out, blocks, threads, stream, nargs, args) result(cuResult)

Launches a CUDA kernel

Arguments

Type IntentOptional Attributes Name
type(CUfunction), intent(in) :: func

Function CUfunction or Kernel CUkernel to launch
type(c_ptr), intent(in), target :: in

Input pointer
type(c_ptr), intent(in), target :: out

Output pointer
type(dim3), intent(in) :: blocks

Grid in blocks
type(dim3), intent(in) :: threads

Thread block
type(dtfft_stream_t), intent(in) :: stream

Stream identifier
integer(kind=int32), intent(in) :: nargs
integer(kind=int32), intent(in), target :: args(MAX_KERNEL_ARGS)

Input parameters of kernel func

Return Value integer(kind=c_int)

Driver result code