diff --git a/dev/.documenter-siteinfo.json b/dev/.documenter-siteinfo.json index effcd626b3..cc096a6118 100644 --- a/dev/.documenter-siteinfo.json +++ b/dev/.documenter-siteinfo.json @@ -1 +1 @@ -{"documenter":{"julia_version":"1.10.4","generation_timestamp":"2024-08-26T07:38:08","documenter_version":"1.4.0"}} \ No newline at end of file +{"documenter":{"julia_version":"1.10.5","generation_timestamp":"2024-09-04T18:14:38","documenter_version":"1.4.0"}} \ No newline at end of file diff --git a/dev/api/array/index.html b/dev/api/array/index.html index 448066458b..bd88a91311 100644 --- a/dev/api/array/index.html +++ b/dev/api/array/index.html @@ -3,4 +3,4 @@ function gtag(){dataLayer.push(arguments);} gtag('js', new Date()); gtag('config', 'UA-154489943-2', {'page_path': location.pathname + location.search + location.hash}); -

Array programming

The CUDA array type, CuArray, generally implements the Base array interface and all of its expected methods.

+

Array programming

The CUDA array type, CuArray, generally implements the Base array interface and all of its expected methods.

diff --git a/dev/api/compiler/index.html b/dev/api/compiler/index.html index abca5edcb0..6dd825b5c4 100644 --- a/dev/api/compiler/index.html +++ b/dev/api/compiler/index.html @@ -3,8 +3,8 @@ function gtag(){dataLayer.push(arguments);} gtag('js', new Date()); gtag('config', 'UA-154489943-2', {'page_path': location.pathname + location.search + location.hash}); -

Compiler

Execution

The main entry-point to the compiler is the @cuda macro:

CUDA.@cuda — Macro
@cuda [kwargs...] func(args...)

High-level interface for executing code on a GPU. The @cuda macro should prefix a call, with func a callable function or object that should return nothing. It will be compiled to a CUDA function upon first use, and to a certain extent arguments will be converted and managed automatically using cudaconvert. Finally, a call to cudacall is performed, scheduling a kernel launch on the current CUDA context.

Several keyword arguments are supported that influence the behavior of @cuda.

  • launch: whether to launch this kernel, defaults to true. If false the returned kernel object should be launched by calling it and passing arguments again.
  • dynamic: use dynamic parallelism to launch device-side kernels, defaults to false.
  • arguments that influence kernel compilation: see cufunction and dynamic_cufunction
  • arguments that influence kernel launch: see CUDA.HostKernel and CUDA.DeviceKernel
source

If needed, you can use a lower-level API that lets you inspect the compiler kernel:

CUDA.cudaconvert — Function
cudaconvert(x)

This function is called for every argument to be passed to a kernel, allowing it to be converted to a GPU-friendly format. By default, the function does nothing and returns the input object x as-is.

Do not add methods to this function, but instead extend the underlying Adapt.jl package and register methods for the the CUDA.KernelAdaptor type.

source
CUDA.cufunction — Function
cufunction(f, tt=Tuple{}; kwargs...)

Low-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. For a higher-level interface, use @cuda.

The following keyword arguments are supported:

  • minthreads: the required number of threads in a thread block
  • maxthreads: the maximum number of threads in a thread block
  • blocks_per_sm: a minimum number of thread blocks to be scheduled on a single multiprocessor
  • maxregs: the maximum number of registers to be allocated to a single thread (only supported on LLVM 4.0+)
  • name: override the name that the kernel will have in the generated code
  • always_inline: inline all function calls in the kernel
  • fastmath: use less precise square roots and flush denormals
  • cap and ptx: to override the compute capability and PTX version to compile for

The output of this function is automatically cached, i.e. you can simply call cufunction in a hot path without degrading performance. New code will be generated automatically, when when function changes, or when different types or keyword arguments are provided.

source
CUDA.HostKernel — Type
(::HostKernel)(args...; kwargs...)
-(::DeviceKernel)(args...; kwargs...)

Low-level interface to call a compiled kernel, passing GPU-compatible arguments in args. For a higher-level interface, use @cuda.

A HostKernel is callable on the host, and a DeviceKernel is callable on the device (created by @cuda with dynamic=true).

The following keyword arguments are supported:

  • threads (default: 1): Number of threads per block, or a 1-, 2- or 3-tuple of dimensions (e.g. threads=(32, 32) for a 2D block of 32×32 threads). Use threadIdx() and blockDim() to query from within the kernel.
  • blocks (default: 1): Number of thread blocks to launch, or a 1-, 2- or 3-tuple of dimensions (e.g. blocks=(2, 4, 2) for a 3D grid of blocks). Use blockIdx() and gridDim() to query from within the kernel.
  • shmem(default: 0): Amount of dynamic shared memory in bytes to allocate per thread block; used by CuDynamicSharedArray.
  • stream (default: stream()): CuStream to launch the kernel on.
  • cooperative (default: false): whether to launch a cooperative kernel that supports grid synchronization (see CG.this_grid and CG.sync). Note that this requires care wrt. the number of blocks launched.
source
CUDA.version — Function
version(k::HostKernel)

Queries the PTX and SM versions a kernel was compiled for. Returns a named tuple.

source
CUDA.maxthreads — Function
maxthreads(k::HostKernel)

Queries the maximum amount of threads a kernel can use in a single block.

source
CUDA.memory — Function
memory(k::HostKernel)

Queries the local, shared and constant memory usage of a compiled kernel in bytes. Returns a named tuple.

source

Reflection

If you want to inspect generated code, you can use macros that resemble functionality from the InteractiveUtils standard library:

@device_code_lowered
+
Compiler
Execution
The main entry-point to the compiler is the @cuda macro:
CUDA.@cuda — Macro
@cuda [kwargs...] func(args...)
High-level interface for executing code on a GPU. The @cuda macro should prefix a call, with func a callable function or object that should return nothing. It will be compiled to a CUDA function upon first use, and to a certain extent arguments will be converted and managed automatically using cudaconvert. Finally, a call to cudacall is performed, scheduling a kernel launch on the current CUDA context.
Several keyword arguments are supported that influence the behavior of @cuda.
  • launch: whether to launch this kernel, defaults to true. If false the returned kernel object should be launched by calling it and passing arguments again.
  • dynamic: use dynamic parallelism to launch device-side kernels, defaults to false.
  • arguments that influence kernel compilation: see cufunction and dynamic_cufunction
  • arguments that influence kernel launch: see CUDA.HostKernel and CUDA.DeviceKernel
source
If needed, you can use a lower-level API that lets you inspect the compiler kernel:
CUDA.cudaconvert — Function
cudaconvert(x)
This function is called for every argument to be passed to a kernel, allowing it to be converted to a GPU-friendly format. By default, the function does nothing and returns the input object x as-is.
Do not add methods to this function, but instead extend the underlying Adapt.jl package and register methods for the the CUDA.KernelAdaptor type.
source
CUDA.cufunction — Function
cufunction(f, tt=Tuple{}; kwargs...)
Low-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. For a higher-level interface, use @cuda.
The following keyword arguments are supported:
  • minthreads: the required number of threads in a thread block
  • maxthreads: the maximum number of threads in a thread block
  • blocks_per_sm: a minimum number of thread blocks to be scheduled on a single multiprocessor
  • maxregs: the maximum number of registers to be allocated to a single thread (only supported on LLVM 4.0+)
  • name: override the name that the kernel will have in the generated code
  • always_inline: inline all function calls in the kernel
  • fastmath: use less precise square roots and flush denormals
  • cap and ptx: to override the compute capability and PTX version to compile for
The output of this function is automatically cached, i.e. you can simply call cufunction in a hot path without degrading performance. New code will be generated automatically, when when function changes, or when different types or keyword arguments are provided.
source
CUDA.HostKernel — Type
(::HostKernel)(args...; kwargs...)
+(::DeviceKernel)(args...; kwargs...)
Low-level interface to call a compiled kernel, passing GPU-compatible arguments in args. For a higher-level interface, use @cuda.
A HostKernel is callable on the host, and a DeviceKernel is callable on the device (created by @cuda with dynamic=true).
The following keyword arguments are supported:
  • threads (default: 1): Number of threads per block, or a 1-, 2- or 3-tuple of dimensions (e.g. threads=(32, 32) for a 2D block of 32×32 threads). Use threadIdx() and blockDim() to query from within the kernel.
  • blocks (default: 1): Number of thread blocks to launch, or a 1-, 2- or 3-tuple of dimensions (e.g. blocks=(2, 4, 2) for a 3D grid of blocks). Use blockIdx() and gridDim() to query from within the kernel.
  • shmem(default: 0): Amount of dynamic shared memory in bytes to allocate per thread block; used by CuDynamicSharedArray.
  • stream (default: stream()): CuStream to launch the kernel on.
  • cooperative (default: false): whether to launch a cooperative kernel that supports grid synchronization (see CG.this_grid and CG.sync). Note that this requires care wrt. the number of blocks launched.
source
CUDA.version — Function
version(k::HostKernel)
Queries the PTX and SM versions a kernel was compiled for. Returns a named tuple.
source
CUDA.maxthreads — Function
maxthreads(k::HostKernel)
Queries the maximum amount of threads a kernel can use in a single block.
source
CUDA.memory — Function
memory(k::HostKernel)
Queries the local, shared and constant memory usage of a compiled kernel in bytes. Returns a named tuple.
source
Reflection
If you want to inspect generated code, you can use macros that resemble functionality from the InteractiveUtils standard library:
@device_code_lowered
 @device_code_typed
 @device_code_warntype
 @device_code_llvm
@@ -14,5 +14,5 @@
 CUDA.code_warntype
 CUDA.code_llvm
 CUDA.code_ptx
-CUDA.code_sass
For more information, please consult the GPUCompiler.jl documentation. Only the code_sass functionality is actually defined in CUDA.jl:
CUDA.code_sass — Function
code_sass([io], f, types; raw=false)
-code_sass(f, [io]; raw=false)
Prints the SASS code corresponding to one or more CUDA modules to io, which defaults to stdout.
If providing both f and types, it is assumed that this uniquely identifies a kernel function, for which SASS code will be generated, and printed to io.
If only providing a callable function f, typically specified using the do syntax, the SASS code for all modules executed during evaluation of f will be printed. This can be convenient to display the SASS code for functions whose source code is not available.
  • raw: dump the assembly like nvdisasm reports it, without post-processing;
  • in the case of specifying f and types: all keyword arguments from cufunction
See also: @device_code_sass
source

+CUDA.code_sass

For more information, please consult the GPUCompiler.jl documentation. Only the code_sass functionality is actually defined in CUDA.jl:

CUDA.code_sass — Function
code_sass([io], f, types; raw=false)
+code_sass(f, [io]; raw=false)

Prints the SASS code corresponding to one or more CUDA modules to io, which defaults to stdout.

If providing both f and types, it is assumed that this uniquely identifies a kernel function, for which SASS code will be generated, and printed to io.

If only providing a callable function f, typically specified using the do syntax, the SASS code for all modules executed during evaluation of f will be printed. This can be convenient to display the SASS code for functions whose source code is not available.

  • raw: dump the assembly like nvdisasm reports it, without post-processing;
  • in the case of specifying f and types: all keyword arguments from cufunction

See also: @device_code_sass

source
diff --git a/dev/api/essentials/index.html b/dev/api/essentials/index.html index 425a0c1452..1ce41181d8 100644 --- a/dev/api/essentials/index.html +++ b/dev/api/essentials/index.html @@ -3,8 +3,8 @@ function gtag(){dataLayer.push(arguments);} gtag('js', new Date()); gtag('config', 'UA-154489943-2', {'page_path': location.pathname + location.search + location.hash}); -

Essentials

Initialization

CUDA.functional — Method
functional(show_reason=false)

Check if the package has been configured successfully and is ready to use.

This call is intended for packages that support conditionally using an available GPU. If you fail to check whether CUDA is functional, actual use of functionality might warn and error.

source
CUDA.has_cuda — Function
has_cuda()::Bool

Check whether the local system provides an installation of the CUDA driver and runtime. Use this function if your code loads packages that require CUDA.jl. ```

source
CUDA.has_cuda_gpu — Function
has_cuda_gpu()::Bool

Check whether the local system provides an installation of the CUDA driver and runtime, and if it contains a CUDA-capable GPU. See has_cuda for more details.

Note that this function initializes the CUDA API in order to check for the number of GPUs.

source

Global state

CUDA.context — Function
context(ptr)

Identify the context memory was allocated in.

source
context()::CuContext

Get or create a CUDA context for the current thread (as opposed to current_context which may return nothing if there is no context bound to the current thread).

source
CUDA.context! — Function
context!(ctx::CuContext)
-context!(ctx::CuContext) do ... end

Bind the current host thread to the context ctx. Returns the previously-bound context. If used with do-block syntax, the change is only temporary.

Note that the contexts used with this call should be previously acquired by calling context, and not arbitrary contexts created by calling the CuContext constructor.

source
CUDA.device! — Function
device!(dev::Integer)
+
Essentials
Initialization
CUDA.functional — Method
functional(show_reason=false)
Check if the package has been configured successfully and is ready to use.
This call is intended for packages that support conditionally using an available GPU. If you fail to check whether CUDA is functional, actual use of functionality might warn and error.
source
CUDA.has_cuda — Function
has_cuda()::Bool
Check whether the local system provides an installation of the CUDA driver and runtime. Use this function if your code loads packages that require CUDA.jl. ```
source
CUDA.has_cuda_gpu — Function
has_cuda_gpu()::Bool
Check whether the local system provides an installation of the CUDA driver and runtime, and if it contains a CUDA-capable GPU. See has_cuda for more details.
Note that this function initializes the CUDA API in order to check for the number of GPUs.
source
Global state
CUDA.context — Function
context(ptr)
Identify the context memory was allocated in.
source
context()::CuContext
Get or create a CUDA context for the current thread (as opposed to current_context which may return nothing if there is no context bound to the current thread).
source
CUDA.context! — Function
context!(ctx::CuContext)
+context!(ctx::CuContext) do ... end
Bind the current host thread to the context ctx. Returns the previously-bound context. If used with do-block syntax, the change is only temporary.
Note that the contexts used with this call should be previously acquired by calling context, and not arbitrary contexts created by calling the CuContext constructor.
source
CUDA.device! — Function
device!(dev::Integer)
 device!(dev::CuDevice)
-device!(dev) do ... end
Sets dev as the current active device for the calling host thread. Devices can be specified by integer id, or as a CuDevice (slightly faster). Both functions can be used with do-block syntax, in which case the device is only changed temporarily, without changing the default device used to initialize new threads or tasks.
Calling this function at the start of a session will make sure CUDA is initialized (i.e., a primary context will be created and activated).
source
CUDA.device_reset! — Function
device_reset!(dev::CuDevice=device())
Reset the CUDA state associated with a device. This call with release the underlying context, at which point any objects allocated in that context will be invalidated.
Note that this does not guarantee to free up all memory allocations, as many are not bound to a context, so it is generally not useful to call this function to free up memory.
Warning
This function is only reliable on CUDA driver >= v12.0, and may lead to crashes if used on older drivers.
source
CUDA.stream — Function
stream()
Get the CUDA stream that should be used as the default one for the currently executing task.
source
CUDA.stream! — Function
stream!(::CuStream)
-stream!(::CuStream) do ... end
Change the default CUDA stream for the currently executing task, temporarily if using the do-block version of this function.
source

+device!(dev) do ... end

Sets dev as the current active device for the calling host thread. Devices can be specified by integer id, or as a CuDevice (slightly faster). Both functions can be used with do-block syntax, in which case the device is only changed temporarily, without changing the default device used to initialize new threads or tasks.

Calling this function at the start of a session will make sure CUDA is initialized (i.e., a primary context will be created and activated).

source
CUDA.device_reset! — Function
device_reset!(dev::CuDevice=device())

Reset the CUDA state associated with a device. This call with release the underlying context, at which point any objects allocated in that context will be invalidated.

Note that this does not guarantee to free up all memory allocations, as many are not bound to a context, so it is generally not useful to call this function to free up memory.

Warning

This function is only reliable on CUDA driver >= v12.0, and may lead to crashes if used on older drivers.

source
CUDA.stream — Function
stream()

Get the CUDA stream that should be used as the default one for the currently executing task.

source
CUDA.stream! — Function
stream!(::CuStream)
+stream!(::CuStream) do ... end

Change the default CUDA stream for the currently executing task, temporarily if using the do-block version of this function.

source
diff --git a/dev/api/kernel/index.html b/dev/api/kernel/index.html index 472e3b8558..9ab2d9dee5 100644 --- a/dev/api/kernel/index.html +++ b/dev/api/kernel/index.html @@ -3,14 +3,14 @@ function gtag(){dataLayer.push(arguments);} gtag('js', new Date()); gtag('config', 'UA-154489943-2', {'page_path': location.pathname + location.search + location.hash}); -

Kernel programming

This section lists the package's public functionality that corresponds to special CUDA functions for use in device code. It is loosely organized according to the C language extensions appendix from the CUDA C programming guide. For more information about certain intrinsics, refer to the aforementioned NVIDIA documentation.

Indexing and dimensions

CUDA.warpsize — Function
warpsize(dev::CuDevice)

Returns the warp size (in threads) of the device.

source
warpsize()::Int32

Returns the warp size (in threads).

source
CUDA.laneid — Function
laneid()::Int32

Returns the thread's lane within the warp.

source
CUDA.active_mask — Function
active_mask()

Returns a 32-bit mask indicating which threads in a warp are active with the current executing thread.

source

Device arrays

CUDA.jl provides a primitive, lightweight array type to manage GPU data organized in an plain, dense fashion. This is the device-counterpart to the CuArray, and implements (part of) the array interface as well as other functionality for use on the GPU:

CUDA.CuDeviceArray — Type
CuDeviceArray{T,N,A}(ptr, dims, [maxsize])

Construct an N-dimensional dense CUDA device array with element type T wrapping a pointer, where N is determined from the length of dims and T is determined from the type of ptr. dims may be a single scalar, or a tuple of integers corresponding to the lengths in each dimension). If the rank N is supplied explicitly as in Array{T,N}(dims), then it must match the length of dims. The same applies to the element type T, which should match the type of the pointer ptr.

source
CUDA.Const — Type
Const(A::CuDeviceArray)

Mark a CuDeviceArray as constant/read-only. The invariant guaranteed is that you will not modify an CuDeviceArray for the duration of the current kernel.

This API can only be used on devices with compute capability 3.5 or higher.

Warning

Experimental API. Subject to change without deprecation.

source

Memory types

Shared memory

CUDA.CuStaticSharedArray — Function
CuStaticSharedArray(T::Type, dims) -> CuDeviceArray{T,N,AS.Shared}

Get an array of type T and dimensions dims (either an integer length or tuple shape) pointing to a statically-allocated piece of shared memory. The type should be statically inferable and the dimensions should be constant, or an error will be thrown and the generator function will be called dynamically.

source
CUDA.CuDynamicSharedArray — Function
CuDynamicSharedArray(T::Type, dims, offset::Integer=0) -> CuDeviceArray{T,N,AS.Shared}

Get an array of type T and dimensions dims (either an integer length or tuple shape) pointing to a dynamically-allocated piece of shared memory. The type should be statically inferable or an error will be thrown and the generator function will be called dynamically.

Note that the amount of dynamic shared memory needs to specified when launching the kernel.

Optionally, an offset parameter indicating how many bytes to add to the base shared memory pointer can be specified. This is useful when dealing with a heterogeneous buffer of dynamic shared memory; in the case of a homogeneous multi-part buffer it is preferred to use view.

source

Texture memory

CUDA.CuDeviceTexture — Type
CuDeviceTexture{T,N,M,NC,I}

N-dimensional device texture with elements of type T. This type is the device-side counterpart of CuTexture{T,N,P}, and can be used to access textures using regular indexing notation. If NC is true, indices used by these accesses should be normalized, i.e., fall into the [0,1) domain. The I type parameter indicates the kind of interpolation that happens when indexing into this texture. The source memory of the texture is specified by the M parameter, either linear memory or a texture array.

