Numba CUDA

https://nvidia.github.io/numba-cuda/user/
Admin
Numba CUDA enables Python developers to write and execute high-performance code on NVIDIA GPUs using...
Tokens:59,602
Snippets:401
Trust Score:8.3
Update:10 months ago
Show doc for...
Tokens:
Raw
### Numba CUDA Random Number Generation Example

Source: https://nvidia.github.io/numba-cuda/user/index

Illustrates how to use Numba's support for random number generation on the GPU. Includes a simple example and a more complex one demonstrating management of RNG state size and 3D grids.

```python
from numba import cuda
from numba.cuda.random import create_xoroshiro128p_states
import numpy as np

@cuda.jit
def rng_kernel(rng_states, data):
    thread_id = cuda.grid(1)
    # Generate a random float between 0 and 1
    random_value = cuda.random.xoroshiro128p_uniform_float64(rng_states, thread_id)
    data[thread_id] = random_value

# Example usage
grid_size = 1024
block_size = 128

# Create RNG states for the grid
rng_states = create_xoroshiro128p_states(grid_size * block_size, seed=1)
output_data = cuda.device_array(grid_size * block_size, dtype=np.float64)

rng_kernel[grid_size, block_size](rng_states, output_data)

# Copy results back to host to view
# host_output = output_data.copy_to_host()
# print(host_output[:10])

```

--------------------------------

### Numba CUDA Installation and Configuration

Source: https://nvidia.github.io/numba-cuda/user/index

Instructions for installing Numba with CUDA support, including requirements for GPUs and CUDA Toolkits. It also covers configuration options for CUDA bindings, driver, and toolkit search paths.

```bash
# Example installation command (may vary based on environment)
pip install numba

```

```python
# Example configuration: Setting CUDA driver and toolkit paths
import os
os.environ['NUMBA_CUDA_DRIVER_PATH'] = '/usr/local/cuda/lib64/stubs/libcuda.so'
os.environ['NUMBA_CUDA_TOOLKIT_PATH'] = '/usr/local/cuda'

```

--------------------------------

### Numba CUDA GPU Reduction Example

Source: https://nvidia.github.io/numba-cuda/user/index

Provides an example of implementing a GPU reduction operation using Numba's `Reduce` class. Reductions are common parallel patterns for aggregating data on the GPU.

```python
from numba import cuda
from numba.cuda.reductions import Reduce
import numpy as np

# Define a reduction operation (e.g., sum)
reduce_sum = Reduce(lambda x, y: x + y)

@cuda.jit
def gpu_reduction_kernel(input_data, output_data):
    # This is a simplified representation. Actual reduction involves shared memory.
    # The Reduce class abstracts much of this complexity.
    pass # Actual implementation uses the Reduce class internally

# Example usage
# data_host = np.array([...])
# data_device = cuda.to_device(data_host)
# result_device = cuda.device_array(1, dtype=data_host.dtype)
# reduce_sum(data_device, result_device, stream=0) # Simplified call
# result_host = result_device.copy_to_host()

```

--------------------------------

### Example Usage of Naive Matrix Multiplication

Source: https://nvidia.github.io/numba-cuda/user/examples

Demonstrates how to use the `matmul` kernel. It involves creating NumPy arrays, transferring them to the CUDA device, configuring and launching the kernel, and copying the result back to the host.

```python
x_h = np.arange(16).reshape([4, 4])
y_h = np.ones([4, 4])
z_h = np.zeros([4, 4])

x_d = cuda.to_device(x_h)
y_d = cuda.to_device(y_h)
z_d = cuda.to_device(z_h)

threadsperblock = (16, 16)
blockspergrid_x = math.ceil(z_h.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(z_h.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)

matmul[blockspergrid, threadsperblock](x_d, y_d, z_d)
z_h = z_d.copy_to_host()
print(z_h)
print(x_h @ y_h)
```

--------------------------------

### Example Usage of Fast Matrix Multiplication

Source: https://nvidia.github.io/numba-cuda/user/examples

Shows how to use the `fast_matmul` kernel with shared memory. This example mirrors the naive usage but launches the optimized kernel, demonstrating the practical application of shared memory optimization.

```python
x_h = np.arange(16).reshape([4, 4])
y_h = np.ones([4, 4])
z_h = np.zeros([4, 4])

x_d = cuda.to_device(x_h)
y_d = cuda.to_device(y_h)
z_d = cuda.to_device(z_h)

threadsperblock = (TPB, TPB)
blockspergrid_x = math.ceil(z_h.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(z_h.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)

fast_matmul[blockspergrid, threadsperblock](x_d, y_d, z_d)
z_h = z_d.copy_to_host()
print(z_h)
print(x_h @ y_h)
```

