KernelInterface

KernelInterface

The KernelInterface (or KI) module defines the API interface for backends to define various lower-level device and host-side functionality. The KI interface is used to define the higher-level device-side functionality in KernelAbstractions.

Both provide APIs for host and device-side functionality, but KI focuses on lower-level functionality that is shared amongst backends, while KernelAbstractions provides higher-level functionality such as writing kernels that work on arrays with an arbitrary number of dimensions, or convenience functions like allocating arrays on a backend.

Backend

Abstract supertype for all KernelAbstractions backends.

Concrete backends (for example CUDABackend from CUDA.jl or CPU from KernelAbstractions) determine where arrays are allocated and where kernels execute. Use get_backend to obtain the backend for an array and allocate to create storage on a backend.

Example

backend = get_backend(A)
kernel = my_kernel(backend, 256)
kernel(A, ndrange=length(A))
synchronize(backend)

Abstract type for all GPU based KernelAbstractions backends.

get_backend(A::AbstractArray)::Backend

Get a Backend instance suitable for array A.

get_global_size()::@NamedTuple{x::Int, y::Int, z::Int}

Return the number of global work-items specified.

@device_override get_global_size()::@NamedTuple{x::Int, y::Int, z::Int}

get_global_id()::@NamedTuple{x::Int, y::Int, z::Int}

Returns the unique global work-item ID.

@device_override get_global_id()::@NamedTuple{x::Int, y::Int, z::Int}

get_local_size()::@NamedTuple{x::Int, y::Int, z::Int}

Return the number of local work-items specified.

@device_override get_local_size()::@NamedTuple{x::Int, y::Int, z::Int}

get_local_id()::@NamedTuple{x::Int, y::Int, z::Int}

Returns the unique local work-item ID.

@device_override get_local_id()::@NamedTuple{x::Int, y::Int, z::Int}

get_num_groups()::@NamedTuple{x::Int, y::Int, z::Int}

Returns the number of groups.

@device_override get_num_groups()::@NamedTuple{x::Int, y::Int, z::Int}

get_group_id()::@NamedTuple{x::Int, y::Int, z::Int}

Returns the unique group ID.

@device_override get_group_id()::@NamedTuple{x::Int, y::Int, z::Int}

get_sub_group_size()::UInt32

Returns the number of work-items in the sub-group.

@device_override get_sub_group_size()::UInt32

get_max_sub_group_size()::UInt32

Returns the maximum sub-group size for sub-groups in the current workgroup.

@device_override get_max_sub_group_size()::UInt32

get_num_sub_groups()::UInt32

Returns the number of sub-groups in the current workgroup.

@device_override get_num_sub_groups()::UInt32

get_sub_group_id()::UInt32

Returns the sub-group ID within the work-group.

@device_override get_sub_group_id()::UInt32

get_sub_group_local_id()::UInt32

Returns the work-item ID within the current sub-group.

@device_override get_sub_group_local_id()::UInt32

barrier()

After a barrier() call, all read and writes to global and local memory from each thread in the workgroup are visible in from all other threads in the workgroup.

This does not guarantee that a write from a thread in a certain workgroup will be visible to a thread in a different workgroup.

@device_override barrier()

sub_group_barrier()

After a sub_group_barrier() call, all read and writes to global and local memory from each thread in the sub-group are visible in from all other threads in the sub-group.

This does not guarantee that a write from a thread in a certain sub-group will be visible to a thread in a different sub-group.

@device_override sub_group_barrier()

localmemory(::Type{T}, dims)

Declare memory that is local to a workgroup.

@device_override localmemory(::Type{T}, ::Val{Dims}) where {T, Dims}

shfl_down(val::T, offset::Integer) where T

Read val from a lane with higher id given by offset.

@device_override shfl_down(val::T, offset::Integer) where T

shfl_down_types(::Backend)::Vector{DataType}

Returns a vector of DataTypes supported on backend

_print(args...)

Overloaded by backends to enable `KernelAbstractions.@print`
functionality.

@device_override _print(args...)

The generic fallback prints on the host, which keeps CPU backends working. Val arguments are unwrapped, since KernelAbstractions.@print uses them to pass literal strings through to backends that require compile-time format strings.

allocate(::Backend, Type, dims...; unified=false)::AbstractArray

Allocate a storage array appropriate for the computational backend. unified=true allocates an array using unified memory if the backend supports it and throws otherwise. Use supports_unified to determine whether it is supported by a backend.

zeros(::Backend, Type, dims...; unified=false)::AbstractArray

Allocate a storage array appropriate for the computational backend filled with zeros. unified=true allocates an array using unified memory if the backend supports it and throws otherwise.

ones(::Backend, Type, dims...; unified=false)::AbstractArray