Device-side texture objects cannot be created directly, but should be created host-side using CuTexture{T,N,P} and passed to the kernel as an argument.

Warning

Experimental API. Subject to change without deprecation.

source

Synchronization

CUDA.sync_threads — Function
sync_threads()

Waits until all threads in the thread block have reached this point and all global and shared memory accesses made by these threads prior to sync_threads() are visible to all threads in the block.

source
CUDA.sync_threads_count — Function
sync_threads_count(predicate)

Identical to sync_threads() with the additional feature that it evaluates predicate for all threads of the block and returns the number of threads for which predicate evaluates to true.

source
CUDA.sync_threads_and — Function
sync_threads_and(predicate)

Identical to sync_threads() with the additional feature that it evaluates predicate for all threads of the block and returns true if and only if predicate evaluates to true for all of them.

source
CUDA.sync_threads_or — Function
sync_threads_or(predicate)

Identical to sync_threads() with the additional feature that it evaluates predicate for all threads of the block and returns true if and only if predicate evaluates to true for any of them.

source
CUDA.sync_warp — Function
sync_warp(mask::Integer=FULL_MASK)

Waits threads in the warp, selected by means of the bitmask mask, have reached this point and all global and shared memory accesses made by these threads prior to sync_warp() are visible to those threads in the warp. The default value for mask selects all threads in the warp.

Note

Requires CUDA >= 9.0 and sm_6.2

source
CUDA.threadfence_block — Function
threadfence_block()

A memory fence that ensures that:

  • All writes to all memory made by the calling thread before the call to threadfence_block() are observed by all threads in the block of the calling thread as occurring before all writes to all memory made by the calling thread after the call to threadfence_block()
  • All reads from all memory made by the calling thread before the call to threadfence_block() are ordered before all reads from all memory made by the calling thread after the call to threadfence_block().
source
CUDA.threadfence — Function
threadfence()

A memory fence that acts as threadfence_block for all threads in the block of the calling thread and also ensures that no writes to all memory made by the calling thread after the call to threadfence() are observed by any thread in the device as occurring before any write to all memory made by the calling thread before the call to threadfence().

Note that for this ordering guarantee to be true, the observing threads must truly observe the memory and not cached versions of it; this is requires the use of volatile loads and stores, which is not available from Julia right now.

source
CUDA.threadfence_system — Function
threadfence_system()

A memory fence that acts as threadfence_block for all threads in the block of the calling thread and also ensures that all writes to all memory made by the calling thread before the call to threadfence_system() are observed by all threads in the device, host threads, and all threads in peer devices as occurring before all writes to all memory made by the calling thread after the call to threadfence_system().

source

Time functions

CUDA.clock — Function
clock(UInt32)

Returns the value of a per-multiprocessor counter that is incremented every clock cycle.

source
clock(UInt64)

Returns the value of a per-multiprocessor counter that is incremented every clock cycle.

source
CUDA.nanosleep — Function
nanosleep(t)

Puts a thread for a given amount t(in nanoseconds).

Note

Requires CUDA >= 10.0 and sm_6.2

source

Warp-level functions

Voting

The warp vote functions allow the threads of a given warp to perform a reduction-and-broadcast operation. These functions take as input a boolean predicate from each thread in the warp and evaluate it. The results of that evaluation are combined (reduced) across the active threads of the warp in one different ways, broadcasting a single return value to each participating thread.

CUDA.vote_all_sync — Function
vote_all_sync(mask::UInt32, predicate::Bool)

Evaluate predicate for all active threads of the warp and return whether predicate is true for all of them.

source
CUDA.vote_any_sync — Function
vote_any_sync(mask::UInt32, predicate::Bool)

Evaluate predicate for all active threads of the warp and return whether predicate is true for any of them.

source
CUDA.vote_uni_sync — Function
vote_uni_sync(mask::UInt32, predicate::Bool)

Evaluate predicate for all active threads of the warp and return whether predicate is the same for any of them.

source
CUDA.vote_ballot_sync — Function
vote_ballot_sync(mask::UInt32, predicate::Bool)

Evaluate predicate for all active threads of the warp and return an integer whose Nth bit is set if and only if predicate is true for the Nth thread of the warp and the Nth thread is active.

source

Shuffle

CUDA.shfl_sync — Function
shfl_sync(threadmask::UInt32, val, lane::Integer, width::Integer=32)

Shuffle a value from a directly indexed lane lane, and synchronize threads according to threadmask.

source
CUDA.shfl_up_sync — Function
shfl_up_sync(threadmask::UInt32, val, delta::Integer, width::Integer=32)

Shuffle a value from a lane with lower ID relative to caller, and synchronize threads according to threadmask.

source
CUDA.shfl_down_sync — Function
shfl_down_sync(threadmask::UInt32, val, delta::Integer, width::Integer=32)

Shuffle a value from a lane with higher ID relative to caller, and synchronize threads according to threadmask.

source
CUDA.shfl_xor_sync — Function
shfl_xor_sync(threadmask::UInt32, val, mask::Integer, width::Integer=32)

Shuffle a value from a lane based on bitwise XOR of own lane ID with mask, and synchronize threads according to threadmask.

source

Formatted Output

CUDA.@cushow — Macro
@cushow(ex)

GPU analog of Base.@show. It comes with the same type restrictions as @cuprintf.

@cushow threadIdx().x
source
CUDA.@cuprint — Macro
@cuprint(xs...)
+
Kernel programming
This section lists the package's public functionality that corresponds to special CUDA functions for use in device code. It is loosely organized according to the C language extensions appendix from the CUDA C programming guide. For more information about certain intrinsics, refer to the aforementioned NVIDIA documentation.
Indexing and dimensions
CUDA.warpsize — Function
warpsize(dev::CuDevice)
Returns the warp size (in threads) of the device.
source
warpsize()::Int32
Returns the warp size (in threads).
source
CUDA.laneid — Function
laneid()::Int32
Returns the thread's lane within the warp.
source
CUDA.active_mask — Function
active_mask()
Returns a 32-bit mask indicating which threads in a warp are active with the current executing thread.
source
Device arrays
CUDA.jl provides a primitive, lightweight array type to manage GPU data organized in an plain, dense fashion. This is the device-counterpart to the CuArray, and implements (part of) the array interface as well as other functionality for use on the GPU:
CUDA.CuDeviceArray — Type
CuDeviceArray{T,N,A}(ptr, dims, [maxsize])
Construct an N-dimensional dense CUDA device array with element type T wrapping a pointer, where N is determined from the length of dims and T is determined from the type of ptr. dims may be a single scalar, or a tuple of integers corresponding to the lengths in each dimension). If the rank N is supplied explicitly as in Array{T,N}(dims), then it must match the length of dims. The same applies to the element type T, which should match the type of the pointer ptr.
source
CUDA.Const — Type
Const(A::CuDeviceArray)
Mark a CuDeviceArray as constant/read-only. The invariant guaranteed is that you will not modify an CuDeviceArray for the duration of the current kernel.
This API can only be used on devices with compute capability 3.5 or higher.
Warning
Experimental API. Subject to change without deprecation.
source
Memory types
Shared memory
CUDA.CuStaticSharedArray — Function
CuStaticSharedArray(T::Type, dims) -> CuDeviceArray{T,N,AS.Shared}
Get an array of type T and dimensions dims (either an integer length or tuple shape) pointing to a statically-allocated piece of shared memory. The type should be statically inferable and the dimensions should be constant, or an error will be thrown and the generator function will be called dynamically.
source
CUDA.CuDynamicSharedArray — Function
CuDynamicSharedArray(T::Type, dims, offset::Integer=0) -> CuDeviceArray{T,N,AS.Shared}
Get an array of type T and dimensions dims (either an integer length or tuple shape) pointing to a dynamically-allocated piece of shared memory. The type should be statically inferable or an error will be thrown and the generator function will be called dynamically.
Note that the amount of dynamic shared memory needs to specified when launching the kernel.
Optionally, an offset parameter indicating how many bytes to add to the base shared memory pointer can be specified. This is useful when dealing with a heterogeneous buffer of dynamic shared memory; in the case of a homogeneous multi-part buffer it is preferred to use view.
source
Texture memory
CUDA.CuDeviceTexture — Type
CuDeviceTexture{T,N,M,NC,I}
N-dimensional device texture with elements of type T. This type is the device-side counterpart of CuTexture{T,N,P}, and can be used to access textures using regular indexing notation. If NC is true, indices used by these accesses should be normalized, i.e., fall into the [0,1) domain. The I type parameter indicates the kind of interpolation that happens when indexing into this texture. The source memory of the texture is specified by the M parameter, either linear memory or a texture array.
Device-side texture objects cannot be created directly, but should be created host-side using CuTexture{T,N,P} and passed to the kernel as an argument.
Warning
Experimental API. Subject to change without deprecation.
source
Synchronization
CUDA.sync_threads — Function
sync_threads()
Waits until all threads in the thread block have reached this point and all global and shared memory accesses made by these threads prior to sync_threads() are visible to all threads in the block.
source
CUDA.sync_threads_count — Function
sync_threads_count(predicate)
Identical to sync_threads() with the additional feature that it evaluates predicate for all threads of the block and returns the number of threads for which predicate evaluates to true.
source
CUDA.sync_threads_and — Function
sync_threads_and(predicate)
Identical to sync_threads() with the additional feature that it evaluates predicate for all threads of the block and returns true if and only if predicate evaluates to true for all of them.
source
CUDA.sync_threads_or — Function
sync_threads_or(predicate)
Identical to sync_threads() with the additional feature that it evaluates predicate for all threads of the block and returns true if and only if predicate evaluates to true for any of them.
source
CUDA.sync_warp — Function
sync_warp(mask::Integer=FULL_MASK)
Waits threads in the warp, selected by means of the bitmask mask, have reached this point and all global and shared memory accesses made by these threads prior to sync_warp() are visible to those threads in the warp. The default value for mask selects all threads in the warp.
Note
Requires CUDA >= 9.0 and sm_6.2
source
CUDA.threadfence_block — Function
threadfence_block()
A memory fence that ensures that:
  • All writes to all memory made by the calling thread before the call to threadfence_block() are observed by all threads in the block of the calling thread as occurring before all writes to all memory made by the calling thread after the call to threadfence_block()
  • All reads from all memory made by the calling thread before the call to threadfence_block() are ordered before all reads from all memory made by the calling thread after the call to threadfence_block().
source
CUDA.threadfence — Function
threadfence()
A memory fence that acts as threadfence_block for all threads in the block of the calling thread and also ensures that no writes to all memory made by the calling thread after the call to threadfence() are observed by any thread in the device as occurring before any write to all memory made by the calling thread before the call to threadfence().
Note that for this ordering guarantee to be true, the observing threads must truly observe the memory and not cached versions of it; this is requires the use of volatile loads and stores, which is not available from Julia right now.
source
CUDA.threadfence_system — Function
threadfence_system()
A memory fence that acts as threadfence_block for all threads in the block of the calling thread and also ensures that all writes to all memory made by the calling thread before the call to threadfence_system() are observed by all threads in the device, host threads, and all threads in peer devices as occurring before all writes to all memory made by the calling thread after the call to threadfence_system().
source
Time functions
CUDA.clock — Function
clock(UInt32)
Returns the value of a per-multiprocessor counter that is incremented every clock cycle.
source
clock(UInt64)
Returns the value of a per-multiprocessor counter that is incremented every clock cycle.
source
CUDA.nanosleep — Function
nanosleep(t)
Puts a thread for a given amount t(in nanoseconds).
Note
Requires CUDA >= 10.0 and sm_6.2
source
Warp-level functions
Voting
The warp vote functions allow the threads of a given warp to perform a reduction-and-broadcast operation. These functions take as input a boolean predicate from each thread in the warp and evaluate it. The results of that evaluation are combined (reduced) across the active threads of the warp in one different ways, broadcasting a single return value to each participating thread.
CUDA.vote_all_sync — Function
vote_all_sync(mask::UInt32, predicate::Bool)
Evaluate predicate for all active threads of the warp and return whether predicate is true for all of them.
source
CUDA.vote_any_sync — Function
vote_any_sync(mask::UInt32, predicate::Bool)
Evaluate predicate for all active threads of the warp and return whether predicate is true for any of them.
source
CUDA.vote_uni_sync — Function
vote_uni_sync(mask::UInt32, predicate::Bool)
Evaluate predicate for all active threads of the warp and return whether predicate is the same for any of them.
source
CUDA.vote_ballot_sync — Function
vote_ballot_sync(mask::UInt32, predicate::Bool)
Evaluate predicate for all active threads of the warp and return an integer whose Nth bit is set if and only if predicate is true for the Nth thread of the warp and the Nth thread is active.
source
Shuffle
CUDA.shfl_sync — Function
shfl_sync(threadmask::UInt32, val, lane::Integer, width::Integer=32)
Shuffle a value from a directly indexed lane lane, and synchronize threads according to threadmask.
source
CUDA.shfl_up_sync — Function
shfl_up_sync(threadmask::UInt32, val, delta::Integer, width::Integer=32)
Shuffle a value from a lane with lower ID relative to caller, and synchronize threads according to threadmask.
source
CUDA.shfl_down_sync — Function
shfl_down_sync(threadmask::UInt32, val, delta::Integer, width::Integer=32)
Shuffle a value from a lane with higher ID relative to caller, and synchronize threads according to threadmask.
source
CUDA.shfl_xor_sync — Function
shfl_xor_sync(threadmask::UInt32, val, mask::Integer, width::Integer=32)
Shuffle a value from a lane based on bitwise XOR of own lane ID with mask, and synchronize threads according to threadmask.
source
Formatted Output
CUDA.@cushow — Macro
@cushow(ex)
GPU analog of Base.@show. It comes with the same type restrictions as @cuprintf.
@cushow threadIdx().x
source
CUDA.@cuprint — Macro
@cuprint(xs...)
 @cuprintln(xs...)
Print a textual representation of values xs to standard output from the GPU. The functionality builds on @cuprintf, and is intended as a more use friendly alternative of that API. However, that also means there's only limited support for argument types, handling 16/32/64 signed and unsigned integers, 32 and 64-bit floating point numbers, Cchars and pointers. For more complex output, use @cuprintf directly.
Limited string interpolation is also possible:
    @cuprint("Hello, World ", 42, "\n")
-    @cuprint "Hello, World $(42)\n"
source
CUDA.@cuprintln — Macro
@cuprint(xs...)
 @cuprintln(xs...)
Print a textual representation of values xs to standard output from the GPU. The functionality builds on @cuprintf, and is intended as a more use friendly alternative of that API. However, that also means there's only limited support for argument types, handling 16/32/64 signed and unsigned integers, 32 and 64-bit floating point numbers, Cchars and pointers. For more complex output, use @cuprintf directly.
Limited string interpolation is also possible:
    @cuprint("Hello, World ", 42, "\n")
-    @cuprint "Hello, World $(42)\n"
source
CUDA.@cuprintf — Macro
@cuprintf("%Fmt", args...)
Print a formatted string in device context on the host standard output.
Note that this is not a fully C-compliant printf implementation; see the CUDA documentation for supported options and inputs.
Also beware that it is an untyped, and unforgiving printf implementation. Type widths need to match, eg. printing a 64-bit Julia integer requires the %ld formatting string.
source
Assertions
CUDA.@cuassert — Macro
@assert cond [text]
Signal assertion failure to the CUDA driver if cond is false. Preferred syntax for writing assertions, mimicking Base.@assert. Message text is optionally displayed upon assertion failure.
Warning
A failed assertion will crash the GPU, so use sparingly as a debugging tool. Furthermore, the assertion might be disabled at various optimization levels, and thus should not cause any side-effects.
source
Atomics
A high-level macro is available to annotate expressions with:
CUDA.@atomic — Macro
@atomic a[I] = op(a[I], val)
-@atomic a[I] ...= val
Atomically perform a sequence of operations that loads an array element a[I], performs the operation op on that value and a second value val, and writes the result back to the array. This sequence can be written out as a regular assignment, in which case the same array element should be used in the left and right hand side of the assignment, or as an in-place application of a known operator. In both cases, the array reference should be pure and not induce any side-effects.
Warn
This interface is experimental, and might change without warning.  Use the lower-level atomic_...! functions for a stable API, albeit one limited to natively-supported ops.
source
If your expression is not recognized, or you need more control, use the underlying functions:
CUDA.atomic_cas! — Function
atomic_cas!(ptr::LLVMPtr{T}, cmp::T, val::T)
Reads the value old located at address ptr and compare with cmp. If old equals to cmp, stores val at the same address. Otherwise, doesn't change the value old. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64. Additionally, on GPU hardware with compute capability 7.0+, values of type UInt16 are supported.
source
CUDA.atomic_xchg! — Function
atomic_xchg!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr and stores val at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_add! — Function
atomic_add!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old + val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32, UInt64, and Float32. Additionally, on GPU hardware with compute capability 6.0+, values of type Float64 are supported.
source
CUDA.atomic_sub! — Function
atomic_sub!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old - val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_and! — Function
atomic_and!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old & val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_or! — Function
atomic_or!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old | val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_xor! — Function
atomic_xor!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old ⊻ val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_min! — Function
atomic_min!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes min(old, val), and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_max! — Function
atomic_max!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes max(old, val), and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_inc! — Function
atomic_inc!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes ((old >= val) ? 0 : (old+1)), and stores the result back to memory at the same address. These three operations are performed in one atomic transaction. The function returns old.
This operation is only supported for values of type Int32.
source
CUDA.atomic_dec! — Function
atomic_dec!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes (((old == 0) | (old > val)) ? val : (old-1) ), and stores the result back to memory at the same address. These three operations are performed in one atomic transaction. The function returns old.
This operation is only supported for values of type Int32.
source
Dynamic parallelism
Similarly to launching kernels from the host, you can use @cuda while passing dynamic=true for launching kernels from the device. A lower-level API is available as well:
CUDA.dynamic_cufunction — Function
dynamic_cufunction(f, tt=Tuple{})
Low-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. Device-side equivalent of CUDA.cufunction.
No keyword arguments are supported.
source
CUDA.DeviceKernel — Type
(::HostKernel)(args...; kwargs...)
-(::DeviceKernel)(args...; kwargs...)
Low-level interface to call a compiled kernel, passing GPU-compatible arguments in args. For a higher-level interface, use @cuda.
A HostKernel is callable on the host, and a DeviceKernel is callable on the device (created by @cuda with dynamic=true).
The following keyword arguments are supported:
  • threads (default: 1): Number of threads per block, or a 1-, 2- or 3-tuple of dimensions (e.g. threads=(32, 32) for a 2D block of 32×32 threads). Use threadIdx() and blockDim() to query from within the kernel.
  • blocks (default: 1): Number of thread blocks to launch, or a 1-, 2- or 3-tuple of dimensions (e.g. blocks=(2, 4, 2) for a 3D grid of blocks). Use blockIdx() and gridDim() to query from within the kernel.
  • shmem(default: 0): Amount of dynamic shared memory in bytes to allocate per thread block; used by CuDynamicSharedArray.
  • stream (default: stream()): CuStream to launch the kernel on.
  • cooperative (default: false): whether to launch a cooperative kernel that supports grid synchronization (see CG.this_grid and CG.sync). Note that this requires care wrt. the number of blocks launched.
source
Cooperative groups
CUDA.CG — Module
CUDA.jl's cooperative groups implementation.
Cooperative groups in CUDA offer a structured approach to synchronize and communicate among threads. They allow developers to define specific groups of threads, providing a means to fine-tune inter-thread communication granularity. By offering a more nuanced alternative to traditional CUDA synchronization methods, cooperative groups enable a more controlled and efficient parallel decomposition in kernel design.
The following functionality is available in CUDA.jl:
  • implicit groups: thread blocks, grid groups, and coalesced groups.
  • synchronization: sync, barrier_arrive, barrier_wait
  • warp collectives for coalesced groups: shuffle and voting
  • data transfer: memcpy_async, wait and wait_prior
Noteworthy missing functionality:
  • implicit groups: clusters, and multi-grid groups (which are deprecated)
  • explicit groups: tiling and partitioning