--------------------------------

### Numba CUDA Array Interface (Version 3) Example

Source: https://nvidia.github.io/numba-cuda/user/index

Demonstrates the use of the CUDA Array Interface (Version 3) for interoperability. Shows how to consume arrays using `from_cuda_array_interface` and produce arrays using `as_cuda_array`.

```python
from numba import cuda

# Assume 'my_array_interface' is a dictionary adhering to the CUDA Array Interface spec
# Example structure:
# my_array_interface = {
#     "version": 3,
#     "shape": (10,),
#     "typestr": "f8",
#     "data": (device_ptr, is_c_contig), # device_ptr is a GPU pointer
#     "strides": (8,),
#     "version": 3
# }

# Consuming an array via the interface
# cuda_array = cuda.from_cuda_array_interface(my_array_interface)

# Producing an array (example of getting the interface from a Numba CUDA array)
# original_cuda_array = cuda.to_device(np.arange(10))
# array_interface_output = original_cuda_array.as_cuda_array_interface()

```

--------------------------------

### Numba CUDA Kernel Launch Example - CuPy

Source: https://nvidia.github.io/numba-cuda/user/cuda_array_interface

Demonstrates launching a Numba CUDA kernel ('add') to perform element-wise addition on CuPy arrays. This example illustrates how CuPy arrays can act as producers for Numba CUDA kernels, utilizing the CUDA Array Interface for data exchange.

```Python
import cupy
from numba import cuda

@cuda.jit
def add(x, y, out):
    start = cuda.grid(1)
    stride = cuda.gridsize(1)
    for i in range(start, x.shape[0], stride):
        out[i] = x[i] + y[i]

a = cupy.arange(10)
b = a * 2
out = cupy.zeros_like(a)

add[1, 32](a, b, out)
```

--------------------------------

### CUDA Kernel Calling NumPy UFunc (np.sin)

Source: https://nvidia.github.io/numba-cuda/user/examples

Illustrates how to call NumPy universal functions (ufuncs) like `np.sin` from within a Numba CUDA kernel. The output array must be passed as a positional argument. This example includes kernel definition, array setup, kernel launch, and a sanity check against NumPy's `np.sin`.

```python
import numpy as np
from numba import cuda

@cuda.jit
def f(r, x):
    # Compute sin(x) with result written to r
    np.sin(x, r)

# Declare input and output arrays
x = np.arange(10, dtype=np.float32) - 5
r = np.zeros_like(x)

# Launch kernel that calls the ufunc
f[1, 1](r, x)

# A quick sanity check demonstrating equality of the sine computed by
# the sin ufunc inside the kernel, and NumPy's sin ufunc
np.testing.assert_allclose(r, np.sin(x))
```

--------------------------------

### Launch CUDA Kernel using `forall`

Source: https://nvidia.github.io/numba-cuda/user/examples

Launches the Numba CUDA kernel `f` using the `forall()` method, which automatically configures a 1D grid based on the data size. The results are then copied back to the host and printed. This is a simple way to launch kernels.

```python
f.forall(len(a))(a, b, c)
print(c.copy_to_host())
```

--------------------------------

### Numba CUDA Cooperative Kernel Launch Example

Source: https://nvidia.github.io/numba-cuda/user/cooperative_groups

Demonstrates launching a Numba CUDA kernel (`sequential_rows`) with cooperative grid and block dimensions. The launch is implicitly cooperative because the kernel uses `g.sync()`.

```python
# Empty input data
A = np.zeros((1024, 1024), dtype=np.int32)
# A somewhat arbitrary choice (one warp), but generally smaller block sizes
# allow more blocks to be launched (noting that other limitations on
# occupancy apply such as shared memory size)
blockdim = 32
griddim = A.shape[1] // blockdim

# Kernel launch - this is implicitly a cooperative launch
sequential_rows[griddim, blockdim](A)

# What do the results look like?
# print(A)
#
# [[   0    0    0 ...    0    0    0]
#  [   1    1    1 ...    1    1    1]
#  [   2    2    2 ...    2    2    2]
#  ...
#  [1021 1021 1021 ... 1021 1021 1021]
#  [1022 1022 1022 ... 1022 1022 1022]
#  [1023 1023 1023 ... 1023 1023 1023]]

```

--------------------------------

### Import NumPy and Numba CUDA

Source: https://nvidia.github.io/numba-cuda/user/examples

Imports necessary libraries, NumPy for array manipulation and Numba's CUDA module for GPU programming. These are foundational for all subsequent examples.

```python
import numpy as np
from numba import cuda
```

--------------------------------

### Writing and Invoking CUDA Kernels with Numba

Source: https://nvidia.github.io/numba-cuda/user/index