Allocate a storage array appropriate for the computational backend filled with ones. unified=true allocates an array using unified memory if the backend supports it and throws otherwise.

copyto!(::Backend, dest::AbstractArray, src::AbstractArray)

Perform an asynchronous copyto! operation that is execution ordered with respect to the back-end.

For most users, Base.copyto! should suffice, performance a simple, synchronous copy. Only when you know you need asynchronicity w.r.t. the host, you should consider using this asynchronous version, which requires additional lifetime guarantees as documented below.

arr = zeros(64)
GC.@preserve arr begin
    copyto!(backend, arr, ...)
    # other operations
    synchronize(backend)
end

pagelock!(::Backend, dest::AbstractArray)::Union{Nothing, Missing}

Pagelock (pin) a host memory buffer for a backend device. This may be necessary for copyto! to perform asynchronously w.r.t to the host/

This function should return nothing; or missing if not implemented.

unsafe_free!(x::AbstractArray)

Release the memory of an array for reuse by future allocations and reduce pressure on the allocator. After releasing the memory of an array, it should no longer be accessed.

synchronize(::Backend)

Synchronize the current backend.

priority!(::Backend, prio::Symbol)::Nothing

Set the priority for the backend stream/queue. This is an optional feature that backends may or may not implement. If a backend shall support priorities it must accept :high, :normal, :low. Where :normal is the default.

device(backend::Backend)::Int

Return the 1-based index of the currently active device for backend.

ndevices(backend::Backend)::Int

Return the number of devices available to backend.

device!(backend::Backend, id::Int)::Nothing

Select the active device for backend. id is a 1-based device index and must satisfy 1 <= id <= ndevices(backend).

Example

device!(CUDABackend(), 2)  # use the second CUDA device

functional(::Backend)::Union{Bool, Missing}

Queries if the provided backend is functional. This may mean different things for different backends, but generally should mean that the necessary drivers and a compute device are available.

This function should return a Bool or missing if not implemented.

versioninfo(io::IO=stdout, backend::Backend)::Nothing

Print information about backend to io. It is up to the backends to determine what is relevant.

supports_unified(::Backend)::Bool

Returns whether unified memory arrays are supported by the backend.

supports_atomics(::Backend)::Bool

Returns whether @atomic operations are supported by the backend.

supports_float64(::Backend)::Bool

Returns whether Float64 values are supported by the backend.

max_work_group_size(backend, kern; [max_work_items::Int])::Int

The maximum workgroup size limit for a kernel as reported by the backend. This function represents a theoretical maximum; kernel_max_work_group_size should be used before launching a kernel as some backends may error if kernel launch with too big a workgroup is attempted.

max_work_group_size(backend::NewBackend)::Int

sub_group_size(backend)::Int

Returns a reasonable sub-group size supported by the currently active device for the specified backend. This would typically be 32, or 64 for devices that don't support 32.

sub_group_size(backend::NewBackend)::Int

multiprocessor_count(backend::NewBackend)::Int

The multiprocessor count for the current device used by backend. Used for certain algorithm optimizations.

multiprocessor_count(backend::NewBackend)::Int

Kernel{Backend, Kern}

Kernel closure struct that is used to represent the backend kernel on the host.

(kernel::Kernel{<:NewBackend})(args...; numworkgroups=(), workgroupsize=(), ndrange=(), max_work_group_size=typemax(Int))

KI.kernel_function(::NewBackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}

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

Currently, kernel_function only supports the name keyword argument as it is the only one by all backends.

Keyword arguments:

• name: override the name that the kernel will have in the generated code

kernel_function(::NewBackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}

kernel_max_work_group_size(kern; [max_work_items::Int])::Int

The maximum workgroup size limit for a kernel as reported by the backend. This function should always be used to determine the workgroup size before launching a kernel.

kernel_max_work_group_size(kern::Kernel{<:NewBackend}; max_work_items::Int=typemax(Int))::Int

argconvert(::NewBackend, arg)

This function is called for every argument to be passed to a kernel, converting them to their device side representation.

argconvert(::NewBackend, arg)

KI.@kernel backend [workgroupsize=... numworkgroups=... ndrange=...] [kwargs...] func(args...)

High-level interface for executing code on a GPU.

The KI.@kernel macro should prefix a call, with func a callable function or object that should return nothing. It will be compiled to a function native to the specified backend upon first use, and to a certain extent arguments will be converted and managed automatically using argconvert. Finally, if launch=true, the newly created callable kernel object is called and launched according to the specified backend.

There are a few keyword arguments that influence the behavior of KI.@kernel:

• 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.
• name: the name of the kernel in the generated code. Defaults to an automatically- generated name.