source
Group construction and properties
CUDA.CG.thread_rank — Function
thread_rank(group)
Returns the linearized rank of the calling thread along the interval [1, num_threads()].
source
CUDA.CG.thread_block — Type
thread_block <: thread_group
Every GPU kernel is executed by a grid of thread blocks, and threads within each block are guaranteed to reside on the same streaming multiprocessor. A thread_block represents a thread block whose dimensions are not known until runtime.
Constructed via this_thread_block
source
CUDA.CG.group_index — Function
group_index(tb::thread_block)
3-Dimensional index of the block within the launched grid.
source
CUDA.CG.grid_group — Type
grid_group <: thread_group
Threads within this this group are guaranteed to be co-resident on the same device within the same launched kernel. To use this group, the kernel must have been launched with @cuda cooperative=true, and the device must support it (queryable device attribute).
Constructed via this_grid.
source
CUDA.CG.coalesced_group — Type
coalesced_group <: thread_group
A group representing the current set of converged threads in a warp. The size of the group is not guaranteed and it may return a group of only one thread (itself).
This group exposes warp-synchronous builtins. Constructed via coalesced_threads.
source
CUDA.CG.meta_group_size — Function
meta_group_size(cg::coalesced_group)
Total number of partitions created out of all CTAs when the group was created.
source
Synchronization
Data transfer
CUDA.CG.memcpy_async — Function
memcpy_async(group, dst, src, bytes)
Perform a group-wide collective memory copy from src to dst of bytes bytes. This operation may be performed asynchronously, so you should wait or wait_prior before using the data. It is only supported by thread blocks and coalesced groups.
For this operation to be performed asynchronously, the following conditions must be met:
  • the source and destination memory should be aligned to 4, 8 or 16 bytes. this will be deduced from the datatype, but can also be specified explicitly using CUDA.align.
  • the source should be global memory, and the destination should be shared memory.
  • the device should have compute capability 8.0 or higher.
source
Math
Many mathematical functions are provided by the libdevice library, and are wrapped by CUDA.jl. These functions are used to implement well-known functions from the Julia standard library and packages like SpecialFunctions.jl, e.g., calling the cos function will automatically use __nv_cos from libdevice if possible.
Some functions do not have a counterpart in the Julia ecosystem, those have to be called directly. For example, to call __nv_logb or __nv_logbf you use CUDA.logb in a kernel.
For a list of available functions, look at src/device/intrinsics/math.jl.
WMMA
Warp matrix multiply-accumulate (WMMA) is a CUDA API to access Tensor Cores, a new hardware feature in Volta GPUs to perform mixed precision matrix multiply-accumulate operations. The interface is split in two levels, both available in the WMMA submodule: low level wrappers around the LLVM intrinsics, and a higher-level API similar to that of CUDA C.
LLVM Intrinsics
Load matrix
CUDA.WMMA.llvm_wmma_load — Function
WMMA.llvm_wmma_load_{matrix}_{layout}_{shape}_{addr_space}_stride_{elem_type}(src_addr, stride)
Wrapper around the LLVM intrinsic @llvm.nvvm.wmma.load.{matrix}.sync.{layout}.{shape}.{addr_space}.stride.{elem_type}.
Arguments
  • src_addr: The memory address to load from.
  • stride: The leading dimension of the matrix, in numbers of elements.
Placeholders
  • {matrix}: The matrix to load. Can be a, b or c.
  • {layout}: The storage layout for the matrix. Can be row or col, for row major (C style) or column major (Julia style), respectively.
  • {shape}: The overall shape of the MAC operation. Valid values are m16n16k16, m32n8k16, and m8n32k16.
  • {addr_space}: The address space of src_addr. Can be empty (generic addressing), shared or global.
  • {elem_type}: The type of each element in the matrix. For a and b matrices, valid values are u8 (byte unsigned integer),               s8 (byte signed integer), and f16 (half precision floating point). For c and d matrices, valid values are               s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
source
Perform multiply-accumulate
CUDA.WMMA.llvm_wmma_mma — Function
WMMA.llvm_wmma_mma_{a_layout}_{b_layout}_{shape}_{d_elem_type}_{c_elem_type}(a, b, c) or
-WMMA.llvm_wmma_mma_{a_layout}_{b_layout}_{shape}_{a_elem_type}(a, b, c)
For floating point operations: wrapper around the LLVM intrinsic @llvm.nvvm.wmma.mma.sync.{a_layout}.{b_layout}.{shape}.{d_elem_type}.{c_elem_type} For all other operations: wrapper around the LLVM intrinsic @llvm.nvvm.wmma.mma.sync.{a_layout}.{b_layout}.{shape}.{a_elem_type}
Arguments
  • a: The WMMA fragment corresponding to the matrix $A$.
  • b: The WMMA fragment corresponding to the matrix $B$.
  • c: The WMMA fragment corresponding to the matrix $C$.
Placeholders
  • {a_layout}: The storage layout for matrix $A$. Can be row or col, for row major (C style) or column major (Julia style), respectively. Note that this must match the layout used in the load operation.
  • {b_layout}: The storage layout for matrix $B$. Can be row or col, for row major (C style) or column major (Julia style), respectively. Note that this must match the layout used in the load operation.
  • {shape}: The overall shape of the MAC operation. Valid values are m16n16k16, m32n8k16, and m8n32k16.
  • {a_elem_type}: The type of each element in the $A$ matrix. Valid values are u8 (byte unsigned integer), s8 (byte signed integer), and f16 (half precision floating point).
  • {d_elem_type}: The type of each element in the resultant $D$ matrix. Valid values are s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
  • {c_elem_type}: The type of each element in the $C$ matrix. Valid values are s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
Warning
Remember that the shape, type and layout of all operations (be it MMA, load or store) MUST match. Otherwise, the behaviour is undefined!
source
Store matrix
CUDA.WMMA.llvm_wmma_store — Function
WMMA.llvm_wmma_store_d_{layout}_{shape}_{addr_space}_stride_{elem_type}(dst_addr, data, stride)
Wrapper around the LLVM intrinsic @llvm.nvvm.wmma.store.d.sync.{layout}.{shape}.{addr_space}.stride.{elem_type}.
Arguments
  • dst_addr: The memory address to store to.
  • data: The $D$ fragment to store.
  • stride: The leading dimension of the matrix, in numbers of elements.
Placeholders
  • {layout}: The storage layout for the matrix. Can be row or col, for row major (C style) or column major (Julia style), respectively.
  • {shape}: The overall shape of the MAC operation. Valid values are m16n16k16, m32n8k16, and m8n32k16.
  • {addr_space}: The address space of src_addr. Can be empty (generic addressing), shared or global.
  • {elem_type}: The type of each element in the matrix. For a and b matrices, valid values are u8 (byte unsigned integer),               s8 (byte signed integer), and f16 (half precision floating point). For c and d matrices, valid values are               s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
source
CUDA C-like API
Fragment
CUDA.WMMA.Unspecified — Type
WMMA.Unspecified
Type that represents a matrix stored in an unspecified order.
Warning
This storage format is not valid for all WMMA operations!
source
CUDA.WMMA.Fragment — Type
WMMA.Fragment
Type that represents per-thread intermediate results of WMMA operations.
You can access individual elements using the x member or [] operator, but beware that the exact ordering of elements is unspecified.
source
WMMA configuration
CUDA.WMMA.Config — Type
WMMA.Config{M, N, K, d_type}
Type that contains all information for WMMA operations that cannot be inferred from the argument's types.
WMMA instructions calculate the matrix multiply-accumulate operation $D = A \cdot B + C$, where $A$ is a $M \times K$ matrix, $B$ a $K \times N$ matrix, and $C$ and $D$ are $M \times N$ matrices.
d_type refers to the type of the elements of matrix $D$, and can be either Float16 or Float32.
All WMMA operations take a Config as their final argument.
Examples
julia> config = WMMA.Config{16, 16, 16, Float32}
-CUDA.WMMA.Config{16, 16, 16, Float32}
source
Load matrix
CUDA.WMMA.load_a — Function
WMMA.load_a(addr, stride, layout, config)
+    @cuprint "Hello, World $(42)\n"
source
CUDA.@cuprintf — Macro
@cuprintf("%Fmt", args...)
Print a formatted string in device context on the host standard output.
Note that this is not a fully C-compliant printf implementation; see the CUDA documentation for supported options and inputs.
Also beware that it is an untyped, and unforgiving printf implementation. Type widths need to match, eg. printing a 64-bit Julia integer requires the %ld formatting string.
source
Assertions
CUDA.@cuassert — Macro
@assert cond [text]
Signal assertion failure to the CUDA driver if cond is false. Preferred syntax for writing assertions, mimicking Base.@assert. Message text is optionally displayed upon assertion failure.
Warning
A failed assertion will crash the GPU, so use sparingly as a debugging tool. Furthermore, the assertion might be disabled at various optimization levels, and thus should not cause any side-effects.
source
Atomics
A high-level macro is available to annotate expressions with:
CUDA.@atomic — Macro
@atomic a[I] = op(a[I], val)
+@atomic a[I] ...= val
Atomically perform a sequence of operations that loads an array element a[I], performs the operation op on that value and a second value val, and writes the result back to the array. This sequence can be written out as a regular assignment, in which case the same array element should be used in the left and right hand side of the assignment, or as an in-place application of a known operator. In both cases, the array reference should be pure and not induce any side-effects.
Warn
This interface is experimental, and might change without warning.  Use the lower-level atomic_...! functions for a stable API, albeit one limited to natively-supported ops.
source
If your expression is not recognized, or you need more control, use the underlying functions:
CUDA.atomic_cas! — Function
atomic_cas!(ptr::LLVMPtr{T}, cmp::T, val::T)
Reads the value old located at address ptr and compare with cmp. If old equals to cmp, stores val at the same address. Otherwise, doesn't change the value old. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64. Additionally, on GPU hardware with compute capability 7.0+, values of type UInt16 are supported.
source
CUDA.atomic_xchg! — Function
atomic_xchg!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr and stores val at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_add! — Function
atomic_add!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old + val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32, UInt64, and Float32. Additionally, on GPU hardware with compute capability 6.0+, values of type Float64 are supported.
source
CUDA.atomic_sub! — Function
atomic_sub!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old - val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_and! — Function
atomic_and!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old & val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_or! — Function
atomic_or!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old | val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_xor! — Function
atomic_xor!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes old ⊻ val, and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_min! — Function
atomic_min!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes min(old, val), and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_max! — Function
atomic_max!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes max(old, val), and stores the result back to memory at the same address. These operations are performed in one atomic transaction. The function returns old.
This operation is supported for values of type Int32, Int64, UInt32 and UInt64.
source
CUDA.atomic_inc! — Function
atomic_inc!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes ((old >= val) ? 0 : (old+1)), and stores the result back to memory at the same address. These three operations are performed in one atomic transaction. The function returns old.
This operation is only supported for values of type Int32.
source
CUDA.atomic_dec! — Function
atomic_dec!(ptr::LLVMPtr{T}, val::T)
Reads the value old located at address ptr, computes (((old == 0) | (old > val)) ? val : (old-1) ), and stores the result back to memory at the same address. These three operations are performed in one atomic transaction. The function returns old.
This operation is only supported for values of type Int32.
source
Dynamic parallelism
Similarly to launching kernels from the host, you can use @cuda while passing dynamic=true for launching kernels from the device. A lower-level API is available as well:
CUDA.dynamic_cufunction — Function
dynamic_cufunction(f, tt=Tuple{})
Low-level interface to compile a function invocation for the currently-active GPU, returning a callable kernel object. Device-side equivalent of CUDA.cufunction.
No keyword arguments are supported.
source
CUDA.DeviceKernel — Type
(::HostKernel)(args...; kwargs...)
+(::DeviceKernel)(args...; kwargs...)
Low-level interface to call a compiled kernel, passing GPU-compatible arguments in args. For a higher-level interface, use @cuda.
A HostKernel is callable on the host, and a DeviceKernel is callable on the device (created by @cuda with dynamic=true).
The following keyword arguments are supported:
  • threads (default: 1): Number of threads per block, or a 1-, 2- or 3-tuple of dimensions (e.g. threads=(32, 32) for a 2D block of 32×32 threads). Use threadIdx() and blockDim() to query from within the kernel.
  • blocks (default: 1): Number of thread blocks to launch, or a 1-, 2- or 3-tuple of dimensions (e.g. blocks=(2, 4, 2) for a 3D grid of blocks). Use blockIdx() and gridDim() to query from within the kernel.
  • shmem(default: 0): Amount of dynamic shared memory in bytes to allocate per thread block; used by CuDynamicSharedArray.
  • stream (default: stream()): CuStream to launch the kernel on.
  • cooperative (default: false): whether to launch a cooperative kernel that supports grid synchronization (see CG.this_grid and CG.sync). Note that this requires care wrt. the number of blocks launched.
source
Cooperative groups
CUDA.CG — Module
CUDA.jl's cooperative groups implementation.
Cooperative groups in CUDA offer a structured approach to synchronize and communicate among threads. They allow developers to define specific groups of threads, providing a means to fine-tune inter-thread communication granularity. By offering a more nuanced alternative to traditional CUDA synchronization methods, cooperative groups enable a more controlled and efficient parallel decomposition in kernel design.
The following functionality is available in CUDA.jl:
  • implicit groups: thread blocks, grid groups, and coalesced groups.
  • synchronization: sync, barrier_arrive, barrier_wait
  • warp collectives for coalesced groups: shuffle and voting
  • data transfer: memcpy_async, wait and wait_prior
Noteworthy missing functionality:
  • implicit groups: clusters, and multi-grid groups (which are deprecated)
  • explicit groups: tiling and partitioning
source
Group construction and properties
CUDA.CG.thread_rank — Function
thread_rank(group)
Returns the linearized rank of the calling thread along the interval [1, num_threads()].
source
CUDA.CG.thread_block — Type
thread_block <: thread_group
Every GPU kernel is executed by a grid of thread blocks, and threads within each block are guaranteed to reside on the same streaming multiprocessor. A thread_block represents a thread block whose dimensions are not known until runtime.
Constructed via this_thread_block
source
CUDA.CG.group_index — Function
group_index(tb::thread_block)
3-Dimensional index of the block within the launched grid.
source
CUDA.CG.grid_group — Type
grid_group <: thread_group
Threads within this this group are guaranteed to be co-resident on the same device within the same launched kernel. To use this group, the kernel must have been launched with @cuda cooperative=true, and the device must support it (queryable device attribute).
Constructed via this_grid.
source
CUDA.CG.coalesced_group — Type
coalesced_group <: thread_group
A group representing the current set of converged threads in a warp. The size of the group is not guaranteed and it may return a group of only one thread (itself).
This group exposes warp-synchronous builtins. Constructed via coalesced_threads.
source
CUDA.CG.meta_group_size — Function
meta_group_size(cg::coalesced_group)
Total number of partitions created out of all CTAs when the group was created.
source
Synchronization
Data transfer
CUDA.CG.memcpy_async — Function
memcpy_async(group, dst, src, bytes)
Perform a group-wide collective memory copy from src to dst of bytes bytes. This operation may be performed asynchronously, so you should wait or wait_prior before using the data. It is only supported by thread blocks and coalesced groups.
For this operation to be performed asynchronously, the following conditions must be met:
  • the source and destination memory should be aligned to 4, 8 or 16 bytes. this will be deduced from the datatype, but can also be specified explicitly using CUDA.align.
  • the source should be global memory, and the destination should be shared memory.
  • the device should have compute capability 8.0 or higher.
source
Math
Many mathematical functions are provided by the libdevice library, and are wrapped by CUDA.jl. These functions are used to implement well-known functions from the Julia standard library and packages like SpecialFunctions.jl, e.g., calling the cos function will automatically use __nv_cos from libdevice if possible.
Some functions do not have a counterpart in the Julia ecosystem, those have to be called directly. For example, to call __nv_logb or __nv_logbf you use CUDA.logb in a kernel.
For a list of available functions, look at src/device/intrinsics/math.jl.
WMMA
Warp matrix multiply-accumulate (WMMA) is a CUDA API to access Tensor Cores, a new hardware feature in Volta GPUs to perform mixed precision matrix multiply-accumulate operations. The interface is split in two levels, both available in the WMMA submodule: low level wrappers around the LLVM intrinsics, and a higher-level API similar to that of CUDA C.
LLVM Intrinsics
Load matrix
CUDA.WMMA.llvm_wmma_load — Function
WMMA.llvm_wmma_load_{matrix}_{layout}_{shape}_{addr_space}_stride_{elem_type}(src_addr, stride)
Wrapper around the LLVM intrinsic @llvm.nvvm.wmma.load.{matrix}.sync.{layout}.{shape}.{addr_space}.stride.{elem_type}.
Arguments
  • src_addr: The memory address to load from.
  • stride: The leading dimension of the matrix, in numbers of elements.
Placeholders
  • {matrix}: The matrix to load. Can be a, b or c.
  • {layout}: The storage layout for the matrix. Can be row or col, for row major (C style) or column major (Julia style), respectively.
  • {shape}: The overall shape of the MAC operation. Valid values are m16n16k16, m32n8k16, and m8n32k16.
  • {addr_space}: The address space of src_addr. Can be empty (generic addressing), shared or global.
  • {elem_type}: The type of each element in the matrix. For a and b matrices, valid values are u8 (byte unsigned integer),               s8 (byte signed integer), and f16 (half precision floating point). For c and d matrices, valid values are               s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
source
Perform multiply-accumulate
CUDA.WMMA.llvm_wmma_mma — Function
WMMA.llvm_wmma_mma_{a_layout}_{b_layout}_{shape}_{d_elem_type}_{c_elem_type}(a, b, c) or
+WMMA.llvm_wmma_mma_{a_layout}_{b_layout}_{shape}_{a_elem_type}(a, b, c)
For floating point operations: wrapper around the LLVM intrinsic @llvm.nvvm.wmma.mma.sync.{a_layout}.{b_layout}.{shape}.{d_elem_type}.{c_elem_type} For all other operations: wrapper around the LLVM intrinsic @llvm.nvvm.wmma.mma.sync.{a_layout}.{b_layout}.{shape}.{a_elem_type}
Arguments
  • a: The WMMA fragment corresponding to the matrix $A$.
  • b: The WMMA fragment corresponding to the matrix $B$.
  • c: The WMMA fragment corresponding to the matrix $C$.
Placeholders
  • {a_layout}: The storage layout for matrix $A$. Can be row or col, for row major (C style) or column major (Julia style), respectively. Note that this must match the layout used in the load operation.
  • {b_layout}: The storage layout for matrix $B$. Can be row or col, for row major (C style) or column major (Julia style), respectively. Note that this must match the layout used in the load operation.
  • {shape}: The overall shape of the MAC operation. Valid values are m16n16k16, m32n8k16, and m8n32k16.
  • {a_elem_type}: The type of each element in the $A$ matrix. Valid values are u8 (byte unsigned integer), s8 (byte signed integer), and f16 (half precision floating point).
  • {d_elem_type}: The type of each element in the resultant $D$ matrix. Valid values are s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
  • {c_elem_type}: The type of each element in the $C$ matrix. Valid values are s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
Warning
Remember that the shape, type and layout of all operations (be it MMA, load or store) MUST match. Otherwise, the behaviour is undefined!
source
Store matrix
CUDA.WMMA.llvm_wmma_store — Function
WMMA.llvm_wmma_store_d_{layout}_{shape}_{addr_space}_stride_{elem_type}(dst_addr, data, stride)
Wrapper around the LLVM intrinsic @llvm.nvvm.wmma.store.d.sync.{layout}.{shape}.{addr_space}.stride.{elem_type}.
Arguments
  • dst_addr: The memory address to store to.
  • data: The $D$ fragment to store.
  • stride: The leading dimension of the matrix, in numbers of elements.