Demonstrates how to declare and invoke CUDA kernels using Numba's `@cuda.jit` decorator. Covers kernel declaration, invocation strategies, and considerations for block size and multi-dimensional grids.

```python
from numba import cuda
import numpy as np

@cuda.jit
def add_kernel(x, y, out):
    idx = cuda.grid(1)
    if idx < x.shape[0]:
        out[idx] = x[idx] + y[idx]

# Example kernel invocation
input_size = 100000
x = np.arange(input_size, dtype=np.float32)
y = np.arange(input_size, dtype=np.float32)
out = np.zeros_like(x)

# Configure kernel launch
threadsperblock = 128
blockspergrid = (input_size + (threadsperblock - 1)) // threadsperblock

# Launch kernel
add_kernel[blockspergrid, threadsperblock](x, y, out)

print(out[:10])

```

--------------------------------

### Example Foreign Device Function Prototype (C/C++)

Source: https://nvidia.github.io/numba-cuda/user/cuda_ffi

An example illustrating the device function ABI for a function that multiplies two 32-bit floats and returns a 32-bit float. This demonstrates the structure required for foreign functions to be correctly called by Numba kernels.

```c++
extern "C"
__device__ int
mul_f32_f32(
  float* return_value,
  float x,
  float y
);

```

--------------------------------

### Execute and Verify CUDA Kernel

Source: https://nvidia.github.io/numba-cuda/user/examples

Launches the CUDA reduction kernel on the device and prints the result stored in the first element of the modified array. It also verifies the result by comparing it with the sum computed on the host.

```python
array_sum[1, nelem](a)
print(a[0])  # 523776
print(sum(np.arange(1024)))  # 523776
```

--------------------------------

### Using @reduce decorator for GPU reduction

Source: https://nvidia.github.io/numba-cuda/user/reduction

Demonstrates how to use the @cuda.reduce decorator to convert a binary operation into a reduction kernel for CUDA GPUs. Includes an example for summing array elements.

```APIDOC
## POST /cuda/reduce

### Description
Applies a reduction operation on a CUDA array using a binary function defined by the `@cuda.reduce` decorator.

### Method
POST (implicitly, as the decorator creates a callable object)

### Endpoint
`/cuda/reduce` (conceptual endpoint for the reduction operation)

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
* **arr** (numpy.ndarray or numba.cuda.cudadrv.devicearray.DeviceNDArray) - Required - The input array (host or device) to be reduced.
* **size** (int) - Optional - The number of elements in `arr` to reduce. If not specified, the entire array is reduced.
* **res** (numba.cuda.cudadrv.devicearray.DeviceNDArray) - Optional - A device array to store the reduction result. The result is written to the first element. If provided, the function returns None.
* **init** (numeric) - Optional - The initial value for the reduction. Its type must match `arr.dtype`. Defaults to 0.
* **stream** (int) - Optional - The CUDA stream in which to perform the reduction. Defaults to 0.

### Request Example
```python
import numpy
from numba import cuda

@cuda.reduce
def sum_reduce(a, b):
    return a + b

A = (numpy.arange(1234, dtype=numpy.float64)) + 1

# Perform reduction
got = sum_reduce(A)
```

### Response
#### Success Response (200)
- **result** (numeric) - The result of the reduction operation if `res` is not specified.

#### Response Example
```json
{
  "result": 772429179.0
}
```
```

--------------------------------

### Execute Numba JIT Function on CPU

Source: https://nvidia.github.io/numba-cuda/user/examples

Demonstrates the standalone execution of the `business_logic` function, which was compiled using `@numba.jit`, on the CPU. It takes integer arguments and prints the resulting float value. This confirms the CPU compatibility of the decorated function.

```python
print(business_logic(1, 2, 3))
```

--------------------------------

### Create and Transfer Arrays to Device Memory

Source: https://nvidia.github.io/numba-cuda/user/examples

Demonstrates how to create NumPy arrays, transfer them to the GPU using `cuda.to_device()`, and allocate an uninitialized device array for results using `cuda.device_array_like()`. This prepares data for GPU computation.

```python
N = 100000
a = cuda.to_device(np.random.random(N))
b = cuda.to_device(np.random.random(N))
c = cuda.device_array_like(a)
```

--------------------------------

### Import Numba CUDA Libraries

Source: https://nvidia.github.io/numba-cuda/user/examples

Imports necessary libraries from Numba and NumPy for CUDA programming, including random number generation utilities.

```python
import numba
import numpy as np
from numba import cuda
from numba.cuda.random import (
    create_xoroshiro128p_states,
    xoroshiro128p_uniform_float32,
)
```

--------------------------------

