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_interface_cuda_runtime dtfft_interface_cuda_runtime module~dtfft_interface_cuda->module~dtfft_interface_cuda_runtime 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_interface_cuda_runtime->iso_c_binding module~dtfft_interface_cuda_runtime->module~dtfft_parameters module~dtfft_interface_cuda_runtime->module~dtfft_utils 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_parameters module~dtfft_utils->mpi_f08

Used by

  • module~~dtfft_interface_cuda~~UsedByGraph module~dtfft_interface_cuda dtfft_interface_cuda module~dtfft_nvrtc_kernel dtfft_nvrtc_kernel module~dtfft_nvrtc_kernel->module~dtfft_interface_cuda module~dtfft_transpose_plan_cuda dtfft_transpose_plan_cuda module~dtfft_transpose_plan_cuda->module~dtfft_interface_cuda module~dtfft_transpose_plan_cuda->module~dtfft_nvrtc_kernel module~dtfft_abstract_backend dtfft_abstract_backend module~dtfft_transpose_plan_cuda->module~dtfft_abstract_backend module~dtfft_abstract_transpose_plan dtfft_abstract_transpose_plan module~dtfft_transpose_plan_cuda->module~dtfft_abstract_transpose_plan module~dtfft_transpose_handle_cuda dtfft_transpose_handle_cuda module~dtfft_transpose_plan_cuda->module~dtfft_transpose_handle_cuda module~dtfft_abstract_backend->module~dtfft_nvrtc_kernel module~dtfft_abstract_transpose_plan->module~dtfft_nvrtc_kernel module~dtfft_abstract_transpose_plan->module~dtfft_abstract_backend module~dtfft_plan dtfft_plan module~dtfft_plan->module~dtfft_nvrtc_kernel module~dtfft_plan->module~dtfft_transpose_plan_cuda module~dtfft_plan->module~dtfft_abstract_transpose_plan module~dtfft_transpose_plan_host dtfft_transpose_plan_host module~dtfft_plan->module~dtfft_transpose_plan_host module~dtfft_transpose_handle_cuda->module~dtfft_nvrtc_kernel module~dtfft_transpose_handle_cuda->module~dtfft_abstract_backend module~dtfft_backend_cufftmp_m dtfft_backend_cufftmp_m module~dtfft_transpose_handle_cuda->module~dtfft_backend_cufftmp_m module~dtfft_backend_mpi dtfft_backend_mpi module~dtfft_transpose_handle_cuda->module~dtfft_backend_mpi module~dtfft_backend_nccl_m dtfft_backend_nccl_m module~dtfft_transpose_handle_cuda->module~dtfft_backend_nccl_m module~dtfft dtfft module~dtfft->module~dtfft_plan module~dtfft_api dtfft_api module~dtfft_api->module~dtfft_plan module~dtfft_backend_cufftmp_m->module~dtfft_abstract_backend module~dtfft_backend_mpi->module~dtfft_abstract_backend module~dtfft_backend_nccl_m->module~dtfft_abstract_backend module~dtfft_transpose_plan_host->module~dtfft_abstract_transpose_plan

Variables

Type Visibility Attributes Name Initial
procedure(cuModuleLoadDataEx_interface), public, pointer :: cuModuleLoadDataEx

Fortran pointer to the cuModuleLoadDataEx 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

Interfaces

interface

Launches a CUDA function CUfunction or a CUDA kernel CUkernel.

  • private function run_cuda_kernel(func, in, out, blocks, threads, stream, args, funptr) result(cuResult) bind(C, name="run_cuda_kernel")

    Wrapper around cuLaunchKernel, since I have to idea how to pass array of pointers to cuLaunchKernel.

    Launches a CUDA function CUfunction or a CUDA kernel CUkernel.

    Arguments

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

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

    Input pointer
    type(c_ptr), value :: out

    Output pointer
    type(dim3) :: blocks

    Grid in blocks
    type(dim3) :: threads

    Thread block
    type(dtfft_stream_t), value :: stream

    Stream identifier
    type(kernelArgs) :: args

    Kernel parameters
    type(c_funptr), value :: funptr

    Pointer to cuLaunchKernel

    Return Value integer(kind=c_int)

    Driver result code


Abstract Interfaces

abstract interface

  • private function cuModuleLoadDataEx_interface(mod, image, numOptions, options, optionValues) 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
    character(len=c_char) :: image(*)

    Module data to load
    integer(kind=c_int), value :: numOptions

    Number of options
    type(c_ptr), value :: options

    Options for JIT
    type(c_ptr) :: optionValues

    Option values for JIT

    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
    character(len=c_char) :: name(*)

    Name of function to retrieve

    Return Value integer(kind=c_int)

    Driver result code


Derived Types

type, public, bind(C) ::  kernelArgs

Arguments passed to nvrtc-compiled kernels

Components

Type Visibility Attributes Name Initial
integer(kind=c_int), public :: n_ints = 0

Number of integers provided
integer(kind=c_int), public :: ints(5)

Integer array
integer(kind=c_int), public :: n_ptrs = 0

Number of pointers provided
type(c_ptr), public :: ptrs(3)

Pointer array

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, args) result(cuResult)

Launches a CUDA function CUfunction or a CUDA kernel CUkernel.

Arguments

Type IntentOptional Attributes Name
type(CUfunction) :: func

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

Input pointer
type(c_ptr) :: out

Output pointer
type(dim3) :: blocks

Grid in blocks
type(dim3) :: threads

Thread block
type(dtfft_stream_t) :: stream

Stream identifier
type(kernelArgs) :: args

Kernel parameters

Return Value integer(kind=c_int)

Driver result code