Placeholders
  • {layout}: The storage layout for the matrix. Can be row or col, for row major (C style) or column major (Julia style), respectively.
  • {shape}: The overall shape of the MAC operation. Valid values are m16n16k16, m32n8k16, and m8n32k16.
  • {addr_space}: The address space of src_addr. Can be empty (generic addressing), shared or global.
  • {elem_type}: The type of each element in the matrix. For a and b matrices, valid values are u8 (byte unsigned integer),               s8 (byte signed integer), and f16 (half precision floating point). For c and d matrices, valid values are               s32 (32-bit signed integer), f16 (half precision floating point), and f32 (full precision floating point).
source
CUDA C-like API
Fragment
CUDA.WMMA.Unspecified — Type
WMMA.Unspecified
Type that represents a matrix stored in an unspecified order.
Warning
This storage format is not valid for all WMMA operations!
source
CUDA.WMMA.Fragment — Type
WMMA.Fragment
Type that represents per-thread intermediate results of WMMA operations.
You can access individual elements using the x member or [] operator, but beware that the exact ordering of elements is unspecified.
source
WMMA configuration
CUDA.WMMA.Config — Type
WMMA.Config{M, N, K, d_type}
Type that contains all information for WMMA operations that cannot be inferred from the argument's types.
WMMA instructions calculate the matrix multiply-accumulate operation $D = A \cdot B + C$, where $A$ is a $M \times K$ matrix, $B$ a $K \times N$ matrix, and $C$ and $D$ are $M \times N$ matrices.
d_type refers to the type of the elements of matrix $D$, and can be either Float16 or Float32.
All WMMA operations take a Config as their final argument.
Examples
julia> config = WMMA.Config{16, 16, 16, Float32}
+CUDA.WMMA.Config{16, 16, 16, Float32}
source
Load matrix
CUDA.WMMA.load_a — Function
WMMA.load_a(addr, stride, layout, config)
 WMMA.load_b(addr, stride, layout, config)
-WMMA.load_c(addr, stride, layout, config)
Load the matrix a, b or c from the memory location indicated by addr, and return the resulting WMMA.Fragment.
Arguments
  • addr: The address to load the matrix from.
  • stride: The leading dimension of the matrix pointed to by addr, specified in number of elements.
  • layout: The storage layout of the matrix. Possible values are WMMA.RowMajor and WMMA.ColMajor.
  • config: The WMMA configuration that should be used for loading this matrix. See WMMA.Config.
See also: WMMA.Fragment, WMMA.FragmentLayout, WMMA.Config
Warning
All threads in a warp MUST execute the load operation in lockstep, and have to use exactly the same arguments. Failure to do so will result in undefined behaviour.
source
WMMA.load_b and WMMA.load_c have the same signature.
Perform multiply-accumulate
CUDA.WMMA.mma — Function
WMMA.mma(a, b, c, conf)
Perform the matrix multiply-accumulate operation $D = A \cdot B + C$.
Arguments
Warning
All threads in a warp MUST execute the mma operation in lockstep, and have to use exactly the same arguments. Failure to do so will result in undefined behaviour.
source
Store matrix
CUDA.WMMA.store_d — Function
WMMA.store_d(addr, d, stride, layout, config)
Store the result matrix d to the memory location indicated by addr.
Arguments
  • addr: The address to store the matrix to.
  • d: The WMMA.Fragment corresponding to the d matrix.
  • stride: The leading dimension of the matrix pointed to by addr, specified in number of elements.
  • layout: The storage layout of the matrix. Possible values are WMMA.RowMajor and WMMA.ColMajor.
  • config: The WMMA configuration that should be used for storing this matrix. See WMMA.Config.
See also: WMMA.Fragment, WMMA.FragmentLayout, WMMA.Config
Warning
All threads in a warp MUST execute the store operation in lockstep, and have to use exactly the same arguments. Failure to do so will result in undefined behaviour.
source
Fill fragment
CUDA.WMMA.fill_c — Function
WMMA.fill_c(value, config)
Return a WMMA.Fragment filled with the value value.
This operation is useful if you want to implement a matrix multiplication (and thus want to set $C = O$).
Arguments
  • value: The value used to fill the fragment. Can be a Float16 or Float32.
  • config: The WMMA configuration that should be used for this WMMA operation. See WMMA.Config.
source
Other
CUDA.align — Type
CUDA.align{N}(obj)
Construct an aligned object, providing alignment information to APIs that require it.
source

+WMMA.load_c(addr, stride, layout, config)

Load the matrix a, b or c from the memory location indicated by addr, and return the resulting WMMA.Fragment.

Arguments

  • addr: The address to load the matrix from.
  • stride: The leading dimension of the matrix pointed to by addr, specified in number of elements.
  • layout: The storage layout of the matrix. Possible values are WMMA.RowMajor and WMMA.ColMajor.
  • config: The WMMA configuration that should be used for loading this matrix. See WMMA.Config.

See also: WMMA.Fragment, WMMA.FragmentLayout, WMMA.Config

Warning

All threads in a warp MUST execute the load operation in lockstep, and have to use exactly the same arguments. Failure to do so will result in undefined behaviour.

source

WMMA.load_b and WMMA.load_c have the same signature.

Perform multiply-accumulate

CUDA.WMMA.mma — Function
WMMA.mma(a, b, c, conf)

Perform the matrix multiply-accumulate operation $D = A \cdot B + C$.

Arguments

Warning

All threads in a warp MUST execute the mma operation in lockstep, and have to use exactly the same arguments. Failure to do so will result in undefined behaviour.

source

Store matrix

CUDA.WMMA.store_d — Function
WMMA.store_d(addr, d, stride, layout, config)

Store the result matrix d to the memory location indicated by addr.

Arguments

  • addr: The address to store the matrix to.
  • d: The WMMA.Fragment corresponding to the d matrix.
  • stride: The leading dimension of the matrix pointed to by addr, specified in number of elements.
  • layout: The storage layout of the matrix. Possible values are WMMA.RowMajor and WMMA.ColMajor.
  • config: The WMMA configuration that should be used for storing this matrix. See WMMA.Config.

See also: WMMA.Fragment, WMMA.FragmentLayout, WMMA.Config

Warning

All threads in a warp MUST execute the store operation in lockstep, and have to use exactly the same arguments. Failure to do so will result in undefined behaviour.

source

Fill fragment

CUDA.WMMA.fill_c — Function
WMMA.fill_c(value, config)

Return a WMMA.Fragment filled with the value value.

This operation is useful if you want to implement a matrix multiplication (and thus want to set $C = O$).

Arguments

  • value: The value used to fill the fragment. Can be a Float16 or Float32.
  • config: The WMMA configuration that should be used for this WMMA operation. See WMMA.Config.
source

Other

CUDA.align — Type
CUDA.align{N}(obj)

Construct an aligned object, providing alignment information to APIs that require it.

source
diff --git a/dev/development/debugging/index.html b/dev/development/debugging/index.html index f539aa0fa8..9adc8ed8ad 100644 --- a/dev/development/debugging/index.html +++ b/dev/development/debugging/index.html @@ -44,4 +44,4 @@ julia> exit() ========= ERROR SUMMARY: 0 errors -Process(`.julia/artifacts/feb6b469b6047f344fec54df2619d65f6b704bdb/cuda/compute-sanitizer/compute-sanitizer --launch-timeout=0 --target-processes=all --report-api-errors=no julia`, ProcessExited(0))

By default, compute-sanitizer launches the memcheck tool, which is great for dealing with memory issues. Other tools can be selected with the --tool argument, e.g., to find thread synchronization hazards use --tool synccheck, racecheck can be used to find shared memory data races, and initcheck is useful for spotting uses of uninitialized device memory.

cuda-gdb

To debug Julia code, you can use the CUDA debugger cuda-gdb. When using this tool, it is recommended to enable Julia debug mode 2 so that debug information is emitted. Do note that the DWARF info emitted by Julia is currently insufficient to e.g. inspect variables, so the debug experience will not be pleasant.