### Numba CUDA JIT Function CPU-GPU Compatibility

Source: https://nvidia.github.io/numba-cuda/user/index

Explains how Numba's JIT compilation can make functions compatible for execution on both the CPU and GPU. This allows for code reuse and easier development for heterogeneous systems.

```python
from numba import cuda, jit
import numpy as np

# Function compiled for both CPU and GPU
@jit(nopython=True)
def cpu_gpu_compatible_add(x, y):
    return x + y

# --- CPU execution ---
cpu_result = cpu_gpu_compatible_add(np.array([1, 2]), np.array([3, 4]))
# print(cpu_result)

# --- GPU execution (requires CUDA enabled Numba) ---
# try:
#     @cuda.jit(device=True)
#     def gpu_kernel_add(x, y, out):
#         idx = cuda.grid(1)
#         out[idx] = cpu_gpu_compatible_add(x[idx], y[idx])
#
#     # ... (kernel invocation setup as shown in other examples) ...
# except Exception as e:
#     print(f"CUDA not available or error during GPU execution: {e}")

```

--------------------------------

### Initialize 3D Array with Random Numbers using Numba CUDA 3D Grid

Source: https://nvidia.github.io/numba-cuda/user/random

This example demonstrates initializing a large 3D array with random numbers using Numba's CUDA capabilities. It employs strided loops and a 3D grid to manage the number of RNG states efficiently, preventing excessive memory usage and improving GPU utilization. The function takes the 3D array and RNG states as input.

```python
from numba import cuda
from numba.cuda.random import (
    create_xoroshiro128p_states,
    xoroshiro128p_uniform_float32,
)
import numpy as np

@cuda.jit
def random_3d(arr, rng_states):
    # Per-dimension thread indices and strides
    startx, starty, startz = cuda.grid(3)
    stridex, stridey, stridez = cuda.gridsize(3)

    # Linearized thread index
    tid = (startz * stridey * stridex) + (starty * stridex) + startx

    # Use strided loops over the array to assign a random value to each entry
    for i in range(startz, arr.shape[0], stridez):
        for j in range(starty, arr.shape[1], stridey):
            for k in range(startx, arr.shape[2], stridex):
                arr[i, j, k] = xoroshiro128p_uniform_float32(
                    rng_states, tid
                )

# Array dimensions
X, Y, Z = 701, 900, 719

# Block and grid dimensions
bx, by, bz = 8, 8, 8
gx, gy, gz = 16, 16, 16

# Total number of threads
nthreads = bx * by * bz * gx * gy * gz

# Initialize a state for each thread
rng_states = create_xoroshiro128p_states(nthreads, seed=1)

# Generate random numbers
arr = cuda.device_array((X, Y, Z), dtype=np.float32)
random_3d[(gx, gy, gz), (bx, by, bz)](arr, rng_states)
```

--------------------------------

### Initialize Data and Device Buffers for Heat Equation

Source: https://nvidia.github.io/numba-cuda/user/examples

Sets up the initial state for the 1D heat equation simulation. It creates a NumPy array representing temperatures, initializes a central point to a high value, and transfers this data to CUDA device arrays (buffers). An auxiliary buffer is also created for synchronization.

```python
# Use an odd problem size.
# This is so there can be an element truly in the "middle" for symmetry.
size = 1001
data = np.zeros(size)

# Middle element is made very hot
data[500] = 10000
buf_0 = cuda.to_device(data)

# This extra array is used for synchronization purposes
buf_1 = cuda.device_array_like(buf_0)

niter = 10000
```

--------------------------------

### Define CUDA Vector Addition Kernel with Numba

Source: https://nvidia.github.io/numba-cuda/user/examples

Defines a Numba CUDA kernel for vector addition. The kernel operates on device arrays and uses `cuda.grid(1)` to get the thread ID. Output is written to a provided array `c`, as Numba kernels do not return values.

```python
@cuda.jit
def f(a, b, c):
    # like threadIdx.x + (blockIdx.x * blockDim.x)
    tid = cuda.grid(1)
    size = len(c)

    if tid < size:
        c[tid] = a[tid] + b[tid]
```

--------------------------------

### PTX return value handling using C ABI for 'add' function

Source: https://nvidia.github.io/numba-cuda/user/cuda_compilation

Demonstrates how the return value of the 'add' function is handled in PTX when compiled with the C ABI. The result is directly stored into the function's return parameter.

```ptx
add.s32        %r3, %r2, %r1;
st.param.b32   [func_retval0+0], %r3;


```

--------------------------------

### Numba CUDA Kernel for Sessionization

Source: https://nvidia.github.io/numba-cuda/user/examples

A Numba CUDA kernel function `sessionize` that takes user IDs, timestamps, and a results array as input. It identifies session boundaries based on changes in user ID or a time gap exceeding `session_timeout`. It then assigns a `session_id` to each data point, marking the start of a new session with the index of the first data point in that session.

```python
@cuda.jit
def sessionize(user_id, timestamp, results):
    gid = cuda.grid(1)
    size = len(user_id)

    if gid >= size:
        return

    # Determine session boundaries
    is_first_datapoint = gid == 0
    if not is_first_datapoint:
        new_user = user_id[gid] != user_id[gid - 1]
        timed_out = (
            timestamp[gid] - timestamp[gid - 1] > session_timeout
        )
        is_sess_boundary = new_user or timed_out
    else:
        is_sess_boundary = True

    # Determine session labels
    if is_sess_boundary:
        # This thread marks the start of a session
        results[gid] = gid

        # Make sure all session boundaries are written
        # before populating the session id
        grid = cuda.cg.this_grid()
        grid.sync()

        look_ahead = 1
        # Check elements 'forward' of this one
        # until a new session boundary is found
        while results[gid + look_ahead] == 0:
            results[gid + look_ahead] = gid
            look_ahead += 1
            # Avoid out-of-bounds accesses by the last thread
            if gid + look_ahead == size - 1:
                results[gid + look_ahead] = gid
                break
```

--------------------------------

### Environment Variable Configuration

Source: https://nvidia.github.io/numba-cuda/user/external-memory

Explains how to configure an EMM Plugin using the NUMBA_CUDA_MEMORY_MANAGER environment variable.

```APIDOC
## Environment Variable Configuration for EMM Plugin

### Description
Sets an environment variable `NUMBA_CUDA_MEMORY_MANAGER` to a module name. Numba will import this module and use its `_numba_memory_manager` global variable as the memory manager class. This is useful for testing Numba with an EMM Plugin.

### Example
```bash
$ NUMBA_CUDA_MEMORY_MANAGER=rmm pytest --pyargs numba.cuda.tests -v
```
```

--------------------------------

### Install Numba CUDA with Conda (CUDA 12)

Source: https://nvidia.github.io/numba-cuda/user/installation

Installs the numba-cuda package and CUDA 12 toolkit into a Conda environment. This is a convenient method for managing CUDA dependencies within a project.

```bash
$ conda install -c conda-forge numba-cuda "cuda-version=12"
```

--------------------------------

### Install Numba CUDA Dependencies with Pip (CUDA 12)

Source: https://nvidia.github.io/numba-cuda/user/installation

Installs Numba CUDA with support for CUDA 12 dependencies directly from PyPI using pip. This method is suitable for environments where Conda is not used.

```bash
$ pip install numba-cuda[cu12]
```

--------------------------------

### PTX prototype using Numba ABI for 'add' function

Source: https://nvidia.github.io/numba-cuda/user/cuda_compilation

Illustrates the PTX (Parallel Thread Execution) assembly code generated when compiling the 'add' function using Numba's internal ABI. It shows the mangled function name and the three parameters required (return value pointer, x, y).

```ptx
.visible .func  (.param .b32 func_retval0) _ZN8__main__3addB2v1B94cw51cXTLSUwv1sCUt9Uw1VEw0NRRQPKzLTg4gaGKFsG2oMQGEYakJSQB1PQBk0Bynm21OiwU1a0UoLGhDpQE8oxrNQE_3dEii(
    .param .b64 _ZN8__main__3addB2v1B94cw51cXTLSUwv1sCUt9Uw1VEw0NRRQPKzLTg4gaGKFsG2oMQGEYakJSQB1PQBk0Bynm21OiwU1a0UoLGhDpQE8oxrNQE_3dEii_param_0,
    .param .b32 _ZN8__main__3addB2v1B94cw51cXTLSUwv1sCUt9Uw1VEw0NRRQPKzLTg4gaGKFsG2oMQGEYakJSQB1PQBk0Bynm21OiwU1a0UoLGhDpQE8oxrNQE_3dEii_param_1,
    .param .b32 _ZN8__main__3addB2v1B94cw51cXTLSUwv1sCUt9Uw1VEw0NRRQPKzLTg4gaGKFsG2oMQGEYakJSQB1PQBk0Bynm21OiwU1a0UoLGhDpQE8oxrNQE_3dEii_param_2
 )


```

--------------------------------

### Install Numba CUDA with Conda (CUDA 13)

Source: https://nvidia.github.io/numba-cuda/user/installation

Installs the numba-cuda package and CUDA 13 toolkit into a Conda environment. This is the recommended approach for users preferring Conda for package management and CUDA toolkit integration.

```bash
$ conda install -c conda-forge numba-cuda "cuda-version=13"
```