If you encounter the CUDBG_ERROR_UNINITIALIZED error, ensure all your devices are supported by cuda-gdb (e.g., Kepler-era devices aren't). If some aren't, re-start Julia with CUDA_VISIBLE_DEVICES set to ignore that device.

+Process(`.julia/artifacts/feb6b469b6047f344fec54df2619d65f6b704bdb/cuda/compute-sanitizer/compute-sanitizer --launch-timeout=0 --target-processes=all --report-api-errors=no julia`, ProcessExited(0))

By default, compute-sanitizer launches the memcheck tool, which is great for dealing with memory issues. Other tools can be selected with the --tool argument, e.g., to find thread synchronization hazards use --tool synccheck, racecheck can be used to find shared memory data races, and initcheck is useful for spotting uses of uninitialized device memory.

cuda-gdb

To debug Julia code, you can use the CUDA debugger cuda-gdb. When using this tool, it is recommended to enable Julia debug mode 2 so that debug information is emitted. Do note that the DWARF info emitted by Julia is currently insufficient to e.g. inspect variables, so the debug experience will not be pleasant.

If you encounter the CUDBG_ERROR_UNINITIALIZED error, ensure all your devices are supported by cuda-gdb (e.g., Kepler-era devices aren't). If some aren't, re-start Julia with CUDA_VISIBLE_DEVICES set to ignore that device.

diff --git a/dev/development/kernel/index.html b/dev/development/kernel/index.html index cd76e3aba4..87b63caadf 100644 --- a/dev/development/kernel/index.html +++ b/dev/development/kernel/index.html @@ -191,4 +191,4 @@ 1 2 3 - 4

The above example waits for the copy to complete before continuing, but it is also possible to have multiple copies in flight using the CG.wait_prior function, which waits for all but the last N copies to complete.

Warp matrix multiply-accumulate

Warp matrix multiply-accumulate (WMMA) is a cooperative operation to perform mixed precision matrix multiply-accumulate on the tensor core hardware of recent GPUs. The CUDA.jl interface is split in two levels, both available in the WMMA submodule: low level wrappers around the LLVM intrinsics, and a higher-level API similar to that of CUDA C.

Terminology

The WMMA operations perform a matrix multiply-accumulate. More concretely, it calculates $D = A \cdot B + C$, where $A$ is a $M \times K$ matrix, $B$ is a $K \times N$ matrix, and $C$ and $D$ are $M \times N$ matrices.

However, not all values of $M$, $N$ and $K$ are allowed. The tuple $(M, N, K)$ is often called the "shape" of the multiply accumulate operation.

The multiply-accumulate consists of the following steps:

Note that WMMA is a warp-wide operation, which means that all threads in a warp must cooperate, and execute the WMMA operations in lockstep. Failure to do so will result in undefined behaviour.

Each thread in a warp will hold a part of the matrix in its registers. In WMMA parlance, this part is referred to as a "fragment". Note that the exact mapping between matrix elements and fragment is unspecified, and subject to change in future versions.

Finally, it is important to note that the resultant $D$ matrix can be used as a $C$ matrix for a subsequent multiply-accumulate. This is useful if one needs to calculate a sum of the form $\sum_{i=0}^{n} A_i B_i$, where $A_i$ and $B_i$ are matrices of the correct dimension.

LLVM Intrinsics

The LLVM intrinsics are accessible by using the one-to-one Julia wrappers. The return type of each wrapper is the Julia type that corresponds closest to the return type of the LLVM intrinsic. For example, LLVM's [8 x <2 x half>] becomes NTuple{8, NTuple{2, VecElement{Float16}}} in Julia. In essence, these wrappers return the SSA values returned by the LLVM intrinsic. Currently, all intrinsics that are available in LLVM 6, PTX 6.0 and SM 70 are implemented.

These LLVM intrinsics are then lowered to the correct PTX instructions by the LLVM NVPTX backend. For more information about the PTX instructions, please refer to the PTX Instruction Set Architecture Manual.

The LLVM intrinsics are subdivided in three categories:

CUDA C-like API

The main difference between the CUDA C-like API and the lower level wrappers, is that the former enforces several constraints when working with WMMA. For example, it ensures that the $A$ fragment argument to the MMA instruction was obtained by a load_a call, and not by a load_b or load_c. Additionally, it makes sure that the data type and storage layout of the load/store operations and the MMA operation match.

The CUDA C-like API heavily uses Julia's dispatch mechanism. As such, the method names are much shorter than the LLVM intrinsic wrappers, as most information is baked into the type of the arguments rather than the method name.

Note that, in CUDA C++, the fragment is responsible for both the storage of intermediate results and the WMMA configuration. All CUDA C++ WMMA calls are function templates that take the resultant fragment as a by-reference argument. As a result, the type of this argument can be used during overload resolution to select the correct WMMA instruction to call.

In contrast, the API in Julia separates the WMMA storage (WMMA.Fragment) and configuration (WMMA.Config). Instead of taking the resultant fragment by reference, the Julia functions just return it. This makes the dataflow clearer, but it also means that the type of that fragment cannot be used for selection of the correct WMMA instruction. Thus, there is still a limited amount of information that cannot be inferred from the argument types, but must nonetheless match for all WMMA operations, such as the overall shape of the MMA. This is accomplished by a separate "WMMA configuration" (see WMMA.Config) that you create once, and then give as an argument to all intrinsics.

Element access and broadcasting

Similar to the CUDA C++ WMMA API, WMMA.Fragments have an x member that can be used to access individual elements. Note that, in contrast to the values returned by the LLVM intrinsics, the x member is flattened. For example, while the Float16 variants of the load_a instrinsics return NTuple{8, NTuple{2, VecElement{Float16}}}, the x member has type NTuple{16, Float16}.

Typically, you will only need to access the x member to perform elementwise operations. This can be more succinctly expressed using Julia's broadcast mechanism. For example, to double each element in a fragment, you can simply use:

frag = 2.0f0 .* frag
+ 4

The above example waits for the copy to complete before continuing, but it is also possible to have multiple copies in flight using the CG.wait_prior function, which waits for all but the last N copies to complete.

Warp matrix multiply-accumulate

Warp matrix multiply-accumulate (WMMA) is a cooperative operation to perform mixed precision matrix multiply-accumulate on the tensor core hardware of recent GPUs. The CUDA.jl interface is split in two levels, both available in the WMMA submodule: low level wrappers around the LLVM intrinsics, and a higher-level API similar to that of CUDA C.

Terminology

The WMMA operations perform a matrix multiply-accumulate. More concretely, it calculates $D = A \cdot B + C$, where $A$ is a $M \times K$ matrix, $B$ is a $K \times N$ matrix, and $C$ and $D$ are $M \times N$ matrices.

However, not all values of $M$, $N$ and $K$ are allowed. The tuple $(M, N, K)$ is often called the "shape" of the multiply accumulate operation.

The multiply-accumulate consists of the following steps:

Note that WMMA is a warp-wide operation, which means that all threads in a warp must cooperate, and execute the WMMA operations in lockstep. Failure to do so will result in undefined behaviour.

Each thread in a warp will hold a part of the matrix in its registers. In WMMA parlance, this part is referred to as a "fragment". Note that the exact mapping between matrix elements and fragment is unspecified, and subject to change in future versions.

Finally, it is important to note that the resultant $D$ matrix can be used as a $C$ matrix for a subsequent multiply-accumulate. This is useful if one needs to calculate a sum of the form $\sum_{i=0}^{n} A_i B_i$, where $A_i$ and $B_i$ are matrices of the correct dimension.

LLVM Intrinsics

The LLVM intrinsics are accessible by using the one-to-one Julia wrappers. The return type of each wrapper is the Julia type that corresponds closest to the return type of the LLVM intrinsic. For example, LLVM's [8 x <2 x half>] becomes NTuple{8, NTuple{2, VecElement{Float16}}} in Julia. In essence, these wrappers return the SSA values returned by the LLVM intrinsic. Currently, all intrinsics that are available in LLVM 6, PTX 6.0 and SM 70 are implemented.

These LLVM intrinsics are then lowered to the correct PTX instructions by the LLVM NVPTX backend. For more information about the PTX instructions, please refer to the PTX Instruction Set Architecture Manual.

The LLVM intrinsics are subdivided in three categories:

CUDA C-like API

The main difference between the CUDA C-like API and the lower level wrappers, is that the former enforces several constraints when working with WMMA. For example, it ensures that the $A$ fragment argument to the MMA instruction was obtained by a load_a call, and not by a load_b or load_c. Additionally, it makes sure that the data type and storage layout of the load/store operations and the MMA operation match.

The CUDA C-like API heavily uses Julia's dispatch mechanism. As such, the method names are much shorter than the LLVM intrinsic wrappers, as most information is baked into the type of the arguments rather than the method name.

Note that, in CUDA C++, the fragment is responsible for both the storage of intermediate results and the WMMA configuration. All CUDA C++ WMMA calls are function templates that take the resultant fragment as a by-reference argument. As a result, the type of this argument can be used during overload resolution to select the correct WMMA instruction to call.

In contrast, the API in Julia separates the WMMA storage (WMMA.Fragment) and configuration (WMMA.Config). Instead of taking the resultant fragment by reference, the Julia functions just return it. This makes the dataflow clearer, but it also means that the type of that fragment cannot be used for selection of the correct WMMA instruction. Thus, there is still a limited amount of information that cannot be inferred from the argument types, but must nonetheless match for all WMMA operations, such as the overall shape of the MMA. This is accomplished by a separate "WMMA configuration" (see WMMA.Config) that you create once, and then give as an argument to all intrinsics.

Element access and broadcasting

Similar to the CUDA C++ WMMA API, WMMA.Fragments have an x member that can be used to access individual elements. Note that, in contrast to the values returned by the LLVM intrinsics, the x member is flattened. For example, while the Float16 variants of the load_a instrinsics return NTuple{8, NTuple{2, VecElement{Float16}}}, the x member has type NTuple{16, Float16}.

Typically, you will only need to access the x member to perform elementwise operations. This can be more succinctly expressed using Julia's broadcast mechanism. For example, to double each element in a fragment, you can simply use:

frag = 2.0f0 .* frag
diff --git a/dev/development/profiling/index.html b/dev/development/profiling/index.html index 76396ee350..e7b669831d 100644 --- a/dev/development/profiling/index.html +++ b/dev/development/profiling/index.html @@ -96,4 +96,4 @@ NVTX.@annotate function foo() ... -end

For more details, refer to the documentation of the NVTX.jl package.

Compiler options

Some tools, like NSight Systems Compute, also make it possible to do source-level profiling. CUDA.jl will by default emit the necessary source line information, which you can disable by launching Julia with -g0. Conversely, launching with -g2 will emit additional debug information, which can be useful in combination with tools like cuda-gdb, but might hurt performance or code size.

+end

For more details, refer to the documentation of the NVTX.jl package.

Compiler options

Some tools, like NSight Systems Compute, also make it possible to do source-level profiling. CUDA.jl will by default emit the necessary source line information, which you can disable by launching Julia with -g0. Conversely, launching with -g2 will emit additional debug information, which can be useful in combination with tools like cuda-gdb, but might hurt performance or code size.

diff --git a/dev/development/troubleshooting/index.html b/dev/development/troubleshooting/index.html index c7953d0696..5dcdd7344f 100644 --- a/dev/development/troubleshooting/index.html +++ b/dev/development/troubleshooting/index.html @@ -45,4 +45,4 @@ • %17 = call CUDA.sin(::Int64)::Union{}

Both from the IR and the list of calls Cthulhu offers to inspect further, we can see that the call to CUDA.sin(::Int64) results in an error: in the IR it is immediately followed by an unreachable, while in the list of calls it is inferred to return Union{}. Now we know where to look, it's easy to figure out what's wrong:

help?> CUDA.sin
   # 2 methods for generic function "sin":
   [1] sin(x::Float32) in CUDA at /home/tim/Julia/pkg/CUDA/src/device/intrinsics/math.jl:13
-  [2] sin(x::Float64) in CUDA at /home/tim/Julia/pkg/CUDA/src/device/intrinsics/math.jl:12

There's no method of CUDA.sin that accepts an Int64, and thus the function was determined to unconditionally throw a method error. For now, we disallow these situations and refuse to compile, but in the spirit of dynamic languages we might change this behavior to just throw an error at run time.

+ [2] sin(x::Float64) in CUDA at /home/tim/Julia/pkg/CUDA/src/device/intrinsics/math.jl:12

There's no method of CUDA.sin that accepts an Int64, and thus the function was determined to unconditionally throw a method error. For now, we disallow these situations and refuse to compile, but in the spirit of dynamic languages we might change this behavior to just throw an error at run time.

diff --git a/dev/faq/index.html b/dev/faq/index.html index c2ff706f10..4215ab8173 100644 --- a/dev/faq/index.html +++ b/dev/faq/index.html @@ -20,4 +20,4 @@ ├─possible versions are: [0.4.1, 0.5.0-0.5.4, 0.6.0-0.6.10, 0.7.0-0.7.3, 0.8.0-0.8.3, 0.9.0, 0.10.0-0.10.4, 0.11.0-0.11.1] or uninstalled ├─restricted to versions * by an explicit requirement, leaving only versions [0.4.1, 0.5.0-0.5.4, 0.6.0-0.6.10, 0.7.0-0.7.3, 0.8.0-0.8.3, 0.9.0, 0.10.0-0.10.4, 0.11.0-0.11.1] └─restricted by compatibility requirements with CUDA [052768ef] to versions: [0.4.1, 0.5.0-0.5.4, 0.6.0-0.6.10, 0.7.0-0.7.3, 0.8.0-0.8.3, 0.9.0, 0.10.0-0.10.4] or uninstalled, leaving only versions: [0.4.1, 0.5.0-0.5.4, 0.6.0-0.6.10, 0.7.0-0.7.3, 0.8.0-0.8.3, 0.9.0, 0.10.0-0.10.4] - └─CUDA [052768ef] log: see above

A common source of these incompatibilities is having both CUDA.jl and the older CUDAnative.jl/CuArrays.jl/CUDAdrv.jl stack installed: These are incompatible, and cannot coexist. You can inspect in the Pkg REPL which exact packages you have installed using the status --manifest option.

Can you wrap this or that CUDA API?

If a certain API isn't wrapped with some high-level functionality, you can always use the underlying C APIs which are always available as unexported methods. For example, you can access the CUDA driver library as cu prefixed, unexported functions like CUDA.cuDriverGetVersion. Similarly, vendor libraries like CUBLAS are available through their exported submodule handles, e.g., CUBLAS.cublasGetVersion_v2.

Any help on designing or implementing high-level wrappers for this low-level functionality is greatly appreciated, so please consider contributing your uses of these APIs on the respective repositories.

When installing CUDA.jl on a cluster, why does Julia stall during precompilation?

If you're working on a cluster, precompilation may stall if you have not requested sufficient memory. You may also wish to make sure you have enough disk space prior to installing CUDA.jl.

+ └─CUDA [052768ef] log: see above

A common source of these incompatibilities is having both CUDA.jl and the older CUDAnative.jl/CuArrays.jl/CUDAdrv.jl stack installed: These are incompatible, and cannot coexist. You can inspect in the Pkg REPL which exact packages you have installed using the status --manifest option.

Can you wrap this or that CUDA API?

If a certain API isn't wrapped with some high-level functionality, you can always use the underlying C APIs which are always available as unexported methods. For example, you can access the CUDA driver library as cu prefixed, unexported functions like CUDA.cuDriverGetVersion. Similarly, vendor libraries like CUBLAS are available through their exported submodule handles, e.g., CUBLAS.cublasGetVersion_v2.

Any help on designing or implementing high-level wrappers for this low-level functionality is greatly appreciated, so please consider contributing your uses of these APIs on the respective repositories.

When installing CUDA.jl on a cluster, why does Julia stall during precompilation?

If you're working on a cluster, precompilation may stall if you have not requested sufficient memory. You may also wish to make sure you have enough disk space prior to installing CUDA.jl.

diff --git a/dev/index.html b/dev/index.html index bb1a529425..c5aaf89fff 100644 --- a/dev/index.html +++ b/dev/index.html @@ -13,4 +13,4 @@ Pkg.test("CUDA") # the test suite takes command-line options that allow customization; pass --help for details: -#Pkg.test("CUDA"; test_args=`--help`)

For more details on the installation process, consult the Installation section. To understand the toolchain in more detail, have a look at the tutorials in this manual. It is highly recommended that new users start with the Introduction tutorial. For an overview of the available functionality, read the Usage section. The following resources may also be of interest:

Acknowledgements

The Julia CUDA stack has been a collaborative effort by many individuals. Significant contributions have been made by the following individuals:

Supporting and Citing

Much of the software in this ecosystem was developed as part of academic research. If you would like to help support it, please star the repository as such metrics may help us secure funding in the future. If you use our software as part of your research, teaching, or other activities, we would be grateful if you could cite our work. The CITATION.bib file in the root of this repository lists the relevant papers.

+#Pkg.test("CUDA"; test_args=`--help`)

For more details on the installation process, consult the Installation section. To understand the toolchain in more detail, have a look at the tutorials in this manual. It is highly recommended that new users start with the Introduction tutorial. For an overview of the available functionality, read the Usage section. The following resources may also be of interest:

Acknowledgements

The Julia CUDA stack has been a collaborative effort by many individuals. Significant contributions have been made by the following individuals:

Supporting and Citing

Much of the software in this ecosystem was developed as part of academic research. If you would like to help support it, please star the repository as such metrics may help us secure funding in the future. If you use our software as part of your research, teaching, or other activities, we would be grateful if you could cite our work. The CITATION.bib file in the root of this repository lists the relevant papers.

diff --git a/dev/installation/conditional/index.html b/dev/installation/conditional/index.html index bab61e2b10..553295f70a 100644 --- a/dev/installation/conditional/index.html +++ b/dev/installation/conditional/index.html @@ -33,4 +33,4 @@ function __init__() use_gpu[] = CUDA.functional() -end

The disadvantage of this approach is the introduction of a type instability.

+end

The disadvantage of this approach is the introduction of a type instability.

diff --git a/dev/installation/overview/index.html b/dev/installation/overview/index.html index b83d61bc18..b4f6504cd7 100644 --- a/dev/installation/overview/index.html +++ b/dev/installation/overview/index.html @@ -27,4 +27,4 @@ julia> CUDA.versioninfo() CUDA runtime 11.8, local installation ...

Calling the above helper function generates the following LocalPreferences.toml file in your active environment:

[CUDA_Runtime_jll]
-local = "true"

This preference not only configures CUDA.jl to use a local toolkit, it also prevents downloading any artifact, so it may be interesting to set this preference before ever importing CUDA.jl (e.g., by putting this preference file in a system-wide depot).

If CUDA.jl doesn't properly detect your local toolkit, it may be that certain libraries or binaries aren't on a globally-discoverable path. For more information, run Julia with the JULIA_DEBUG environment variable set to CUDA_Runtime_Discovery.

Note that using a local toolkit instead of artifacts any CUDA-related JLL, not just of CUDA_Runtime_jll. Any package that depends on such a JLL needs to inspect CUDA.local_toolkit, and if set use CUDA_Runtime_Discovery to detect libraries and binaries instead.

Precompiling CUDA.jl without CUDA

CUDA.jl can be precompiled and imported on systems without a GPU or CUDA installation. This simplifies the situation where an application optionally uses CUDA. However, when CUDA.jl is precompiled in such an environment, it cannot be used to run GPU code. This is a result of artifacts being selected at precompile time.

In some cases, e.g. with containers or HPC log-in nodes, you may want to precompile CUDA.jl on a system without CUDA, yet still want to have it download the necessary artifacts and/or produce a precompilation image that can be used on a system with CUDA. This can be achieved by informing CUDA.jl which CUDA toolkit to run time by calling CUDA.set_runtime_version!.

When using artifacts, that's as simple as e.g. calling CUDA.set_runtime_version!(v"11.8"), and afterwards re-starting Julia and re-importing CUDA.jl in order to trigger precompilation again and download the necessary artifacts. If you want to use a local CUDA installation, you also need to set the local_toolkit keyword argument, e.g., by calling CUDA.set_runtime_version!(v"11.8"; local_toolkit=true). Note that the version specified here needs to match what will be available at run time. In both cases, i.e. when using artifacts or a local toolkit, the chosen version needs to be compatible with the available driver.

Finally, in such a scenario you may also want to call CUDA.precompile_runtime() to ensure that the GPUCompiler runtime library is precompiled as well. This and all of the above is demonstrated in the Dockerfile that's part of the CUDA.jl repository.

+local = "true"

This preference not only configures CUDA.jl to use a local toolkit, it also prevents downloading any artifact, so it may be interesting to set this preference before ever importing CUDA.jl (e.g., by putting this preference file in a system-wide depot).

If CUDA.jl doesn't properly detect your local toolkit, it may be that certain libraries or binaries aren't on a globally-discoverable path. For more information, run Julia with the JULIA_DEBUG environment variable set to CUDA_Runtime_Discovery.

Note that using a local toolkit instead of artifacts any CUDA-related JLL, not just of CUDA_Runtime_jll. Any package that depends on such a JLL needs to inspect CUDA.local_toolkit, and if set use CUDA_Runtime_Discovery to detect libraries and binaries instead.

Precompiling CUDA.jl without CUDA

CUDA.jl can be precompiled and imported on systems without a GPU or CUDA installation. This simplifies the situation where an application optionally uses CUDA. However, when CUDA.jl is precompiled in such an environment, it cannot be used to run GPU code. This is a result of artifacts being selected at precompile time.

In some cases, e.g. with containers or HPC log-in nodes, you may want to precompile CUDA.jl on a system without CUDA, yet still want to have it download the necessary artifacts and/or produce a precompilation image that can be used on a system with CUDA. This can be achieved by informing CUDA.jl which CUDA toolkit to run time by calling CUDA.set_runtime_version!.

When using artifacts, that's as simple as e.g. calling CUDA.set_runtime_version!(v"11.8"), and afterwards re-starting Julia and re-importing CUDA.jl in order to trigger precompilation again and download the necessary artifacts. If you want to use a local CUDA installation, you also need to set the local_toolkit keyword argument, e.g., by calling CUDA.set_runtime_version!(v"11.8"; local_toolkit=true). Note that the version specified here needs to match what will be available at run time. In both cases, i.e. when using artifacts or a local toolkit, the chosen version needs to be compatible with the available driver.

Finally, in such a scenario you may also want to call CUDA.precompile_runtime() to ensure that the GPUCompiler runtime library is precompiled as well. This and all of the above is demonstrated in the Dockerfile that's part of the CUDA.jl repository.

diff --git a/dev/installation/troubleshooting/index.html b/dev/installation/troubleshooting/index.html index 4a7fff20cc..6107664c91 100644 --- a/dev/installation/troubleshooting/index.html +++ b/dev/installation/troubleshooting/index.html @@ -3,4 +3,4 @@ function gtag(){dataLayer.push(arguments);} gtag('js', new Date()); gtag('config', 'UA-154489943-2', {'page_path': location.pathname + location.search + location.hash}); -

Troubleshooting

UndefVarError: libcuda not defined

This means that CUDA.jl could not find a suitable CUDA driver. For more information, re-run with the JULIA_DEBUG environment variable set to CUDA_Driver_jll.

UNKNOWN_ERROR(999)

If you encounter this error, there are several known issues that may be causing it:

  • a mismatch between the CUDA driver and driver library: on Linux, look for clues in dmesg
  • the CUDA driver is in a bad state: this can happen after resume. Try rebooting.

Generally though, it's impossible to say what's the reason for the error, but Julia is likely not to blame. Make sure your set-up works (e.g., try executing nvidia-smi, a CUDA C binary, etc), and if everything looks good file an issue.

NVML library not found (on Windows)

Check and make sure the NVSMI folder is in your PATH. By default it may not be. Look in C:\Program Files\NVIDIA Corporation for the NVSMI folder - you should see nvml.dll within it. You can add this folder to your PATH and check that nvidia-smi runs properly.

The specified module could not be found (on Windows)

Ensure the Visual C++ Redistributable is installed.

+

Troubleshooting

UndefVarError: libcuda not defined

This means that CUDA.jl could not find a suitable CUDA driver. For more information, re-run with the JULIA_DEBUG environment variable set to CUDA_Driver_jll.

UNKNOWN_ERROR(999)

If you encounter this error, there are several known issues that may be causing it:

  • a mismatch between the CUDA driver and driver library: on Linux, look for clues in dmesg
  • the CUDA driver is in a bad state: this can happen after resume. Try rebooting.

Generally though, it's impossible to say what's the reason for the error, but Julia is likely not to blame. Make sure your set-up works (e.g., try executing nvidia-smi, a CUDA C binary, etc), and if everything looks good file an issue.

NVML library not found (on Windows)

Check and make sure the NVSMI folder is in your PATH. By default it may not be. Look in C:\Program Files\NVIDIA Corporation for the NVSMI folder - you should see nvml.dll within it. You can add this folder to your PATH and check that nvidia-smi runs properly.

The specified module could not be found (on Windows)

Ensure the Visual C++ Redistributable is installed.

diff --git a/dev/lib/driver/index.html b/dev/lib/driver/index.html index d38493cf26..324d4ceea0 100644 --- a/dev/lib/driver/index.html +++ b/dev/lib/driver/index.html @@ -3,24 +3,24 @@ function gtag(){dataLayer.push(arguments);} gtag('js', new Date()); gtag('config', 'UA-154489943-2', {'page_path': location.pathname + location.search + location.hash}); -

CUDA driver

This section lists the package's public functionality that directly corresponds to functionality of the CUDA driver API. In general, the abstractions stay close to those of the CUDA driver API, so for more information on certain library calls you can consult the CUDA driver API reference.

The documentation is grouped according to the modules of the driver API.

Error Handling

CUDA.name — Method
name(err::CuError)

Gets the string representation of an error code.

julia> err = CuError(CUDA.cudaError_enum(1))
+
CUDA driver
This section lists the package's public functionality that directly corresponds to functionality of the CUDA driver API. In general, the abstractions stay close to those of the CUDA driver API, so for more information on certain library calls you can consult the CUDA driver API reference.
The documentation is grouped according to the modules of the driver API.
Error Handling
CUDA.name — Method
name(err::CuError)
Gets the string representation of an error code.
julia> err = CuError(CUDA.cudaError_enum(1))
 CuError(CUDA_ERROR_INVALID_VALUE)
 
 julia> name(err)
-"ERROR_INVALID_VALUE"
source
Version Management
CUDA.set_runtime_version! — Function
CUDA.set_runtime_version!([version::VersionNumber]; [local_toolkit::Bool])
Configures the active project to use a specific CUDA toolkit version from a specific source.
If local_toolkit is set, the CUDA toolkit will be used from the local system, otherwise it will be downloaded from an artifact source. In the case of a local toolkit, version informs CUDA.jl which version that is (this may be useful if auto-detection fails). In the case of artifact sources, version controls which version will be downloaded and used.
When not specifying either the version or the local_toolkit argument, the default behavior will be used, which is to use the most recent compatible runtime available from an artifact source. Note that this will override any Preferences that may be configured in a higher-up depot; to clear preferences nondestructively, use CUDA.reset_runtime_version! instead.
source
CUDA.reset_runtime_version! — Function
CUDA.reset_runtime_version!()
Resets the CUDA version preferences in the active project to the default, which is to use the most recent compatible runtime available from an artifact source, unless a higher-up depot has configured a different preference. To force use of the default behavior for the local project, use CUDA.set_runtime_version! with no arguments.
source
Device Management
CUDA.current_device — Function
current_device()
Returns the current device.
Warning
This is a low-level API, returning the current device as known to the CUDA driver. For most users, it is recommended to use the device method instead.
source
CUDA.name — Method
name(dev::CuDevice)
Returns an identifier string for the device.
source
CUDA.totalmem — Method
totalmem(dev::CuDevice)
Returns the total amount of memory (in bytes) on the device.
source
CUDA.attribute — Function
attribute(dev::CuDevice, code)
Returns information about the device.
source
attribute(X, pool::CuMemoryPool, attr)
Returns attribute attr about pool. The type of the returned value depends on the attribute, and as such must be passed as the X parameter.
source
attribute(X, ptr::Union{Ptr,CuPtr}, attr)
Returns attribute attr about pointer ptr. The type of the returned value depends on the attribute, and as such must be passed as the X parameter.
source
Certain common attributes are exposed by additional convenience functions:
CUDA.warpsize — Method
warpsize(dev::CuDevice)
Returns the warp size (in threads) of the device.
source
Context Management
CUDA.CuContext — Type
CuContext(dev::CuDevice, flags=CTX_SCHED_AUTO)
-CuContext(f::Function, ...)
Create a CUDA context for device. A context on the GPU is analogous to a process on the CPU, with its own distinct address space and allocated resources. When a context is destroyed, the system cleans up the resources allocated to it.
When you are done using the context, call CUDA.unsafe_destroy! to mark it for deletion, or use do-block syntax with this constructor.
source
CUDA.unsafe_destroy! — Method
unsafe_destroy!(ctx::CuContext)
Immediately destroy a context, freeing up all resources associated with it. This does not respect any users of the context, and might make other objects unusable.
source
CUDA.current_context — Function
current_context()
Returns the current context. Throws an undefined reference error if the current thread has no context bound to it, or if the bound context has been destroyed.
Warning
This is a low-level API, returning the current context as known to the CUDA driver. For most users, it is recommended to use the context method instead.
source
CUDA.activate — Method
activate(ctx::CuContext)
Binds the specified CUDA context to the calling CPU thread.
source
CUDA.synchronize — Method
synchronize(ctx::Context)
Block for the all operations on ctx to complete. This is a heavyweight operation, typically you only need to call synchronize which only synchronizes the stream associated with the current task.
source
CUDA.device_synchronize — Function
device_synchronize()
Block for the all operations on ctx to complete. This is a heavyweight operation, typically you only need to call synchronize which only synchronizes the stream associated with the current task.
On the device, device_synchronize acts as a synchronization point for child grids in the context of dynamic parallelism.
source
Primary Context Management
CUDA.CuPrimaryContext — Type
CuPrimaryContext(dev::CuDevice)
Create a primary CUDA context for a given device.
Each primary context is unique per device and is shared with CUDA runtime API. It is meant for interoperability with (applications using) the runtime API.
source
CUDA.CuContext — Method
CuContext(pctx::CuPrimaryContext)
Derive a context from a primary context.
Calling this function increases the reference count of the primary context. The returned context should not be free with the unsafe_destroy! function that's used with ordinary contexts. Instead, the refcount of the primary context should be decreased by calling unsafe_release!, or set to zero by calling unsafe_reset!. The easiest way to do this is by using the do-block syntax.
source
CUDA.isactive — Method
isactive(pctx::CuPrimaryContext)
Query whether a primary context is active.
source
CUDA.flags — Method
flags(pctx::CuPrimaryContext)
Query the flags of a primary context.
source
CUDA.unsafe_reset! — Method
unsafe_reset!(pctx::CuPrimaryContext)
Explicitly destroys and cleans up all resources associated with a device's primary context in the current process. Note that this forcibly invalidates all contexts derived from this primary context, and as a result outstanding resources might become invalid.
source
CUDA.unsafe_release! — Method
CUDA.unsafe_release!(pctx::CuPrimaryContext)
Lower the refcount of a context, possibly freeing up all resources associated with it. This does not respect any users of the context, and might make other objects unusable.
source
Module Management
CUDA.CuModule — Type
CuModule(data, options::Dict{CUjit_option,Any})
-CuModuleFile(path, options::Dict{CUjit_option,Any})
Create a CUDA module from a data, or a file containing data. The data may be PTX code, a CUBIN, or a FATBIN.
The options is an optional dictionary of JIT options and their respective value.
source
Function Management
CUDA.CuFunction — Type
CuFunction(mod::CuModule, name::String)
Acquires a function handle from a named function in a module.
source
Global Variable Management
CUDA.CuGlobal — Type
CuGlobal{T}(mod::CuModule, name::String)
Acquires a typed global variable handle from a named global in a module.
source
Base.eltype — Method
eltype(var::CuGlobal)
Return the element type of a global variable object.
source
Base.getindex — Method
Base.getindex(var::CuGlobal)
Return the current value of a global variable.
source
Base.setindex! — Method
Base.setindex(var::CuGlobal{T}, val::T)
Set the value of a global variable to val
source
Linker
CUDA.add_data! — Function
add_data!(link::CuLink, name::String, code::String)
Add PTX code to a pending link operation.
source
add_data!(link::CuLink, name::String, data::Vector{UInt8})
Add object code to a pending link operation.
source
CUDA.add_file! — Function
add_file!(link::CuLink, path::String, typ::CUjitInputType)
Add data from a file to a link operation. The argument typ indicates the type of the contained data.
source
CUDA.CuLinkImage — Type
The result of a linking operation.
This object keeps its parent linker object alive, as destroying a linker destroys linked images too.
source
CUDA.complete — Function
complete(link::CuLink)
Complete a pending linker invocation, returning an output image.
source
CUDA.CuModule — Method
CuModule(img::CuLinkImage, ...)
Create a CUDA module from a completed linking operation. Options from CuModule apply.
source
Memory Management
Different kinds of memory objects can be created, representing different kinds of memory that the CUDA toolkit supports. Each of these memory objects can be allocated by calling alloc with the type of memory as first argument, and freed by calling free. Certain kinds of memory have specific methods defined.
Device memory
This memory is accessible only by the GPU, and is the most common kind of memory used in CUDA programming.
CUDA.alloc — Method
alloc(DeviceMemory, bytesize::Integer;
-      [async=false], [stream::CuStream], [pool::CuMemoryPool])
Allocate bytesize bytes of memory on the device. This memory is only accessible on the GPU, and requires explicit calls to unsafe_copyto!, which wraps cuMemcpy, for access on the CPU.
source
Unified memory
Unified memory is accessible by both the CPU and the GPU, and is managed by the CUDA runtime. It is automatically migrated between the CPU and the GPU as needed, which simplifies programming but can lead to performance issues if not used carefully.
CUDA.alloc — Method
alloc(UnifiedMemory, bytesize::Integer, [flags::CUmemAttach_flags])
Allocate bytesize bytes of unified memory. This memory is accessible from both the CPU and GPU, with the CUDA driver automatically copying upon first access.
source
CUDA.prefetch — Method
prefetch(::UnifiedMemory, [bytes::Integer]; [device::CuDevice], [stream::CuStream])
Prefetches memory to the specified destination device.
source
CUDA.advise — Method
advise(::UnifiedMemory, advice::CUDA.CUmem_advise, [bytes::Integer]; [device::CuDevice])
Advise about the usage of a given memory range.
source
Host memory
Host memory resides on the CPU, but is accessible by the GPU via the PCI bus. This is the slowest kind of memory, but is useful for communicating between running kernels and the host (e.g., to update counters or flags).
CUDA.HostMemory — Type
HostMemory
Pinned memory residing on the CPU, possibly accessible on the GPU.
source
CUDA.alloc — Method
alloc(HostMemory, bytesize::Integer, [flags])
Allocate bytesize bytes of page-locked memory on the host. This memory is accessible from the CPU, and makes it possible to perform faster memory copies to the GPU. Furthermore, if flags is set to MEMHOSTALLOC_DEVICEMAP the memory is also accessible from the GPU. These accesses are direct, and go through the PCI bus. If flags is set to MEMHOSTALLOC_PORTABLE, the memory is considered mapped by all CUDA contexts, not just the one that created the memory, which is useful if the memory needs to be accessed from multiple devices. Multiple flags can be set at one time using a bytewise OR:
flags = MEMHOSTALLOC_PORTABLE | MEMHOSTALLOC_DEVICEMAP
source
CUDA.register — Method
register(HostMemory, ptr::Ptr, bytesize::Integer, [flags])
Page-lock the host memory pointed to by ptr. Subsequent transfers to and from devices will be faster, and can be executed asynchronously. If the MEMHOSTREGISTER_DEVICEMAP flag is specified, the buffer will also be accessible directly from the GPU. These accesses are direct, and go through the PCI bus. If the MEMHOSTREGISTER_PORTABLE flag is specified, any CUDA context can access the memory.
source
Array memory
Array memory is a special kind of memory that is optimized for 2D and 3D access patterns. The memory is opaquely managed by the CUDA runtime, and is typically only used on combination with texture intrinsics.
CUDA.ArrayMemory — Type
ArrayMemory
Array memory residing on the GPU, possibly in a specially-formatted way.
source
CUDA.alloc — Method
alloc(ArrayMemory, dims::Dims)
Allocate array memory with dimensions dims. The memory is accessible on the GPU, but can only be used in conjunction with special intrinsics (e.g., texture intrinsics).
source
Pointers
To work with these buffers, you need to convert them to a Ptr, CuPtr, or in the case of ArrayMemory an CuArrayPtr. You can then use common Julia methods on these pointers, such as unsafe_copyto!. CUDA.jl also provides some specialized functionality that does not match standard Julia functionality:
Version Management
CUDA.set_runtime_version! — Function
CUDA.set_runtime_version!([version::VersionNumber]; [local_toolkit::Bool])
Configures the active project to use a specific CUDA toolkit version from a specific source.
If local_toolkit is set, the CUDA toolkit will be used from the local system, otherwise it will be downloaded from an artifact source. In the case of a local toolkit, version informs CUDA.jl which version that is (this may be useful if auto-detection fails). In the case of artifact sources, version controls which version will be downloaded and used.
When not specifying either the version or the local_toolkit argument, the default behavior will be used, which is to use the most recent compatible runtime available from an artifact source. Note that this will override any Preferences that may be configured in a higher-up depot; to clear preferences nondestructively, use CUDA.reset_runtime_version! instead.
source
CUDA.reset_runtime_version! — Function
CUDA.reset_runtime_version!()
Resets the CUDA version preferences in the active project to the default, which is to use the most recent compatible runtime available from an artifact source, unless a higher-up depot has configured a different preference. To force use of the default behavior for the local project, use CUDA.set_runtime_version! with no arguments.
source
Device Management
CUDA.current_device — Function
current_device()
Returns the current device.
Warning
This is a low-level API, returning the current device as known to the CUDA driver. For most users, it is recommended to use the device method instead.
source
CUDA.name — Method
name(dev::CuDevice)
Returns an identifier string for the device.
source
CUDA.totalmem — Method
totalmem(dev::CuDevice)
Returns the total amount of memory (in bytes) on the device.
source
CUDA.attribute — Function
attribute(dev::CuDevice, code)
Returns information about the device.
source
attribute(X, pool::CuMemoryPool, attr)
Returns attribute attr about pool. The type of the returned value depends on the attribute, and as such must be passed as the X parameter.
source
attribute(X, ptr::Union{Ptr,CuPtr}, attr)
Returns attribute attr about pointer ptr. The type of the returned value depends on the attribute, and as such must be passed as the X parameter.
source
Certain common attributes are exposed by additional convenience functions:
CUDA.warpsize — Method
warpsize(dev::CuDevice)
Returns the warp size (in threads) of the device.
source
Context Management
CUDA.CuContext — Type
CuContext(dev::CuDevice, flags=CTX_SCHED_AUTO)
+CuContext(f::Function, ...)
Create a CUDA context for device. A context on the GPU is analogous to a process on the CPU, with its own distinct address space and allocated resources. When a context is destroyed, the system cleans up the resources allocated to it.
When you are done using the context, call CUDA.unsafe_destroy! to mark it for deletion, or use do-block syntax with this constructor.
source
CUDA.unsafe_destroy! — Method
unsafe_destroy!(ctx::CuContext)
Immediately destroy a context, freeing up all resources associated with it. This does not respect any users of the context, and might make other objects unusable.
source
CUDA.current_context — Function
current_context()
Returns the current context. Throws an undefined reference error if the current thread has no context bound to it, or if the bound context has been destroyed.
Warning
This is a low-level API, returning the current context as known to the CUDA driver. For most users, it is recommended to use the context method instead.
source
CUDA.activate — Method
activate(ctx::CuContext)
Binds the specified CUDA context to the calling CPU thread.
source
CUDA.synchronize — Method
synchronize(ctx::Context)
Block for the all operations on ctx to complete. This is a heavyweight operation, typically you only need to call synchronize which only synchronizes the stream associated with the current task.
source
CUDA.device_synchronize — Function
device_synchronize()
Block for the all operations on ctx to complete. This is a heavyweight operation, typically you only need to call synchronize which only synchronizes the stream associated with the current task.
On the device, device_synchronize acts as a synchronization point for child grids in the context of dynamic parallelism.
source
Primary Context Management
CUDA.CuPrimaryContext — Type
CuPrimaryContext(dev::CuDevice)
Create a primary CUDA context for a given device.
Each primary context is unique per device and is shared with CUDA runtime API. It is meant for interoperability with (applications using) the runtime API.
source
CUDA.CuContext — Method
CuContext(pctx::CuPrimaryContext)
Derive a context from a primary context.
Calling this function increases the reference count of the primary context. The returned context should not be free with the unsafe_destroy! function that's used with ordinary contexts. Instead, the refcount of the primary context should be decreased by calling unsafe_release!, or set to zero by calling unsafe_reset!. The easiest way to do this is by using the do-block syntax.
source
CUDA.isactive — Method
isactive(pctx::CuPrimaryContext)
Query whether a primary context is active.
source
CUDA.flags — Method
flags(pctx::CuPrimaryContext)
Query the flags of a primary context.
source
CUDA.unsafe_reset! — Method
unsafe_reset!(pctx::CuPrimaryContext)
Explicitly destroys and cleans up all resources associated with a device's primary context in the current process. Note that this forcibly invalidates all contexts derived from this primary context, and as a result outstanding resources might become invalid.
source
CUDA.unsafe_release! — Method
CUDA.unsafe_release!(pctx::CuPrimaryContext)
Lower the refcount of a context, possibly freeing up all resources associated with it. This does not respect any users of the context, and might make other objects unusable.
source
Module Management
CUDA.CuModule — Type
CuModule(data, options::Dict{CUjit_option,Any})
+CuModuleFile(path, options::Dict{CUjit_option,Any})
Create a CUDA module from a data, or a file containing data. The data may be PTX code, a CUBIN, or a FATBIN.
The options is an optional dictionary of JIT options and their respective value.
source
Function Management
CUDA.CuFunction — Type
CuFunction(mod::CuModule, name::String)
Acquires a function handle from a named function in a module.
source
Global Variable Management
CUDA.CuGlobal — Type
CuGlobal{T}(mod::CuModule, name::String)
Acquires a typed global variable handle from a named global in a module.
source
Base.eltype — Method
eltype(var::CuGlobal)
Return the element type of a global variable object.
source
Base.getindex — Method
Base.getindex(var::CuGlobal)
Return the current value of a global variable.
source
Base.setindex! — Method
Base.setindex(var::CuGlobal{T}, val::T)
Set the value of a global variable to val
source
Linker
CUDA.add_data! — Function
add_data!(link::CuLink, name::String, code::String)
Add PTX code to a pending link operation.
source
add_data!(link::CuLink, name::String, data::Vector{UInt8})
Add object code to a pending link operation.
source
CUDA.add_file! — Function
add_file!(link::CuLink, path::String, typ::CUjitInputType)
Add data from a file to a link operation. The argument typ indicates the type of the contained data.
source
CUDA.CuLinkImage — Type
The result of a linking operation.
This object keeps its parent linker object alive, as destroying a linker destroys linked images too.
source
CUDA.complete — Function
complete(link::CuLink)
Complete a pending linker invocation, returning an output image.
source
CUDA.CuModule — Method
CuModule(img::CuLinkImage, ...)
Create a CUDA module from a completed linking operation. Options from CuModule apply.
source
Memory Management
Different kinds of memory objects can be created, representing different kinds of memory that the CUDA toolkit supports. Each of these memory objects can be allocated by calling alloc with the type of memory as first argument, and freed by calling free. Certain kinds of memory have specific methods defined.
Device memory
This memory is accessible only by the GPU, and is the most common kind of memory used in CUDA programming.
CUDA.alloc — Method
alloc(DeviceMemory, bytesize::Integer;
+      [async=false], [stream::CuStream], [pool::CuMemoryPool])
Allocate bytesize bytes of memory on the device. This memory is only accessible on the GPU, and requires explicit calls to unsafe_copyto!, which wraps cuMemcpy, for access on the CPU.
source
Unified memory
Unified memory is accessible by both the CPU and the GPU, and is managed by the CUDA runtime. It is automatically migrated between the CPU and the GPU as needed, which simplifies programming but can lead to performance issues if not used carefully.
CUDA.alloc — Method
alloc(UnifiedMemory, bytesize::Integer, [flags::CUmemAttach_flags])
Allocate bytesize bytes of unified memory. This memory is accessible from both the CPU and GPU, with the CUDA driver automatically copying upon first access.
source
CUDA.prefetch — Method
prefetch(::UnifiedMemory, [bytes::Integer]; [device::CuDevice], [stream::CuStream])
Prefetches memory to the specified destination device.
source
CUDA.advise — Method
advise(::UnifiedMemory, advice::CUDA.CUmem_advise, [bytes::Integer]; [device::CuDevice])
Advise about the usage of a given memory range.
source
Host memory
Host memory resides on the CPU, but is accessible by the GPU via the PCI bus. This is the slowest kind of memory, but is useful for communicating between running kernels and the host (e.g., to update counters or flags).
CUDA.HostMemory — Type
HostMemory
Pinned memory residing on the CPU, possibly accessible on the GPU.
source
CUDA.alloc — Method
alloc(HostMemory, bytesize::Integer, [flags])
Allocate bytesize bytes of page-locked memory on the host. This memory is accessible from the CPU, and makes it possible to perform faster memory copies to the GPU. Furthermore, if flags is set to MEMHOSTALLOC_DEVICEMAP the memory is also accessible from the GPU. These accesses are direct, and go through the PCI bus. If flags is set to MEMHOSTALLOC_PORTABLE, the memory is considered mapped by all CUDA contexts, not just the one that created the memory, which is useful if the memory needs to be accessed from multiple devices. Multiple flags can be set at one time using a bytewise OR:
flags = MEMHOSTALLOC_PORTABLE | MEMHOSTALLOC_DEVICEMAP
source
CUDA.register — Method
register(HostMemory, ptr::Ptr, bytesize::Integer, [flags])
Page-lock the host memory pointed to by ptr. Subsequent transfers to and from devices will be faster, and can be executed asynchronously. If the MEMHOSTREGISTER_DEVICEMAP flag is specified, the buffer will also be accessible directly from the GPU. These accesses are direct, and go through the PCI bus. If the MEMHOSTREGISTER_PORTABLE flag is specified, any CUDA context can access the memory.
source
Array memory
Array memory is a special kind of memory that is optimized for 2D and 3D access patterns. The memory is opaquely managed by the CUDA runtime, and is typically only used on combination with texture intrinsics.
CUDA.ArrayMemory — Type
ArrayMemory
Array memory residing on the GPU, possibly in a specially-formatted way.
source
CUDA.alloc — Method
alloc(ArrayMemory, dims::Dims)
Allocate array memory with dimensions dims. The memory is accessible on the GPU, but can only be used in conjunction with special intrinsics (e.g., texture intrinsics).
source
Pointers
To work with these buffers, you need to convert them to a Ptr, CuPtr, or in the case of ArrayMemory an CuArrayPtr. You can then use common Julia methods on these pointers, such as unsafe_copyto!. CUDA.jl also provides some specialized functionality that does not match standard Julia functionality:
CUDA.unsafe_copy2d! — Function
unsafe_copy2d!(dst, dstTyp, src, srcTyp, width, height=1;
                dstPos=(1,1), dstPitch=0,
                srcPos=(1,1), srcPitch=0,
-               async=false, stream=nothing)
Perform a 2D memory copy between pointers src and dst, at respectively position srcPos and dstPos (1-indexed). Pitch can be specified for both the source and destination; consult the CUDA documentation for more details. This call is executed asynchronously if async is set, otherwise stream is synchronized.
source
CUDA.unsafe_copy3d! — Function
unsafe_copy3d!(dst, dstTyp, src, srcTyp, width, height=1, depth=1;
+               async=false, stream=nothing)
Perform a 2D memory copy between pointers src and dst, at respectively position srcPos and dstPos (1-indexed). Pitch can be specified for both the source and destination; consult the CUDA documentation for more details. This call is executed asynchronously if async is set, otherwise stream is synchronized.
source
CUDA.unsafe_copy3d! — Function
unsafe_copy3d!(dst, dstTyp, src, srcTyp, width, height=1, depth=1;
                dstPos=(1,1,1), dstPitch=0, dstHeight=0,
                srcPos=(1,1,1), srcPitch=0, srcHeight=0,
-               async=false, stream=nothing)
Perform a 3D memory copy between pointers src and dst, at respectively position srcPos and dstPos (1-indexed). Both pitch and height can be specified for both the source and destination; consult the CUDA documentation for more details. This call is executed asynchronously if async is set, otherwise stream is synchronized.
source
CUDA.memset — Function
memset(mem::CuPtr, value::Union{UInt8,UInt16,UInt32}, len::Integer; [stream::CuStream])
Initialize device memory by copying val for len times.
source
Other
CUDA.free_memory — Function
free_memory()
Returns the free amount of memory (in bytes), available for allocation by the CUDA context.
source
CUDA.total_memory — Function
total_memory()
Returns the total amount of memory (in bytes), available for allocation by the CUDA context.
source
Stream Management
CUDA.CuStream — Type
CuStream(; flags=STREAM_DEFAULT, priority=nothing)
Create a CUDA stream.
source
CUDA.isdone — Method
isdone(s::CuStream)
Return false if a stream is busy (has task running or queued) and true if that stream is free.
source
CUDA.priority_range — Function
priority_range()
Return the valid range of stream priorities as a StepRange (with step size  1). The lower bound of the range denotes the least priority (typically 0), with the upper bound representing the greatest possible priority (typically -1).
source
CUDA.synchronize — Method
synchronize([stream::CuStream])
Wait until stream has finished executing, with stream defaulting to the stream associated with the current Julia task.
See also: device_synchronize
source
CUDA.@sync — Macro
@sync [blocking=false] ex
Run expression ex and synchronize the GPU afterwards.
The blocking keyword argument determines how synchronization is performed. By default, non-blocking synchronization will be used, which gives other Julia tasks a chance to run while waiting for the GPU to finish. This may increase latency, so for short operations, or when benchmaring code that does not use multiple tasks, it may be beneficial to use blocking synchronization instead by setting blocking=true. Blocking synchronization can also be enabled globally by changing the nonblocking_synchronization preference.
See also: synchronize.
source
For specific use cases, special streams are available:
CUDA.default_stream — Function
default_stream()
Return the default stream.
Note
It is generally better to use stream() to get a stream object that's local to the current task. That way, operations scheduled in other tasks can overlap.
source
CUDA.legacy_stream — Function
legacy_stream()
Return a special object to use use an implicit stream with legacy synchronization behavior.
You can use this stream to perform operations that should block on all streams (with the exception of streams created with STREAM_NON_BLOCKING). This matches the old pre-CUDA 7 global stream behavior.
source
CUDA.per_thread_stream — Function
per_thread_stream()
Return a special object to use an implicit stream with per-thread synchronization behavior. This stream object is normally meant to be used with APIs that do not have per-thread versions of their APIs (i.e. without a ptsz or ptds suffix).
Note
It is generally not needed to use this type of stream. With CUDA.jl, each task already gets its own non-blocking stream, and multithreading in Julia is typically accomplished using tasks.
source
Event Management
CUDA.record — Function
record(e::CuEvent, [stream::CuStream])
Record an event on a stream.
source
CUDA.isdone — Method
isdone(e::CuEvent)
Return false if there is outstanding work preceding the most recent call to record(e) and true if all captured work has been completed.
source
CUDA.elapsed — Function
elapsed(start::CuEvent, stop::CuEvent)
Computes the elapsed time between two events (in seconds).
source
CUDA.@elapsed — Macro
@elapsed [blocking=false] ex
A macro to evaluate an expression, discarding the resulting value, instead returning the number of seconds it took to execute on the GPU, as a floating-point number.
See also: @sync.
source
Execution Control
CUDA.CuDim3 — Type
CuDim3(x)
+               async=false, stream=nothing)
Perform a 3D memory copy between pointers src and dst, at respectively position srcPos and dstPos (1-indexed). Both pitch and height can be specified for both the source and destination; consult the CUDA documentation for more details. This call is executed asynchronously if async is set, otherwise stream is synchronized.
source
CUDA.memset — Function
memset(mem::CuPtr, value::Union{UInt8,UInt16,UInt32}, len::Integer; [stream::CuStream])
Initialize device memory by copying val for len times.
source
Other
CUDA.free_memory — Function
free_memory()
Returns the free amount of memory (in bytes), available for allocation by the CUDA context.
source
CUDA.total_memory — Function
total_memory()
Returns the total amount of memory (in bytes), available for allocation by the CUDA context.
source
Stream Management
CUDA.CuStream — Type
CuStream(; flags=STREAM_DEFAULT, priority=nothing)
Create a CUDA stream.
source
CUDA.isdone — Method
isdone(s::CuStream)
Return false if a stream is busy (has task running or queued) and true if that stream is free.
source
CUDA.priority_range — Function
priority_range()
Return the valid range of stream priorities as a StepRange (with step size  1). The lower bound of the range denotes the least priority (typically 0), with the upper bound representing the greatest possible priority (typically -1).
source
CUDA.synchronize — Method
synchronize([stream::CuStream])
Wait until stream has finished executing, with stream defaulting to the stream associated with the current Julia task.
See also: device_synchronize
source
CUDA.@sync — Macro
@sync [blocking=false] ex
Run expression ex and synchronize the GPU afterwards.
The blocking keyword argument determines how synchronization is performed. By default, non-blocking synchronization will be used, which gives other Julia tasks a chance to run while waiting for the GPU to finish. This may increase latency, so for short operations, or when benchmaring code that does not use multiple tasks, it may be beneficial to use blocking synchronization instead by setting blocking=true. Blocking synchronization can also be enabled globally by changing the nonblocking_synchronization preference.
See also: synchronize.
source
For specific use cases, special streams are available:
CUDA.default_stream — Function
default_stream()
Return the default stream.
Note
It is generally better to use stream() to get a stream object that's local to the current task. That way, operations scheduled in other tasks can overlap.
source
CUDA.legacy_stream — Function
legacy_stream()
Return a special object to use use an implicit stream with legacy synchronization behavior.
You can use this stream to perform operations that should block on all streams (with the exception of streams created with STREAM_NON_BLOCKING). This matches the old pre-CUDA 7 global stream behavior.
source
CUDA.per_thread_stream — Function
per_thread_stream()
Return a special object to use an implicit stream with per-thread synchronization behavior. This stream object is normally meant to be used with APIs that do not have per-thread versions of their APIs (i.e. without a ptsz or ptds suffix).
Note
It is generally not needed to use this type of stream. With CUDA.jl, each task already gets its own non-blocking stream, and multithreading in Julia is typically accomplished using tasks.
source
Event Management
CUDA.record — Function
record(e::CuEvent, [stream::CuStream])
Record an event on a stream.
source
CUDA.isdone — Method
isdone(e::CuEvent)
Return false if there is outstanding work preceding the most recent call to record(e) and true if all captured work has been completed.
source
CUDA.elapsed — Function
elapsed(start::CuEvent, stop::CuEvent)
Computes the elapsed time between two events (in seconds).
source
CUDA.@elapsed — Macro
@elapsed [blocking=false] ex
A macro to evaluate an expression, discarding the resulting value, instead returning the number of seconds it took to execute on the GPU, as a floating-point number.
See also: @sync.
source
Execution Control
CUDA.CuDim3 — Type
CuDim3(x)
 
 CuDim3((x,))
 CuDim3((x, y))
-CuDim3((x, y, x))
A type used to specify dimensions, consisting of 3 integers for respectively the x, y and z dimension. Unspecified dimensions default to 1.
Often accepted as argument through the CuDim type alias, eg. in the case of cudacall or CUDA.launch, allowing to pass dimensions as a plain integer or a tuple without having to construct an explicit CuDim3 object.
source
CUDA.cudacall — Function
cudacall(f, types, values...; blocks::CuDim, threads::CuDim,
+CuDim3((x, y, x))
A type used to specify dimensions, consisting of 3 integers for respectively the x, y and z dimension. Unspecified dimensions default to 1.
Often accepted as argument through the CuDim type alias, eg. in the case of cudacall or CUDA.launch, allowing to pass dimensions as a plain integer or a tuple without having to construct an explicit CuDim3 object.
source
CUDA.cudacall — Function
cudacall(f, types, values...; blocks::CuDim, threads::CuDim,
          cooperative=false, shmem=0, stream=stream())
ccall-like interface for launching a CUDA function f on a GPU.
For example:
vadd = CuFunction(md, "vadd")
 a = rand(Float32, 10)
 b = rand(Float32, 10)
@@ -32,13 +32,13 @@
 cd = alloc(CUDA.DeviceMemory, 10*sizeof(Float32))
 
 cudacall(vadd, (CuPtr{Cfloat},CuPtr{Cfloat},CuPtr{Cfloat}), ad, bd, cd; threads=10)
-unsafe_copyto!(convert(Ptr{Cvoid}, c), cd, 10*sizeof(Float32)))
The blocks and threads arguments control the launch configuration, and should both consist of either an integer, or a tuple of 1 to 3 integers (omitted dimensions default to 1). The types argument can contain both a tuple of types, and a tuple type, the latter being slightly faster.
source
CUDA.launch — Function
launch(f::CuFunction; args...; blocks::CuDim=1, threads::CuDim=1,
-       cooperative=false, shmem=0, stream=stream())
Low-level call to launch a CUDA function f on the GPU, using blocks and threads as respectively the grid and block configuration. Dynamic shared memory is allocated according to shmem, and the kernel is launched on stream stream.
Arguments to a kernel should either be bitstype, in which case they will be copied to the internal kernel parameter buffer, or a pointer to device memory.
This is a low-level call, prefer to use cudacall instead.
source
launch(exec::CuGraphExec, [stream::CuStream])
Launches an executable graph, by default in the currently-active stream.
source
Profiler Control
CUDA.@profile — Macro
@profile [trace=false] [raw=false] code...
-@profile external=true code...
Profile the GPU execution of code.
There are two modes of operation, depending on whether external is true or false. The default value depends on whether Julia is being run under an external profiler.
Integrated profiler (external=false, the default)
In this mode, CUDA.jl will profile the execution of code and display the result. By default, a summary of host and device-side execution will be show, including any NVTX events. To display a chronological trace of the captured activity instead, trace can be set to true. Trace output will include an ID column that can be used to match host-side and device-side activity. If raw is true, all data will always be included, even if it may not be relevant. The output will be written to io, which defaults to stdout.
Slow operations will be highlighted in the output: Entries colored in yellow are among the slowest 25%, while entries colored in red are among the slowest 5% of all operations.
!!! compat "Julia 1.9" This functionality is only available on Julia 1.9 and later.
!!! compat "CUDA 11.2" Older versions of CUDA, before 11.2, contain bugs that may prevent     the CUDA.@profile macro to work. It is recommended to use a newer runtime.
External profilers (external=true, when an external profiler is detected)
For more advanced profiling, it is possible to use an external profiling tool, such as NSight Systems or NSight Compute. When doing so, it is often advisable to only enable the profiler for the specific code region of interest. This can be done by wrapping the code with CUDA.@profile external=true, which used to be the only way to use this macro.
source
CUDA.Profile.start — Function
start()
Enables profile collection by the active profiling tool for the current context. If profiling is already enabled, then this call has no effect.
source
CUDA.Profile.stop — Function
stop()
Disables profile collection by the active profiling tool for the current context. If profiling is already disabled, then this call has no effect.
source
Texture Memory
Textures are represented by objects of type CuTexture which are bound to some underlying memory, either CuArrays or CuTextureArrays:
CUDA.CuTexture — Type
CuTexture{T,N,P}
N-dimensional texture object with elements of type T. These objects do not store data themselves, but are bounds to another source of device memory. Texture objects can be passed to CUDA kernels, where they will be accessible through the CuDeviceTexture type.
Warning
Experimental API. Subject to change without deprecation.
source
CUDA.CuTexture — Method
CuTexture{T,N,P}(parent::P; address_mode, filter_mode, normalized_coordinates)
Construct a N-dimensional texture object with elements of type T as stored in parent.
Several keyword arguments alter the behavior of texture objects:
  • address_mode (wrap, clamp, mirror): how out-of-bounds values are accessed. Can be specified as a value for all dimensions, or as a tuple of N entries.
  • interpolation (nearest neighbour, linear, bilinear): how non-integral indices are fetched. Nearest-neighbour fetches a single value, others interpolate between multiple.
  • normalized_coordinates (true, false): whether indices are expected to fall in the normalized [0:1) range.
!!! warning Experimental API. Subject to change without deprecation.
source
CuTexture(x::CuTextureArray{T,N})
Create a N-dimensional texture object withelements of type T that will be read from x.
Warning
Experimental API. Subject to change without deprecation.
source
CuTexture(x::CuArray{T,N})
Create a N-dimensional texture object that reads from a CuArray.
Note that it is necessary the their memory is well aligned and strided (good pitch). Currently, that is not being enforced.
Warning
Experimental API. Subject to change without deprecation.
source
You can create CuTextureArray objects from both host and device memory:
CUDA.CuTextureArray — Type
CuTextureArray{T,N}(undef, dims)
N-dimensional dense texture array with elements of type T. These arrays are optimized for texture fetching, and are only meant to be used as a source for CuTexture{T,N,P} objects.
Warning
Experimental API. Subject to change without deprecation.
source
CUDA.CuTextureArray — Method
CuTextureArray(A::AbstractArray)
Allocate and initialize a texture array from host memory in A.
Warning
Experimental API. Subject to change without deprecation.
source
CuTextureArray(A::CuArray)
Allocate and initialize a texture array from device memory in A.
Warning
Experimental API. Subject to change without deprecation.
source
Occupancy API
The occupancy API can be used to figure out an appropriate launch configuration for a compiled kernel (represented as a CuFunction) on the current device:
CUDA.launch_configuration — Function
launch_configuration(fun::CuFunction; shmem=0, max_threads=0)
Calculate a suggested launch configuration for kernel fun requiring shmem bytes of dynamic shared memory. Returns a tuple with a suggested amount of threads, and the minimal amount of blocks to reach maximal occupancy. Optionally, the maximum amount of threads can be constrained using max_threads.
In the case of a variable amount of shared memory, pass a callable object for shmem instead, taking a single integer representing the block size and returning the amount of dynamic shared memory for that configuration.
source
CUDA.active_blocks — Function
active_blocks(fun::CuFunction, threads; shmem=0)
Calculate the maximum number of active blocks per multiprocessor when running threads threads of a kernel fun requiring shmem bytes of dynamic shared memory.
source
CUDA.occupancy — Function
occupancy(fun::CuFunction, threads; shmem=0)
Calculate the theoretical occupancy of launching threads threads of a kernel fun requiring shmem bytes of dynamic shared memory.
source
Graph Execution
CUDA graphs can be easily recorded and executed using the high-level @captured macro:
CUDA.@captured — Macro
for ...
+unsafe_copyto!(convert(Ptr{Cvoid}, c), cd, 10*sizeof(Float32)))
The blocks and threads arguments control the launch configuration, and should both consist of either an integer, or a tuple of 1 to 3 integers (omitted dimensions default to 1). The types argument can contain both a tuple of types, and a tuple type, the latter being slightly faster.
source
CUDA.launch — Function
launch(f::CuFunction; args...; blocks::CuDim=1, threads::CuDim=1,
+       cooperative=false, shmem=0, stream=stream())
Low-level call to launch a CUDA function f on the GPU, using blocks and threads as respectively the grid and block configuration. Dynamic shared memory is allocated according to shmem, and the kernel is launched on stream stream.
Arguments to a kernel should either be bitstype, in which case they will be copied to the internal kernel parameter buffer, or a pointer to device memory.
This is a low-level call, prefer to use cudacall instead.
source
launch(exec::CuGraphExec, [stream::CuStream])
Launches an executable graph, by default in the currently-active stream.
source
Profiler Control
CUDA.@profile — Macro
@profile [trace=false] [raw=false] code...
+@profile external=true code...
Profile the GPU execution of code.
There are two modes of operation, depending on whether external is true or false. The default value depends on whether Julia is being run under an external profiler.
Integrated profiler (external=false, the default)
In this mode, CUDA.jl will profile the execution of code and display the result. By default, a summary of host and device-side execution will be show, including any NVTX events. To display a chronological trace of the captured activity instead, trace can be set to true. Trace output will include an ID column that can be used to match host-side and device-side activity. If raw is true, all data will always be included, even if it may not be relevant. The output will be written to io, which defaults to stdout.
Slow operations will be highlighted in the output: Entries colored in yellow are among the slowest 25%, while entries colored in red are among the slowest 5% of all operations.
!!! compat "Julia 1.9" This functionality is only available on Julia 1.9 and later.
!!! compat "CUDA 11.2" Older versions of CUDA, before 11.2, contain bugs that may prevent     the CUDA.@profile macro to work. It is recommended to use a newer runtime.
External profilers (external=true, when an external profiler is detected)
For more advanced profiling, it is possible to use an external profiling tool, such as NSight Systems or NSight Compute. When doing so, it is often advisable to only enable the profiler for the specific code region of interest. This can be done by wrapping the code with CUDA.@profile external=true, which used to be the only way to use this macro.
source
CUDA.Profile.start — Function
start()
Enables profile collection by the active profiling tool for the current context. If profiling is already enabled, then this call has no effect.
source
CUDA.Profile.stop — Function
stop()
Disables profile collection by the active profiling tool for the current context. If profiling is already disabled, then this call has no effect.
source
Texture Memory
Textures are represented by objects of type CuTexture which are bound to some underlying memory, either CuArrays or CuTextureArrays:
CUDA.CuTexture — Type
CuTexture{T,N,P}
N-dimensional texture object with elements of type T. These objects do not store data themselves, but are bounds to another source of device memory. Texture objects can be passed to CUDA kernels, where they will be accessible through the CuDeviceTexture type.
Warning
Experimental API. Subject to change without deprecation.
source
CUDA.CuTexture — Method
CuTexture{T,N,P}(parent::P; address_mode, filter_mode, normalized_coordinates)
Construct a N-dimensional texture object with elements of type T as stored in parent.
Several keyword arguments alter the behavior of texture objects:
  • address_mode (wrap, clamp, mirror): how out-of-bounds values are accessed. Can be specified as a value for all dimensions, or as a tuple of N entries.
  • interpolation (nearest neighbour, linear, bilinear): how non-integral indices are fetched. Nearest-neighbour fetches a single value, others interpolate between multiple.
  • normalized_coordinates (true, false): whether indices are expected to fall in the normalized [0:1) range.
!!! warning Experimental API. Subject to change without deprecation.
source
CuTexture(x::CuTextureArray{T,N})
Create a N-dimensional texture object withelements of type T that will be read from x.
Warning
Experimental API. Subject to change without deprecation.
source
CuTexture(x::CuArray{T,N})
Create a N-dimensional texture object that reads from a CuArray.
Note that it is necessary the their memory is well aligned and strided (good pitch). Currently, that is not being enforced.
Warning
Experimental API. Subject to change without deprecation.
source
You can create CuTextureArray objects from both host and device memory:
CUDA.CuTextureArray — Type
CuTextureArray{T,N}(undef, dims)
N-dimensional dense texture array with elements of type T. These arrays are optimized for texture fetching, and are only meant to be used as a source for CuTexture{T,N,P} objects.
Warning
Experimental API. Subject to change without deprecation.
source
CUDA.CuTextureArray — Method
CuTextureArray(A::AbstractArray)
Allocate and initialize a texture array from host memory in A.
Warning
Experimental API. Subject to change without deprecation.
source
CuTextureArray(A::CuArray)
Allocate and initialize a texture array from device memory in A.
Warning
Experimental API. Subject to change without deprecation.
source
Occupancy API
The occupancy API can be used to figure out an appropriate launch configuration for a compiled kernel (represented as a CuFunction) on the current device:
CUDA.launch_configuration — Function
launch_configuration(fun::CuFunction; shmem=0, max_threads=0)
Calculate a suggested launch configuration for kernel fun requiring shmem bytes of dynamic shared memory. Returns a tuple with a suggested amount of threads, and the minimal amount of blocks to reach maximal occupancy. Optionally, the maximum amount of threads can be constrained using max_threads.
In the case of a variable amount of shared memory, pass a callable object for shmem instead, taking a single integer representing the block size and returning the amount of dynamic shared memory for that configuration.
source
CUDA.active_blocks — Function
active_blocks(fun::CuFunction, threads; shmem=0)
Calculate the maximum number of active blocks per multiprocessor when running threads threads of a kernel fun requiring shmem bytes of dynamic shared memory.
source
CUDA.occupancy — Function
occupancy(fun::CuFunction, threads; shmem=0)
Calculate the theoretical occupancy of launching threads threads of a kernel fun requiring shmem bytes of dynamic shared memory.
source
Graph Execution
CUDA graphs can be easily recorded and executed using the high-level @captured macro:
CUDA.@captured — Macro
for ...
     @captured begin
         # code that executes several kernels or CUDA operations
     end