--------------------------------

### Install Numba CUDA Dependencies with Pip (CUDA 13)

Source: https://nvidia.github.io/numba-cuda/user/installation

Installs Numba CUDA with support for CUDA 13 dependencies directly from PyPI using pip. This is an alternative to Conda for users who prefer pip for managing Python packages and their CUDA-related requirements.

```bash
$ pip install numba-cuda[cu13]
```

--------------------------------

### Launch CUDA Kernel with Manual Grid Configuration

Source: https://nvidia.github.io/numba-cuda/user/examples

Launches the Numba CUDA kernel `f` by manually specifying the grid and thread block dimensions using Numba's subscripting syntax (`[nblocks, nthreads]`). This provides more control over kernel launch configuration. Results are copied back and printed.

```python
# Enough threads per block for several warps per block
nthreads = 256
# Enough blocks to cover the entire vector depending on its length
nblocks = (len(a) // nthreads) + 1
f[nblocks, nthreads](a, b, c)
print(c.copy_to_host())
```

--------------------------------

### CUDA Reduction Kernel using Lambda Function (Python)

Source: https://nvidia.github.io/numba-cuda/user/reduction

This example shows an alternative way to define a CUDA reduction kernel using a lambda function with Numba's reduce decorator. It achieves the same functionality as the previous example with a more concise syntax.

```python
sum_reduce = cuda.reduce(lambda a, b: a + b)

```

--------------------------------

### Memory Pointer Classes

Source: https://nvidia.github.io/numba-cuda/user/external-memory

EMM Plugins should construct memory pointer instances that represent their allocations, for return to Numba. The appropriate memory pointer class to use depends on how the memory was allocated and whether it is mapped into the device memory space.

```APIDOC
## Memory Pointer Classes

EMM Plugins should construct memory pointer instances that represent their allocations, for return to Numba. The appropriate memory pointer class to use in each method is:

* `MemoryPointer`: returned from `memalloc`
* `MappedMemory`: returned from `memhostalloc` or `mempin` when the host memory is mapped into the device memory space.
* `PinnedMemory`: return from `memhostalloc` or `mempin` when the host memory is not mapped into the device memory space.

Memory pointers can take a finalizer, which is a function that is called when the buffer is no longer needed. Usually the finalizer will make a call to the memory management library (either internal to Numba, or external if allocated by an EMM Plugin) to inform it that the memory is no longer required, and that it could potentially be freed and/or unpinned. The memory manager may choose to defer actually cleaning up the memory to any later time after the finalizer runs - it is not required to free the buffer immediately.

### Class: `numba.cuda.MemoryPointer`

#### Description
A memory pointer that owns a buffer, with an optional finalizer. Memory pointers provide reference counting, and instances are initialized with a reference count of 1.

The base `MemoryPointer` class does not use the reference count for managing the buffer lifetime. Instead, the buffer lifetime is tied to the memory pointer instance’s lifetime: 
* When the instance is deleted, the finalizer will be called.
* When the reference count drops to 0, no action is taken.

Subclasses of `MemoryPointer` may modify these semantics, for example to tie the buffer lifetime to the reference count, so that the buffer is freed when there are no more references.

#### Parameters

* **context** (_Context_) – The context in which the pointer was allocated.
* **pointer** (_ctypes.c_void_p_) – The address of the buffer.
* **size** (_int_) – The size of the allocation in bytes.
* **owner** (_NoneType_) – The owner is sometimes set by the internals of this class, or used for Numba’s internal memory management. It should not be provided by an external user of the `MemoryPointer` class (e.g. from within an EMM Plugin); the default of None should always suffice.
* **finalizer** (_function_) – A function that is called when the buffer is to be freed.

### Class: `numba.cuda.cudadrv.driver.AutoFreePointer`

#### Description
Modifies the ownership semantic of the MemoryPointer so that the instance lifetime is directly tied to the number of references. When the reference count reaches zero, the finalizer is invoked. Constructor arguments are the same as for `MemoryPointer`.

### Class: `numba.cuda.MappedMemory`

#### Description
A memory pointer that refers to a buffer on the host that is mapped into device memory.

#### Parameters

* **context** (_Context_) – The context in which the pointer was mapped.
* **pointer** (_ctypes.c_void_p_) – The address of the buffer.
* **size** (_int_) – The size of the buffer in bytes.
* **owner** (_NoneType_) – The owner is sometimes set by the internals of this class, or used for Numba’s internal memory management. It should not be provided by an external user of the `MappedMemory` class (e.g. from within an EMM Plugin); the default of None should always suffice.
* **finalizer** (_function_) – A function that is called when the buffer is to be freed.

### Class: `numba.cuda.PinnedMemory`

#### Description
A pointer to a pinned buffer on the host.

#### Parameters

* **context** (_Context_) – The context in which the pointer was mapped.
* **pointer** (_ctypes.c_void_p_) – The address of the buffer.
* **size** (_int_) – The size of the buffer in bytes.
* **owner** – An object owning the buffer that has been pinned. For EMM plugin implementation, the default of `None` suffices for memory allocated in `memhostalloc` - for `mempin`, it should be the owner passed in to the `mempin` method.
* **finalizer** (_function_) – A function that is called when the buffer is to be freed.
```

--------------------------------

### GetIpcHandleMixin

Source: https://nvidia.github.io/numba-cuda/user/external-memory

Provides a default implementation for `get_ipc_handle()` using the driver API.

```APIDOC
## Class: GetIpcHandleMixin

### Description
A class that provides a default implementation of `get_ipc_handle()`.

### Methods

#### `get_ipc_handle(_memory_)`

*   **Description**: Opens an IPC memory handle by using `cuMemGetAddressRange` to determine the base pointer of the allocation. An IPC handle of type `cu_ipc_mem_handle` is constructed and initialized with `cuIpcGetMemHandle`. A `numba.cuda.IpcHandle` is returned, populated with the underlying `ipc_mem_handle`.
*   **Parameters**:
    *   `memory` - The memory object for which to get the IPC handle.
*   **Returns**: A `numba.cuda.IpcHandle` object.
```

--------------------------------

### Define CUDA Setup and Teardown Callbacks

Source: https://nvidia.github.io/numba-cuda/user/cuda_ffi

Defines placeholder functions for setup and teardown callbacks used with Numba CUDA modules. These functions are invoked during module loading and unloading, respectively, and receive a handle to the CUDA module.

```python
def setup_callback(mod: cuda.cudadrv.drvapi.cu_module):
    ...
def teardown_callback(mod: cuda.cudadrv.drvapi.cu_module):
    ...
```

--------------------------------

### CUDA Kernel Incrementing Data via Pointer

Source: https://nvidia.github.io/numba-cuda/user/examples

Defines and launches a CUDA kernel (`add_one`) that modifies data through a pointer. The kernel increments each element of the data pointed to by `x`, up to the specified length `n`. This example demonstrates obtaining a pointer from a Numba device array using the CUDA Array Interface.

```python
import numpy as np
from numba import cuda
from numba.cuda import types

sig = types.void(types.CPointer(types.uint8), types.uint32)

@cuda.jit(sig)
def add_one(x, n):
    i = cuda.grid(1)
    if i < n:
        x[i] += 1

x = cuda.to_device(np.arange(10, dtype=np.uint8))

# Print initial values of x
print(x.copy_to_host())  # [0 1 2 3 4 5 6 7 8 9]

# Obtain a pointer to the data from from the CUDA Array Interface
x_ptr = x.__cuda_array_interface__["data"][0]
x_len = len(x)

# Launch the kernel with the pointer and length
add_one[1, 32](x_ptr, x_len)

# Demonstrate that the data was updated by the kernel
print(x.copy_to_host())  # [ 1  2  3  4  5  6  7  8  9 10]
```

--------------------------------

### CUDA Non-Square Matrix Multiplication Example

Source: https://nvidia.github.io/numba-cuda/user/examples

Demonstrates how to perform non-square matrix multiplication on the GPU using Numba CUDA. It involves transferring data to the device, calculating grid and block dimensions, launching a kernel, and copying results back. This approach generalizes square matrix multiplication by adjusting block and grid sizes.

```python
import math
import numpy as np
from numba import cuda

# Assume TPB is defined elsewhere, e.g., TPB = 16

# Placeholder for the actual fast_matmul kernel definition
# @cuda.jit
def fast_matmul(x, y, z):
    # Kernel logic for matrix multiplication
    pass

x_h = np.arange(115).reshape([5, 23])
y_h = np.ones([23, 7])
z_h = np.zeros([5, 7])

x_d = cuda.to_device(x_h)
y_d = cuda.to_device(y_h)
z_d = cuda.to_device(z_h)

threadsperblock = (TPB, TPB)
grid_y_max = max(x_h.shape[0], y_h.shape[0])
grid_x_max = max(x_h.shape[1], y_h.shape[1])
blockspergrid_x = math.ceil(grid_x_max / threadsperblock[0])
blockspergrid_y = math.ceil(grid_y_max / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)

fast_matmul[blockspergrid, threadsperblock](x_d, y_d, z_d)
z_h = z_d.copy_to_host()
print(z_h)
print(x_h @ y_h)
```

--------------------------------

### Compile PTX using C ABI for 'add' function