-end
A convenience macro for recording a graph of CUDA operations and automatically cache and update the execution. This can improve performance when executing kernels in a loop, where the launch overhead might dominate the execution.
Warning
For this to be effective, the kernels and operations executed inside of the captured region should not signficantly change across iterations of the loop. It is allowed to, e.g., change kernel arguments or inputs to operations, as this will be processed by updating the cached executable graph. However, significant changes will result in an instantiation of the graph from scratch, which is an expensive operation.
See also: capture.
source
Low-level operations are available too:
CUDA.CuGraph — Type
CuGraph([flags])
Create an empty graph for use with low-level graph operations. If you want to create a graph while directly recording operations, use capture. For a high-level interface that also automatically executes the graph, use the @captured macro.
source
CUDA.capture — Function
capture([flags], [throw_error::Bool=true]) do
+end
A convenience macro for recording a graph of CUDA operations and automatically cache and update the execution. This can improve performance when executing kernels in a loop, where the launch overhead might dominate the execution.
Warning
For this to be effective, the kernels and operations executed inside of the captured region should not signficantly change across iterations of the loop. It is allowed to, e.g., change kernel arguments or inputs to operations, as this will be processed by updating the cached executable graph. However, significant changes will result in an instantiation of the graph from scratch, which is an expensive operation.
See also: capture.
source
Low-level operations are available too:
CUDA.CuGraph — Type
CuGraph([flags])
Create an empty graph for use with low-level graph operations. If you want to create a graph while directly recording operations, use capture. For a high-level interface that also automatically executes the graph, use the @captured macro.
source
CUDA.capture — Function
capture([flags], [throw_error::Bool=true]) do
     ...
-end
Capture a graph of CUDA operations. The returned graph can then be instantiated and executed repeatedly for improved performance.
Note that many operations, like initial kernel compilation or memory allocations, cannot be captured. To work around this, you can set the throw_error keyword to false, which will cause this function to return nothing if such a failure happens. You can then try to evaluate the function in a regular way, and re-record afterwards.
See also: instantiate.
source
CUDA.instantiate — Function
instantiate(graph::CuGraph)
Creates an executable graph from a graph. This graph can then be launched, or updated with an other graph.
See also: launch, update.
source
CUDA.launch — Method
launch(f::CuFunction; args...; blocks::CuDim=1, threads::CuDim=1,
-       cooperative=false, shmem=0, stream=stream())
Low-level call to launch a CUDA function f on the GPU, using blocks and threads as respectively the grid and block configuration. Dynamic shared memory is allocated according to shmem, and the kernel is launched on stream stream.
Arguments to a kernel should either be bitstype, in which case they will be copied to the internal kernel parameter buffer, or a pointer to device memory.
This is a low-level call, prefer to use cudacall instead.
source
launch(exec::CuGraphExec, [stream::CuStream])
Launches an executable graph, by default in the currently-active stream.
source
CUDA.update — Function
update(exec::CuGraphExec, graph::CuGraph; [throw_error::Bool=true])
Check whether an executable graph can be updated with a graph and perform the update if possible. Returns a boolean indicating whether the update was successful. Unless throw_error is set to false, also throws an error if the update failed.
source

+end

Capture a graph of CUDA operations. The returned graph can then be instantiated and executed repeatedly for improved performance.

Note that many operations, like initial kernel compilation or memory allocations, cannot be captured. To work around this, you can set the throw_error keyword to false, which will cause this function to return nothing if such a failure happens. You can then try to evaluate the function in a regular way, and re-record afterwards.

See also: instantiate.

source
CUDA.instantiate — Function
instantiate(graph::CuGraph)

Creates an executable graph from a graph. This graph can then be launched, or updated with an other graph.

See also: launch, update.

source
CUDA.launch — Method
launch(f::CuFunction; args...; blocks::CuDim=1, threads::CuDim=1,
+       cooperative=false, shmem=0, stream=stream())

Low-level call to launch a CUDA function f on the GPU, using blocks and threads as respectively the grid and block configuration. Dynamic shared memory is allocated according to shmem, and the kernel is launched on stream stream.

Arguments to a kernel should either be bitstype, in which case they will be copied to the internal kernel parameter buffer, or a pointer to device memory.

This is a low-level call, prefer to use cudacall instead.

source
launch(exec::CuGraphExec, [stream::CuStream])

Launches an executable graph, by default in the currently-active stream.

source
CUDA.update — Function
update(exec::CuGraphExec, graph::CuGraph; [throw_error::Bool=true])

Check whether an executable graph can be updated with a graph and perform the update if possible. Returns a boolean indicating whether the update was successful. Unless throw_error is set to false, also throws an error if the update failed.

source
diff --git a/dev/tutorials/custom_structs/index.html b/dev/tutorials/custom_structs/index.html index 6236189e34..49b006f95f 100644 --- a/dev/tutorials/custom_structs/index.html +++ b/dev/tutorials/custom_structs/index.html @@ -32,4 +32,4 @@ Interpolate(xs, ys) end

Now our struct plays nicely with CUDA.jl:

result = itp.(pts)
2-element CuArray{Float64, 1, CUDA.DeviceMemory}:
  20.0
- 30.0

It works, we get the same result as on the CPU.

@assert CuArray(result_cpu) == result

Alternatively instead of defining Adapt.adapt_structure explictly, we could have done

Adapt.@adapt_structure Interpolate

which expands to the same code that we wrote manually.

This page was generated using Literate.jl.

+ 30.0

It works, we get the same result as on the CPU.

@assert CuArray(result_cpu) == result

Alternatively instead of defining Adapt.adapt_structure explictly, we could have done

Adapt.@adapt_structure Interpolate

which expands to the same code that we wrote manually.

This page was generated using Literate.jl.

diff --git a/dev/tutorials/introduction/index.html b/dev/tutorials/introduction/index.html index 8d801466b0..11b00ad0a3 100644 --- a/dev/tutorials/introduction/index.html +++ b/dev/tutorials/introduction/index.html @@ -89,45 +89,45 @@ @cuda gpu_add1!(y, x) end end
bench_gpu1! (generic function with 1 method)
@btime bench_gpu1!($y_d, $x_d)
  119.783 ms (47 allocations: 1.23 KiB)

That's a lot slower than the version above based on broadcasting. What happened?

Profiling

When you don't get the performance you expect, usually your first step should be to profile the code and see where it's spending its time:

bench_gpu1!(y_d, x_d)  # run it once to force compilation
-CUDA.@profile bench_gpu1!(y_d, x_d)
Profiler ran for 77.68 ms, capturing 804 events.
+CUDA.@profile bench_gpu1!(y_d, x_d)
Profiler ran for 75.97 ms, capturing 804 events.
 
-Host-side activity: calling CUDA APIs took 76.45 ms (98.43% of the trace)
+Host-side activity: calling CUDA APIs took 74.96 ms (98.67% of the trace)
 ┌──────────┬────────────┬───────┬─────────────────────┐
 │ Time (%) │ Total time │ Calls │ Name                │
 ├──────────┼────────────┼───────┼─────────────────────┤
-│   98.42% │   76.45 ms │     1 │ cuStreamSynchronize │
-│    0.05% │   37.91 µs │     1 │ cuLaunchKernel      │
-│    0.00% │    2.86 µs │     1 │ cuCtxSetCurrent     │
-│    0.00% │  715.26 ns │     1 │ cuCtxGetDevice      │
-│    0.00% │  476.84 ns │     1 │ cuDeviceGetCount    │
+│   98.67% │   74.95 ms │     1 │ cuStreamSynchronize │
+│    0.06% │   43.15 µs │     1 │ cuLaunchKernel      │
+│    0.00% │    3.34 µs │     1 │ cuCtxSetCurrent     │
+│    0.00% │  953.67 ns │     1 │ cuCtxGetDevice      │
+│    0.00% │  953.67 ns │     1 │ cuDeviceGetCount    │
 └──────────┴────────────┴───────┴─────────────────────┘
 
-Device-side activity: GPU was busy for 77.28 ms (99.50% of the trace)
+Device-side activity: GPU was busy for 75.84 ms (99.83% of the trace)
 ┌──────────┬────────────┬───────┬───────────────────────────────────────────────
 │ Time (%) │ Total time │ Calls │ Name                                         ⋯
 ├──────────┼────────────┼───────┼───────────────────────────────────────────────
-│   99.50% │   77.28 ms │     1 │ _Z9gpu_add1_13CuDeviceArrayI7Float32Ll1ELl1E ⋯
+│   99.83% │   75.84 ms │     1 │ _Z9gpu_add1_13CuDeviceArrayI7Float32Ll1ELl1E ⋯
 └──────────┴────────────┴───────┴───────────────────────────────────────────────
                                                                 1 column omitted
-

You can see that almost all of the time was spent in ptxcall_gpu_add1__1, the name of the kernel that CUDA.jl assigned when compiling gpu_add1! for these inputs. (Had you created arrays of multiple data types, e.g., xu_d = CUDA.fill(0x01, N), you might have also seen ptxcall_gpu_add1__2 and so on. Like the rest of Julia, you can define a single method and it will be specialized at compile time for the particular data types you're using.)

For further insight, run the profiling with the option trace=true

CUDA.@profile trace=true bench_gpu1!(y_d, x_d)
Profiler ran for 107.91 ms, capturing 804 events.
+

You can see that almost all of the time was spent in ptxcall_gpu_add1__1, the name of the kernel that CUDA.jl assigned when compiling gpu_add1! for these inputs. (Had you created arrays of multiple data types, e.g., xu_d = CUDA.fill(0x01, N), you might have also seen ptxcall_gpu_add1__2 and so on. Like the rest of Julia, you can define a single method and it will be specialized at compile time for the particular data types you're using.)

For further insight, run the profiling with the option trace=true

CUDA.@profile trace=true bench_gpu1!(y_d, x_d)
Profiler ran for 107.86 ms, capturing 804 events.
 
-Host-side activity: calling CUDA APIs took 107.04 ms (99.19% of the trace)
+Host-side activity: calling CUDA APIs took 106.85 ms (99.06% of the trace)
 ┌─────┬───────────┬───────────┬────────┬─────────────────────┐
 │  ID │     Start │      Time │ Thread │ Name                │
 ├─────┼───────────┼───────────┼────────┼─────────────────────┤
-│  21 │  76.53 µs │  34.09 µs │      1 │ cuLaunchKernel      │
-│ 795 │ 822.07 µs │   2.62 µs │      2 │ cuCtxSetCurrent     │
-│ 796 │ 829.46 µs │ 715.26 ns │      2 │ cuCtxGetDevice      │
-│ 797 │  837.8 µs │ 715.26 ns │      2 │ cuDeviceGetCount    │
-│ 800 │ 846.15 µs │ 107.04 ms │      2 │ cuStreamSynchronize │
+│  21 │  46.97 µs │   36.0 µs │      1 │ cuLaunchKernel      │
+│ 795 │ 939.85 µs │   3.58 µs │      2 │ cuCtxSetCurrent     │
+│ 796 │ 950.34 µs │   1.19 µs │      2 │ cuCtxGetDevice      │
+│ 797 │ 963.45 µs │ 953.67 ns │      2 │ cuDeviceGetCount    │
+│ 800 │ 980.38 µs │ 106.84 ms │      2 │ cuStreamSynchronize │
 └─────┴───────────┴───────────┴────────┴─────────────────────┘
 
-Device-side activity: GPU was busy for 107.76 ms (99.86% of the trace)
-┌────┬───────────┬───────────┬─────────┬────────┬──────┬────────────────────────
-│ ID │     Start │      Time │ Threads │ Blocks │ Regs │ Name                  ⋯
-├────┼───────────┼───────────┼─────────┼────────┼──────┼────────────────────────
-│ 21 │ 113.01 µs │ 107.76 ms │       1 │      1 │   19 │ _Z9gpu_add1_13CuDevic ⋯
-└────┴───────────┴───────────┴─────────┴────────┴──────┴────────────────────────
+Device-side activity: GPU was busy for 107.73 ms (99.88% of the trace)
+┌────┬──────────┬───────────┬─────────┬────────┬──────┬─────────────────────────
+│ ID │    Start │      Time │ Threads │ Blocks │ Regs │ Name                   ⋯
+├────┼──────────┼───────────┼─────────┼────────┼──────┼─────────────────────────
+│ 21 │ 83.68 µs │ 107.73 ms │       1 │      1 │   19 │ _Z9gpu_add1_13CuDevice ⋯
+└────┴──────────┴───────────┴─────────┴────────┴──────┴─────────────────────────
                                                                 1 column omitted

The key thing to note here is that we are only using a single block with a single thread. These terms will be explained shortly, but for now, suffice it to say that this is an indication that this computation ran sequentially. Of note, sequential processing with GPUs is much slower than with CPUs; where GPUs shine is with large-scale parallelism.

Writing a parallel GPU kernel

To speed up the kernel, we want to parallelize it, which means assigning different tasks to different threads. To facilitate the assignment of work, each CUDA thread gets access to variables that indicate its own unique identity, much as Threads.threadid() does for CPU threads. The CUDA analogs of threadid and nthreads are called threadIdx and blockDim, respectively; one difference is that these return a 3-dimensional structure with fields x, y, and z to simplify cartesian indexing for up to 3-dimensional arrays. Consequently we can assign unique work in the following way:

function gpu_add2!(y, x)
     index = threadIdx().x    # this example only requires linear indexing, so just use `x`
@@ -162,21 +162,21 @@
     CUDA.@sync begin
         @cuda threads=256 blocks=numblocks gpu_add3!(y, x)
     end
-end
bench_gpu3! (generic function with 1 method)
@btime bench_gpu3!($y_d, $x_d)
  67.268 μs (52 allocations: 1.31 KiB)

Finally, we've achieved the similar performance to what we got with the broadcasted version. Let's profile again to confirm this launch configuration:

CUDA.@profile trace=true bench_gpu3!(y_d, x_d)
Profiler ran for 13.86 ms, capturing 302 events.
+end
bench_gpu3! (generic function with 1 method)
@btime bench_gpu3!($y_d, $x_d)
  67.268 μs (52 allocations: 1.31 KiB)

Finally, we've achieved the similar performance to what we got with the broadcasted version. Let's profile again to confirm this launch configuration:

CUDA.@profile trace=true bench_gpu3!(y_d, x_d)
Profiler ran for 14.76 ms, capturing 296 events.
 
-Host-side activity: calling CUDA APIs took 98.47 µs (0.71% of the trace)
+Host-side activity: calling CUDA APIs took 101.57 µs (0.69% of the trace)
 ┌─────┬──────────┬──────────┬─────────────────────┐
 │  ID │    Start │     Time │ Name                │
 ├─────┼──────────┼──────────┼─────────────────────┤
-│  21 │ 13.66 ms │ 41.72 µs │ cuLaunchKernel      │
-│ 298 │ 13.84 ms │  5.48 µs │ cuStreamSynchronize │
+│  21 │ 14.55 ms │ 46.73 µs │ cuLaunchKernel      │
+│ 292 │ 14.74 ms │  5.25 µs │ cuStreamSynchronize │
 └─────┴──────────┴──────────┴─────────────────────┘
 
-Device-side activity: GPU was busy for 131.37 µs (0.95% of the trace)
+Device-side activity: GPU was busy for 131.61 µs (0.89% of the trace)
 ┌────┬─────────┬───────────┬─────────┬────────┬──────┬──────────────────────────
 │ ID │   Start │      Time │ Threads │ Blocks │ Regs │ Name                    ⋯
 ├────┼─────────┼───────────┼─────────┼────────┼──────┼──────────────────────────
-│ 21 │ 13.7 ms │ 131.37 µs │     256 │   4096 │   40 │ _Z9gpu_add3_13CuDeviceA ⋯
+│ 21 │ 14.6 ms │ 131.61 µs │     256 │   4096 │   40 │ _Z9gpu_add3_13CuDeviceA ⋯
 └────┴─────────┴───────────┴─────────┴────────┴──────┴──────────────────────────
                                                                 1 column omitted

In the previous example, the number of threads was hard-coded to 256. This is not ideal, as using more threads generally improves performance, but the maximum number of allowed threads to launch depends on your GPU as well as on the kernel. To automatically select an appropriate number of threads, it is recommended to use the launch configuration API. This API takes a compiled (but not launched) kernel, returns a tuple with an upper bound on the number of threads, and the minimum number of blocks that are required to fully saturate the GPU:

kernel = @cuda launch=false gpu_add3!(y_d, x_d)
@@ -227,4 +227,4 @@
  [1] throw_boundserror at abstractarray.jl:484
  [2] checkbounds at abstractarray.jl:449
  [3] setindex! at /home/tbesard/Julia/CUDA/src/device/array.jl:79
- [4] some_kernel at /tmp/tmpIMYANH:6
Warning

On older GPUs (with a compute capability below sm_70) these errors are fatal, and effectively kill the CUDA environment. On such GPUs, it's often a good idea to perform your "sanity checks" using code that runs on the CPU and only turn over the computation to the GPU once you've deemed it to be safe.

Summary

Keep in mind that the high-level functionality of CUDA often means that you don't need to worry about writing kernels at such a low level. However, there are many cases where computations can be optimized using clever low-level manipulations. Hopefully, you now feel comfortable taking the plunge.

This page was generated using Literate.jl.

+ [4] some_kernel at /tmp/tmpIMYANH:6
Warning

On older GPUs (with a compute capability below sm_70) these errors are fatal, and effectively kill the CUDA environment. On such GPUs, it's often a good idea to perform your "sanity checks" using code that runs on the CPU and only turn over the computation to the GPU once you've deemed it to be safe.

Summary

Keep in mind that the high-level functionality of CUDA often means that you don't need to worry about writing kernels at such a low level. However, there are many cases where computations can be optimized using clever low-level manipulations. Hopefully, you now feel comfortable taking the plunge.

This page was generated using Literate.jl.

diff --git a/dev/tutorials/performance/index.html b/dev/tutorials/performance/index.html index ccd4c51492..dc8dc5f121 100644 --- a/dev/tutorials/performance/index.html +++ b/dev/tutorials/performance/index.html @@ -53,4 +53,4 @@ blocks = cld(length(y), threads) CUDA.@sync kernel(y, x; threads, blocks) -end
bench_gpu5! (generic function with 1 method)
@btime bench_gpu4!($y_d, $x_d)
  76.149 ms (57 allocations: 3.70 KiB)
@btime bench_gpu5!($y_d, $x_d)
  75.732 ms (58 allocations: 3.73 KiB)

This benchmark shows there is a only a small performance benefit for this kernel however we can see a big difference in the amount of registers used, recalling that 28 registers were used when using a StepRange:

CUDA.registers(@cuda gpu_add5!(y_d, x_d))
  12

This page was generated using Literate.jl.

+end
bench_gpu5! (generic function with 1 method)
@btime bench_gpu4!($y_d, $x_d)
  76.149 ms (57 allocations: 3.70 KiB)
@btime bench_gpu5!($y_d, $x_d)
  75.732 ms (58 allocations: 3.73 KiB)

This benchmark shows there is a only a small performance benefit for this kernel however we can see a big difference in the amount of registers used, recalling that 28 registers were used when using a StepRange:

CUDA.registers(@cuda gpu_add5!(y_d, x_d))
  12

This page was generated using Literate.jl.

diff --git a/dev/usage/array/index.html b/dev/usage/array/index.html index 6a9fb27230..a5b97fd9f4 100644 --- a/dev/usage/array/index.html +++ b/dev/usage/array/index.html @@ -250,4 +250,4 @@ julia> fft(a) 2×2 CuArray{ComplexF32, 2, CUDA.DeviceMemory}: 2.6692+0.0im 0.65323+0.0im - -1.11072+0.0im 0.749168+0.0im + -1.11072+0.0im 0.749168+0.0im diff --git a/dev/usage/memory/index.html b/dev/usage/memory/index.html index 8ab12518f2..34a0ed7a47 100644 --- a/dev/usage/memory/index.html +++ b/dev/usage/memory/index.html @@ -91,4 +91,4 @@ println("Batch $batch: ", a .+ b) end Batch 1: [3] -Batch 2: [7]

For each batch, every argument (assumed to be an array-like) is uploaded to the GPU using the adapt mechanism from above. Afterwards, the memory is eagerly put back in the CUDA memory pool using unsafe_free! to lower GC pressure.

+Batch 2: [7]

For each batch, every argument (assumed to be an array-like) is uploaded to the GPU using the adapt mechanism from above. Afterwards, the memory is eagerly put back in the CUDA memory pool using unsafe_free! to lower GC pressure.

diff --git a/dev/usage/multigpu/index.html b/dev/usage/multigpu/index.html index efb6ef7080..a16859a3cf 100644 --- a/dev/usage/multigpu/index.html +++ b/dev/usage/multigpu/index.html @@ -46,4 +46,4 @@ using Test c = Array(d_c) -@test a+b ≈ c +@test a+b ≈ c diff --git a/dev/usage/multitasking/index.html b/dev/usage/multitasking/index.html index c4f202afd2..5a25503755 100644 --- a/dev/usage/multitasking/index.html +++ b/dev/usage/multitasking/index.html @@ -73,4 +73,4 @@ # comparison results[1] == results[2] -end

By using the Threads.@spawn macro, the tasks will be scheduled to be run on different CPU threads. This can be useful when you are calling a lot of operations that "block" in CUDA, e.g., memory copies to or from unpinned memory. The same result will occur when using a Threads.@threads for ... end block. Generally, though, operations that synchronize GPU execution (including the call to synchronize itself) are implemented in a way that they yield back to the Julia scheduler, to enable concurrent execution without requiring the use of different CPU threads.

Warning

Use of multiple threads with CUDA.jl is a recent addition, and there may still be bugs or performance issues.

+end

By using the Threads.@spawn macro, the tasks will be scheduled to be run on different CPU threads. This can be useful when you are calling a lot of operations that "block" in CUDA, e.g., memory copies to or from unpinned memory. The same result will occur when using a Threads.@threads for ... end block. Generally, though, operations that synchronize GPU execution (including the call to synchronize itself) are implemented in a way that they yield back to the Julia scheduler, to enable concurrent execution without requiring the use of different CPU threads.

Warning

Use of multiple threads with CUDA.jl is a recent addition, and there may still be bugs or performance issues.

diff --git a/dev/usage/overview/index.html b/dev/usage/overview/index.html index 774c630475..448673767b 100644 --- a/dev/usage/overview/index.html +++ b/dev/usage/overview/index.html @@ -31,4 +31,4 @@ @show capability(device) end

If such high-level wrappers are missing, you can always access the underling C API (functions and structures prefixed with cu) without having to ever exit Julia:

version = Ref{Cint}()
 CUDA.cuDriverGetVersion(version)
-@show version[]
+@show version[] diff --git a/dev/usage/workflow/index.html b/dev/usage/workflow/index.html index 74ab937460..0beb275192 100644 --- a/dev/usage/workflow/index.html +++ b/dev/usage/workflow/index.html @@ -35,4 +35,4 @@ 2 julia> CUDA.@allowscalar a[1] += 1 -3 +3