Source: https://nvidia.github.io/numba-cuda/user/cuda_compilation

Compiles the 'add' Python function for CUDA device execution using the C ABI. This results in a predictable function name ('add') and standard parameter handling (x, y), simplifying integration with C/C++ code.

```python
ptx, resty = cuda.compile_ptx(add, int32(int32, int32), device=True, abi="c")

```

--------------------------------

### Import Numba CUDA Libraries

Source: https://nvidia.github.io/numba-cuda/user/examples

Imports necessary libraries from NumPy and Numba for CUDA operations, including data types.

```python
import numpy as np
from numba import cuda
from numba.cuda.types import int32
```

--------------------------------

### PTX prototype using C ABI for 'add' function

Source: https://nvidia.github.io/numba-cuda/user/cuda_compilation

Shows the PTX assembly code generated when compiling the 'add' function using the C ABI. The function name is preserved ('add'), and it takes exactly two parameters corresponding to the Python function's arguments.

```ptx
.visible .func  (.param .b32 func_retval0) add(
    .param .b32 add_param_0,
    .param .b32 add_param_1
)


```

--------------------------------

### Generate CUDA Device Data

Source: https://nvidia.github.io/numba-cuda/user/examples

Creates a one-dimensional NumPy array and transfers it to the CUDA device for processing. This data will be used as input for the CUDA kernel.

```python
# generate data
a = cuda.to_device(np.arange(1024))
nelem = len(a)
```

--------------------------------

### Compile PTX with custom C ABI name for 'add' function

Source: https://nvidia.github.io/numba-cuda/user/cuda_compilation

Compiles the 'add' function using the C ABI and specifies a custom ABI name ('add_f32') via the 'abi_info' dictionary. This allows distinguishing between different compiled versions of the same Python function, for example, with different data types.

```python
ptx, resty = cuda.compile_ptx(add, float32(float32, float32), device=True,
                              abi="c", abi_info={"abi_name": "add_f32"})


```

--------------------------------

### Import Numba CUDA and NumPy Modules

Source: https://nvidia.github.io/numba-cuda/user/examples

Imports necessary libraries for Numba CUDA programming and numerical operations. This includes `cuda` and `float32` from Numba, and `numpy` for array manipulation.

```python
from numba import cuda
from numba.cuda import float32
import numpy as np
import math
```

--------------------------------

### Compile PTX using Numba ABI for 'add' function

Source: https://nvidia.github.io/numba-cuda/user/cuda_compilation

Compiles the 'add' Python function for CUDA device execution using Numba's internal ABI. This compilation results in mangled function names and specific parameter handling, as shown in the generated PTX.

```python
ptx, resty = cuda.compile_ptx(add, int32(int32, int32), device=True)

```

--------------------------------

### Set Number of Samples for Monte Carlo Integration

Source: https://nvidia.github.io/numba-cuda/user/examples

Defines the number of random samples to be drawn for the Monte Carlo integration. A higher number of samples generally leads to a more accurate approximation of the integral.

```python
# number of samples, higher will lead to a more accurate answer
nsamps = 1000000
```

--------------------------------

### CUDA Kernel for Shared Memory Reduction

Source: https://nvidia.github.io/numba-cuda/user/examples

A Numba-decorated CUDA kernel that performs a reduction (summation) on input data using shared memory. It synchronizes threads and iteratively combines partial sums.

```python
@cuda.jit
def array_sum(data):
    tid = cuda.threadIdx.x
    size = len(data)
    if tid < size:
        i = cuda.grid(1)

        # Declare an array in shared memory
        shr = cuda.shared.array(nelem, int32)
        shr[tid] = data[i]

        # Ensure writes to shared memory are visible
        # to all threads before reducing
        cuda.syncthreads()

        s = 1
        while s < cuda.blockDim.x:
            if tid % (2 * s) == 0:
                # Stride by `s` and add
                shr[tid] += shr[tid + s]
            s *= 2
            cuda.syncthreads()

        # After the loop, the zeroth  element contains the sum
        if tid == 0:
            data[tid] = shr[tid]
```

--------------------------------

### Prepare Data for GPU Kernel with Numba CUDA

Source: https://nvidia.github.io/numba-cuda/user/examples

Prepares input data (X, Y, Z) and an output array (results) on the GPU using `numba.cuda.to_device`. This function is a prerequisite for launching a CUDA kernel that will utilize the `business_logic` function. It involves transferring data from host memory to device memory.

```python
from numba import cuda

X = cuda.to_device([1, 10, 234])
Y = cuda.to_device([2, 2, 4014])
Z = cuda.to_device([3, 14, 2211])
results = cuda.to_device([0.0, 0.0, 0.0])
```
Tokens:
Raw