Numba CUDA

https://nvidia.github.io/numba-cuda
Admin
Numba-CUDA provides a CUDA target for the Numba Python JIT Compiler, enabling the writing of SIMT...
Tokens:79,350
Snippets:937
Trust Score:8.3
Update:2 months ago
Show doc for...
Tokens:
Raw
### Load Pretty Printer Extension in launch.json

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Add a setup command to launch.json to automatically import the Numba CUDA pretty-printer extension when starting the debugger.

```json
{
    "setupCommands": [
        {
            "description": "Load the Numba CUDA cuda-gdb pretty-printer extension",
            "text": "python import gdb_print_extension",
            "ignoreFailures": true
        }
    ]
}
```

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

### Initialize Heat Equation Data

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

Setup the problem domain, initial temperature conditions, and device buffers.

```python
 1# Use an odd problem size.
 2# This is so there can be an element truly in the "middle" for symmetry.
 3size = 1001
 4data = np.zeros(size)
 5
 6# Middle element is made very hot
 7data[500] = 10000
 8buf_0 = cuda.to_device(data)
 9
10# This extra array is used for synchronization purposes
11buf_1 = cuda.device_array_like(buf_0)
12
13niter = 10000
```

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

### Importing Numba and Consumer Library

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

Initial setup required to demonstrate cross-library array consumption.

```python
from numba import cuda, int32, void
import other_cai_library
```

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

### Clone Numba CUDA Repository

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Initial step to access debugging examples and pretty printer extensions.

```bash
git clone https://github.com/NVIDIA/numba-cuda.git
```

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

### Define Setup and Teardown Callbacks

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

Callbacks for module-specific initialization and cleanup tasks, invoked by the CUDA driver module handle.

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

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

### Configuring and Launching a Cooperative Kernel

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

Setup input data, define dimensions, and launch the kernel. The launch is implicitly cooperative when a grid group is synchronized.

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

```python
 1# Kernel launch - this is implicitly a cooperative launch
 2sequential_rows[griddim, blockdim](A)
 3
 4# What do the results look like?
 5# print(A)
 6#
 7# [[   0    0    0 ...    0    0    0]
 8#  [   1    1    1 ...    1    1    1]
 9#  [   2    2    2 ...    2    2    2]
10#  ...
11#  [1021 1021 1021 ... 1021 1021 1021]
12#  [1022 1022 1022 ... 1022 1022 1022]
13#  [1023 1023 1023 ... 1023 1023 1023]]
```

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

### Implementing a Cooperative Kernel

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

A kernel example demonstrating grid synchronization to ensure visibility of writes across all threads in the grid.

```python
 1from numba import cuda
 2from numba.cuda import int32
 3import numpy as np
 4
 5sig = (int32[:, ::1],)
 6
 7@cuda.jit(sig)
 8def sequential_rows(M):
 9    col = cuda.grid(1)
10    g = cuda.cg.this_grid()
11
12    rows = M.shape[0]
13    cols = M.shape[1]
14
15    for row in range(1, rows):
16        opposite = cols - col - 1
17        # Each row's elements are one greater than the previous row
18        M[row, col] = M[row - 1, opposite] + 1
19        # Wait until all threads have written their column element,
20        # and that the write is visible to all other threads
21        g.sync()
```

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

### Install Numba CUDA with pip

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

Use these commands to install Numba CUDA and its dependencies from PyPI for specific CUDA versions.

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

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

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

### Naive Matrix Multiplication Usage

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

Example usage of the naive matrix multiplication kernel with host-to-device data transfer.

```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)
```

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

### Basic CUDA Vectorize Example

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/ufunc.html

Demonstrates defining a vectorized function for CUDA using the @vectorize decorator with specified type signatures.

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

@vectorize(['float32(float32, float32, float32)',
            'float64(float64, float64, float64)'],
           target='cuda')
def cu_discriminant(a, b, c):
    return math.sqrt(b ** 2 - 4 * a * c)

N = 10000
dtype = np.float32

# prepare the input
A = np.array(np.random.sample(N), dtype=dtype)
B = np.array(np.random.sample(N) + 10, dtype=dtype)
C = np.array(np.random.sample(N), dtype=dtype)

D = cu_discriminant(A, B, C)

print(D)  # print result
```

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

### Usage of Naive Matrix Multiplication

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

Example of how to invoke the naive matrix multiplication kernel on device data.

```python
 1x_h = np.arange(16).reshape([4, 4])
 2y_h = np.ones([4, 4])
 3z_h = np.zeros([4, 4])
 4
 5x_d = cuda.to_device(x_h)
 6y_d = cuda.to_device(y_h)
 7z_d = cuda.to_device(z_h)
 8
 9threadsperblock = (16, 16)
10blockspergrid_x = math.ceil(z_h.shape[0] / threadsperblock[0])
11blockspergrid_y = math.ceil(z_h.shape[1] / threadsperblock[1])
12blockspergrid = (blockspergrid_x, blockspergrid_y)
13
14matmul[blockspergrid, threadsperblock](x_d, y_d, z_d)
15z_h = z_d.copy_to_host()
16print(z_h)
17print(x_h @ y_h)

```

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

### Generalized CUDA Ufunc Example

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/ufunc.html

Illustrates the use of @guvectorize to perform generalized operations like matrix multiplication on the GPU.

```python
from numba import guvectorize

@guvectorize(['void(float32[:,:], float32[:,:], float32[:,:])'],
             '(m,n),(n,p)->(m,p)', target='cuda')
def matmulcore(A, B, C):
    ...
```

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

### Install Numba CUDA MLIR with CUDA 13

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

Installs the package with dependencies for CUDA Toolkit 13.

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

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

### Install Numba CUDA MLIR with CUDA 12

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

Installs the package with dependencies for CUDA Toolkit 12.

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

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

### Registering a Global Typing Template

Source: https://nvidia.github.io/numba-cuda-mlir/latest/extending/low-level.html

Example of implementing an AbstractTemplate to type a global operator, registered via the typing_registry.

```python
from numba_cuda_mlir.extending import typing_registry
from numba_cuda_mlir.numba_cuda.typing.templates import (
    AbstractTemplate, signature,
)
from numba_cuda_mlir import types
import operator

@typing_registry.register_global(operator.setitem)
class Uint64PointerSetitemTemplate(AbstractTemplate):
    def generic(self, args, kws):
        if len(args) != 3:
            return None
        array, idx, value = args
        if (isinstance(array, types.Array)
                and array.dtype == types.uint64
                and isinstance(idx, types.Integer)
                and value is my_pointer_type):
            return signature(types.none, array, idx, value)
        return None
```

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

### Stacking lowering decorators

Source: https://nvidia.github.io/numba-cuda-mlir/latest/extending/low-level.html

Example of registering a single lowering implementation against multiple function signatures.

```python
@lower(operator.not_, types.Number)
@lower(operator.not_, types.Boolean)
def lower_not(builder, target, args, kwargs):
    ...
```

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

### Matrix Multiplication Output

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

Expected console output after running the optimized matrix multiplication example.

```bash
$ python fast_matmul.py
[[ 6.  6.  6.  6.]
[22. 22. 22. 22.]
[38. 38. 38. 38.]
[54. 54. 54. 54.]]
[[ 6.  6.  6.  6.]
[22. 22. 22. 22.]
[38. 38. 38. 38.]
[54. 54. 54. 54.]]
```

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

### Install Numba CUDA with Conda

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

Use these commands to install Numba CUDA and a specific CUDA toolkit version into a Conda environment.

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

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

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

### Disjoint Views Output

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/memory.html

The output resulting from the disjoint views example.

```text
3.140000
1
```

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

### Importing Numba CUDA Modules

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

Required imports for Numba CUDA matrix multiplication examples.

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

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

### Demonstrate Aliasing in Dynamic Shared Memory

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/memory.html

Example showing how multiple dynamic shared memory arrays alias the same memory region if not partitioned correctly.

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

@cuda.jit
def f():
   f32_arr = cuda.shared.array(0, dtype=np.float32)
   i32_arr = cuda.shared.array(0, dtype=np.int32)
   f32_arr[0] = 3.14
   print(f32_arr[0])
   print(i32_arr[0])

f[1, 1, 0, 4]()
cuda.synchronize()
```

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

### Define a C++ device function prototype

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

Example C++ prototype for a device function intended to be called from a Numba kernel.

```cpp
extern "C"
__device__ int
sum_reduce(
  float* return_value,
  float* array,
  int n
);
```

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

### Use Disjoint Views for Dynamic Shared Memory

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/memory.html

Example showing how to use disjoint views to prevent aliasing when using multiple dynamic shared memory arrays.

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

@cuda.jit
def f_with_view():
   f32_arr = cuda.shared.array(0, dtype=np.float32)
   i32_arr = cuda.shared.array(0, dtype=np.int32)[1:] # 1 int32 = 4 bytes
   f32_arr[0] = 3.14
   i32_arr[0] = 1
   print(f32_arr[0])
   print(i32_arr[0])

f_with_view[1, 1, 0, 8]()
cuda.synchronize()
```

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

### Start a cuda-gdb debugging session

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Initialize a debugging session for a Python script using cuda-gdb and set a breakpoint on kernel launch.

```bash
$ cuda-gdb python3
NVIDIA (R) cuda-gdb 13.1
Portions Copyright (C) 2007-2025 NVIDIA Corporation
Based on GNU gdb 14.2
Copyright (C) 2023 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.
Type "show copying" and "show warranty" for details.
This CUDA-GDB was configured as "x86_64-pc-linux-gnu".
Type "show configuration" for configuration details.
For bug reporting instructions, please see:
<https://forums.developer.nvidia.com/c/developer-tools/cuda-developer-tools/cuda-gdb>.
Find the CUDA-GDB manual and other documentation resources online at:
    <https://docs.nvidia.com/cuda/cuda-gdb/index.html>.

For help, type "help".
Type "apropos word" to search for commands related to "word"...
Reading symbols from python3...
(cuda-gdb) set cuda break_on_launch application
(cuda-gdb) run hello.py
Starting program: /home/mmason/anaconda3/envs/Numba-GA-python3.12/bin/python3 hello.py
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[Detaching after fork from child process 1994689]
[Detaching after fork from child process 1994698]
input: [0 1 2 3 4 5 6 7]
/home/mmason/PythonDebugging/NUMBA/numba-cuda-numba-debug/numba_cuda/numba/cuda/dispatcher.py:690: NumbaPerformanceWarning: Grid size 1 will likely result in GPU under-utilization due to low occupancy.
  warn(errors.NumbaPerformanceWarning(msg))
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]

CUDA thread hit application kernel entry function breakpoint, hello<<<(1,1,1),(8,1,1)>>> (input=..., output=...) at hello.py:18
18     gid = cuda.grid(1)
(cuda-gdb) list
13 # Editor window, and are shown as red dots next to the line numbers. They can also be set
14 # from the CUDA GDB command line using the filename:line-number syntax, or by kernel or function name.
15 #
16 @cuda.jit(debug=True, opt=False)
17 def hello(input, output):
18     gid = cuda.grid(1)
19     size = len(input)
20     if gid >= size:
21         return
22
(cuda-gdb) break 20
Breakpoint 1 at 0x7ffd436f83c0: file hello.py, line 20.
(cuda-gdb) continue
Continuing.

CUDA thread hit Breakpoint 1, hello<<<(1,1,1),(8,1,1)>>> (input=..., output=...) at hello.py:20
20     if gid >= size:
(cuda-gdb) print gid
$1 = 0
(cuda-gdb) p size
$2 = 8
```

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

### Start a CUDA GDB Debugging Session

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Commands to launch the Python interpreter within CUDA GDB and begin execution of a Numba CUDA script.

```bash
$ cuda-gdb python3
(cuda-gdb) set cuda break_on_launch application
(cuda-gdb) run /path/to/numba-cuda/examples/debugging/hello.py
```

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

### Configure and Launch CUDA Kernels

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Demonstrates the syntax for configuring a kernel launch using subscripting and calling it with arguments.

```python
# func is some function decorated with @cuda.jit
func[griddim, blockdim, stream, sharedmem]
```

```python
configured = func[griddim, blockdim, stream, sharedmem]
configured(x, y, z)
```

```python
func[griddim, blockdim, stream, sharedmem](x, y, z)
```

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

### Launch kernel using forall

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

Launch the kernel using the forall() method to automatically configure the 1D grid.

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

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

### Launch kernel with manual grid configuration

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

Configure the grid manually using subscripting syntax to specify blocks and threads.

```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())
```

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

### Dynamic Shared Memory Output

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/memory.html

The output resulting from the aliasing example.

```text
3.140000
1078523331
```

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

### Register pre-launch callbacks

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Attach callbacks to a configured launch to execute logic immediately before the kernel is launched.

```python
cfg = f[1, 1]

def log_launch(kernel, cfg):
    print(cfg.griddim, cfg.blockdim)

cfg.pre_launch_callbacks.append(log_launch)
cfg(arr)
```

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

### Optimized Matrix Multiplication Usage

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

Example usage of the shared-memory optimized matrix multiplication kernel.

```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)
```

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

### Kernel Launch Configuration Syntax

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/memory.html

The standard configuration parameters for launching a CUDA kernel.

```python
kernel_func[grid_dim, block_dim, stream, dyn_shared_mem_size]
```

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

### Load and use the pretty printer extension

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Import the pretty printer extension to display array variables in a human-readable format.

```bash
(cuda-gdb) python import gdb_print_extension
(cuda-gdb) print shared_array
$4 = [-7 -6 -5 -4 -3 -2 -1  0]
```

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

### Execute Heat Equation Kernel

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

Launch the kernel using the forall method to process the simulation over the defined number of iterations.

```python
1solve_heat_equation.forall(len(data))(buf_0, buf_1, niter, 0.25)
```

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

### Execute and verify kernel

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

Run the kernel and compare the result with host-side summation.

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

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

### Launch VSCode Workspace

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Open the provided workspace file to load the debugging environment.

```bash
cd numba-cuda/examples/debugging
code debugging.code-workspace
```

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

### IpcArrayHandle.open()

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

Creates a new DeviceNDArray object that shares the allocation from the original process.

```APIDOC
## IpcArrayHandle.open()

### Description
Returns a new DeviceNDArray that shares the allocation from the original process. This method must not be used on the original process.
```

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

### Registering a function lowering

Source: https://nvidia.github.io/numba-cuda-mlir/latest/extending/low-level.html

Example of using the lowering_registry to register a custom lowering implementation for a specific function and argument type.

```python
from numba_cuda_mlir.extending import lowering_registry
from numba_cuda_mlir.lowering_utilities import convert
from numba_cuda_mlir._mlir import ir as mlir_ir
from numba_cuda_mlir._mlir.dialects import llvm

@lowering_registry.lower(make_boxed_int, types.Integer)
def _lower_make_boxed_int(builder, target, args, kwargs):
    val = builder.load_var(args[0])
    struct_ty = builder.get_mlir_type(my_boxed_int)
    i64_ty = mlir_ir.IntegerType.get_signless(64)
    val = convert(val, i64_ty)
    undef = llvm.UndefOp(struct_ty)
    result = llvm.insertvalue(
        container=undef,
        value=val,
        position=mlir_ir.DenseI64ArrayAttr.get([0]),
    )
    builder.store_var(target, result)
```

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

### Compare Numba and CUDA C/C++ launch syntax

Source: https://nvidia.github.io/numba-cuda-mlir/latest/reference/kernel.html

Note that the order of stream and sharedmem parameters is reversed compared to standard CUDA C/C++ syntax.

```cpp
func<<<griddim, blockdim, sharedmem, stream>>>(x, y, z)
```

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

### forall(ntasks, tpb=0, stream=0, sharedmem=0)

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Returns a 1D-configured dispatcher for a given number of tasks, assuming a 1-1 mapping between Global Thread ID and tasks.

```APIDOC
## forall(ntasks, tpb=0, stream=0, sharedmem=0)

### Description
Returns a 1D-configured dispatcher for a given number of tasks. This assumes the kernel maps the Global Thread ID to tasks on a 1-1 basis.

### Parameters
- **ntasks** (int) - Required - The number of tasks.
- **tpb** (int) - Optional - The size of a block.
- **stream** (cuda.stream) - Optional - The stream on which the dispatcher will be launched.
- **sharedmem** (int) - Optional - The number of bytes of dynamic shared memory required.
```

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

### numba.cuda.launchconfig.current_launch_config()

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Retrieves the active launch configuration if available, otherwise returns None.

```APIDOC
## numba.cuda.launchconfig.current_launch_config()

### Description
Returns the current launch configuration object if one exists, or None if called outside of a launch-triggered compilation context.
```

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

### Configure a CUDA kernel dispatcher

Source: https://nvidia.github.io/numba-cuda-mlir/latest/reference/kernel.html

Use subscripting on a @cuda.jit decorated function to define the launch configuration.

```python
# func is some function decorated with @cuda.jit
func[griddim, blockdim, stream, sharedmem]
```

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

### Copying Device Data to Host

Source: https://nvidia.github.io/numba-cuda/reference/memory.html

Demonstrates transferring data from a GPU device array back to a host NumPy array after kernel execution.

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

arr = np.arange(1000)
d_arr = cuda.to_device(arr)

my_kernel[100, 100](d_arr)

result_array = d_arr.copy_to_host()
```

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

### Inspect Variables with Pretty Printer

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

Demonstrates the difference between raw memory inspection and using the pretty printer extension for array variables.

```bash
Thread 1 "python3" received signal SIGTRAP, Trace/breakpoint trap.
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
hello<<<(1,1,1),(8,1,1)>>> (input=..., output=...) at hello.py:50
50    shared_array[gid] = - input[size - gid - 1]
(cuda-gdb) next
53    cuda.syncthreads()
(cuda-gdb) print shared_array
$3 = {meminfo = 0x0, parent = 0x0, nitems = 8, itemsize = 8, data = 0x7ffe00000400, shape = {8}, strides = {8}}
```

```bash
(cuda-gdb) python import gdb_print_extension
(cuda-gdb) print shared_array
$4 = [-7 -6 -5 -4 -3 -2 -1  0]
```

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

### numba.cuda.Reduce.__init__

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

Initializes a reduction object using a binary function that is compiled as a CUDA device function.

```APIDOC
## numba.cuda.Reduce(functor)

### Description
Creates a reduction object that reduces values using a given binary function. The binary function is compiled once and cached inside this object.

### Parameters
- **functor** (callable) - Required - A function implementing a binary operation for reduction. It will be compiled as a CUDA device function using `cuda.jit(device=True)`.
```

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

### Create Numba CUDA Debugging Environment

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Commands to initialize a new Anaconda virtual environment for Numba CUDA development.

```bash
conda create --name numba-cuda-debug python=3.12
conda activate numba-cuda-debug
```

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

### Configure PYTHONPATH for Pretty Printer

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/debugging.html

Set the PYTHONPATH in launch.json to ensure the GDB pretty printer extension is accessible to the debugger.

```json
{
    "environment": [
        {
            "name": "PYTHONPATH",
            "value": "${workspaceFolder}/../../misc:${env:PYTHONPATH}"
        }
    ]
}
```

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

### Calling Device Functions from CUDA Ufuncs

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/ufunc.html

Shows how to define a device function using @cuda.jit and call it from within a @vectorize kernel.

```python
from numba import vectorize, cuda

# define a device function
@cuda.jit('float32(float32, float32, float32)', device=True, inline="always")
def cu_device_fn(x, y, z):
    return x ** y / z

# define a ufunc that calls our device function
@vectorize(['float32(float32, float32, float32)'], target='cuda')
def cu_ufunc(x, y, z):
    return cu_device_fn(x, y, z)
```

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

### inspect_asm(signature=None)

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Return this kernel’s PTX assembly code for the device in the current context.

```APIDOC
## inspect_asm(signature=None)

### Description
Return this kernel’s PTX assembly code for the device in the current context.

### Parameters
- **signature** (tuple) - Optional - A tuple of argument types.

### Returns
- The PTX code for the given signature, or a dict of PTX codes for all previously-encountered signatures.
```

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

### Execute Integration

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

Define the target function and compute the integral between specified limits.

```python
1# define a function to integrate
2@numba.jit
3def func(x):
4    return 1.0 / x
5
6mc_integrate(1, 2, nsamps)  # array(0.6929643, dtype=float32)
7mc_integrate(2, 3, nsamps)  # array(0.4054021, dtype=float32)
```

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

### Dispatcher Launch Configuration

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Configures a CUDA kernel launch with grid dimensions, block dimensions, stream, and shared memory size.

```APIDOC
## Dispatcher Launch Configuration

### Description
Configures a kernel launch using subscripting syntax. The kernel is launched with the specified grid dimensions, block dimensions, stream, and dynamic shared memory size.

### Syntax
`func[griddim, blockdim, stream, sharedmem](*args)`

### Parameters
- **griddim** (int or tuple) - Required - Size of the grid.
- **blockdim** (int or tuple) - Required - Size of the thread blocks.
- **stream** (cuda.stream) - Optional - The CUDA stream for the launch.
- **sharedmem** (int) - Optional - Size of dynamic shared memory in bytes.
```

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

### Define JIT-compiled business logic

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

Define a function using @numba.jit to enable compilation for CPU execution.

```python
1@numba.jit
2def business_logic(x, y, z):
3    return 4 * z * (2 * x - (4 * y) / 2 * pi)
4
```

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

### Use IpcArrayHandle as a context manager

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

Demonstrates using the context manager interface to automatically open and close an IPC array handle.

```python
with the_ipc_array_handle as ipc_array:
    # use ipc_array here as a normal gpu array object
    some_code(ipc_array)
# ipc_array is dead at this point
```

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

### Prepare GPU device arrays

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

Transfer data to the GPU device using cuda.to_device.

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

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

### numba.cuda.launchconfig.capture_compile_config(dispatcher)

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

A context manager for observing the compile-time launch configuration for a specific dispatcher.

```APIDOC
## numba.cuda.launchconfig.capture_compile_config(dispatcher)

### Description
Captures the launch configuration used during the compilation of a kernel. Useful for testing, debugging, or external introspection.

### Usage
```python
with launchconfig.capture_compile_config(f) as capture:
    f[1, 1](arr)
cfg = capture["config"]
```
```

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

### Legacy Link Support in @cuda.jit

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

Demonstrates the deprecated method of passing linkable files directly to the @cuda.jit decorator.

```python
@cuda.jit(link=['functions.cu'])
def multiply_vectors(r, x, y):
    i = cuda.grid(1)

    if i < len(r):
        r[i] = mul(x[i], y[i])
```

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

### Constructing Vectors from Components

Source: https://nvidia.github.io/numba-cuda-mlir/latest/reference/types.html

Shows how to build larger vector types by combining primitive types and smaller vector types.

```python
zero = uint32(0)
u2 = uint32x2(1, 2)
# Construct a 3-component vector with primitive type and a 2-component vector
u3 = uint32x3(zero, u2)
# Construct a 4-component vector with 2 2-component vectors
u4 = uint32x4(u2, u2)
```

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

### Invoke CUDA kernel

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

Launch the kernel using the forall method and print the results.

```python
1f.forall(len(X))(results, X, Y, Z)
2print(results)
3# [-126.79644737231007, 416.28324559588634, -218912930.2987788]
```

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

### specialize(*args)

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Create a new instance of this dispatcher specialized for the given args.

```APIDOC
## specialize(*args)

### Description
Create a new instance of this dispatcher specialized for the given args.

### Parameters
- **args** (tuple) - Required - The argument types to specialize for.
```

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

### inspect_llvm(signature=None)

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Return the LLVM IR for this kernel.

```APIDOC
## inspect_llvm(signature=None)

### Description
Return the LLVM IR for this kernel.

### Parameters
- **signature** (tuple) - Optional - A tuple of argument types.

### Returns
- The LLVM IR for the given signature, or a dict of LLVM IR for all previously-encountered signatures.
```

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

### Launch a configured CUDA kernel

Source: https://nvidia.github.io/numba-cuda-mlir/latest/reference/kernel.html

Kernel configuration can be stored as an object or executed in a single statement.

```python
configured = func[griddim, blockdim, stream, sharedmem]
configured(x, y, z)
```

```python
func[griddim, blockdim, stream, sharedmem](x, y, z)
```

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

### Prepare data for CUDA kernel

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

Transfer host data to the device for processing.

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

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

### Execute JIT function on CPU

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

Run the compiled function standalone on the CPU.

```python
1print(business_logic(1, 2, 3))  # -126.79644737231007
```

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

### Configure Kernel Execution Grid

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/kernels.html

Manually calculate the grid dimensions based on array shape and threads per block before launching the kernel.

```python
threadsperblock = (16, 16)
blockspergrid_x = math.ceil(an_array.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(an_array.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)
increment_a_2D_array[blockspergrid, threadsperblock](an_array)
```

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

### Define a reduction using a lambda

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

Shows how to create a reduction kernel using a lambda function instead of a decorated function.

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

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

### Demonstrating constant capture of global variables

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/globals.html

Shows that host-side global variables are captured as constants during the first kernel compilation, rendering subsequent modifications ineffective.

```python
 1import numpy as np
 2from numba import cuda
 3
 4TAX_RATE = 0.08
 5PRICES = np.array([10.0, 25.0, 5.0, 15.0, 30.0], dtype=np.float64)
 6
 7@cuda.jit
 8def compute_totals(quantities, totals):
 9    i = cuda.grid(1)
10    if i < totals.size:
11        totals[i] = quantities[i] * PRICES[i] * (1 + TAX_RATE)
12
13d_quantities = cuda.to_device(np.array([1, 2, 3, 4, 5], dtype=np.float64))
14d_totals = cuda.device_array(5, dtype=np.float64)
15
16# First kernel call - compiles and captures values
17compute_totals[1, 32](d_quantities, d_totals)
18print("Value of d_totals:", d_totals.copy_to_host())
19
20# These modifications have no effect on subsequent kernel calls
21TAX_RATE = 0.10  # noqa: F841
22PRICES[:] = [20.0, 50.0, 10.0, 30.0, 60.0]
23
24# Second kernel call still uses the original values
25compute_totals[1, 32](d_quantities, d_totals)
26print("Value of d_totals:", d_totals.copy_to_host())
```

```text
Value of d_totals: [ 10.8  54.   16.2  64.8 162. ]
Value of d_totals: [ 10.8  54.   16.2  64.8 162. ]
```

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

### Define and use a reduction kernel

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

Demonstrates defining a reduction function using the @reduce decorator and applying it to a NumPy array.

```python
import numpy
from numba_cuda_mlir import cuda

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

A = (numpy.arange(1234, dtype=numpy.float64)) + 1
expect = A.sum()      # NumPy sum reduction
got = sum_reduce(A)   # cuda sum reduction
assert expect == got
```

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

### Import required modules

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

Initial imports for NumPy and Numba CUDA types.

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

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

### Generate Test Data for Sessionization

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

Prepares input arrays and result buffers on the device for testing the sessionization kernel.

```python
 1# Generate data
 2ids = cuda.to_device(
 3    np.array(
 4        [
 5            1,
 6            1,
 7            1,
 8            1,
 9            1,
10            1,
11            2,
12            2,
13            2,
14            3,
15            3,
16            3,
17            3,
18            3,
19            3,
20            3,
21            3,
22            3,
23            3,
24            4,
25            4,
26            4,
27            4,
28            4,
29            4,
30            4,
31            4,
32            4,
33        ]
34    )
35)
36sec = cuda.to_device(
37    np.array(
38        [
39            1,
40            2,
41            3,
42            5000,
43            5001,
44            5002,
45            1,
46            2,
47            3,
48            1,
49            2,
50            5000,
51            5001,
52            10000,
53            10001,
54            10002,
55            10003,
56            15000,
57            150001,
58            1,
59            5000,
60            50001,
61            15000,
62            20000,
63            25000,
64            25001,
65            25002,
66            25003,
67        ],
68        dtype="datetime64[ns]",
69    ).astype("int64")  # Cast to int64 for compatibility
70)
71# Create a vector to hold the results
72results = cuda.to_device(np.zeros(len(ids)))
```

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

### inspect_sass(signature=None)

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Return this kernel’s SASS assembly code for the device in the current context.

```APIDOC
## inspect_sass(signature=None)

### Description
Return this kernel’s SASS assembly code for the device in the current context. Requires nvdisasm to be available on the PATH.

### Parameters
- **signature** (tuple) - Optional - A tuple of argument types.

### Returns
- The SASS code for the given signature, or a dict of SASS codes for all previously-encountered signatures.
```

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

### numba.cuda.detect()

Source: https://nvidia.github.io/numba-cuda/reference/host.html

Detects supported CUDA hardware and prints a summary to the console.

```APIDOC
## numba.cuda.detect()

### Description
Detect supported CUDA hardware and print a summary of the detected hardware.

### Returns
- **bool** - True if any supported devices were detected, False otherwise.
```

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

### Constructing a Vector Type

Source: https://nvidia.github.io/numba-cuda-mlir/latest/reference/types.html

Demonstrates direct construction of a float32x3 vector type using its constructor.

```python
from numba.cuda import float32x3

# In kernel
f3 = float32x3(0.0, -1.0, 1.0)
```

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

### numba.cuda.is_available()

Source: https://nvidia.github.io/numba-cuda/reference/host.html

Checks if a CUDA GPU is available for use. This function will initialize the driver if it has not been initialized yet.

```APIDOC
## numba.cuda.is_available()

### Description
Returns a boolean to indicate the availability of a CUDA GPU.

### Returns
- **bool** - True if a CUDA GPU is available, False otherwise.
```

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

### Kernel Dispatcher Configuration and Launch

Source: https://nvidia.github.io/numba-cuda-mlir/latest/reference/kernel.html

Configures a CUDA kernel dispatcher with grid, block, stream, and shared memory parameters and executes it with kernel arguments.

```APIDOC
## Kernel Launch Configuration

### Description
Configures and launches a CUDA kernel using the dispatcher object created by the @cuda.jit decorator.

### Syntax
`func[griddim, blockdim, stream, sharedmem](arg1, arg2, ...)`

### Parameters
- **griddim** (int or tuple) - Required - The size of the grid (up to 3 dimensions).
- **blockdim** (int or tuple) - Required - The size of the thread blocks (up to 3 dimensions).
- **stream** (optional) - Optional - The CUDA stream on which the kernel will be launched.
- **sharedmem** (int) - Optional - The size of dynamic shared memory in bytes.

### Usage Example
```python
# Configure and launch in a single statement
func[griddim, blockdim, stream, sharedmem](x, y, z)
```
```

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

### Launching Kernel on a Non-Default Stream

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

Executing the kernel in a stream different from the array's default stream.

```python
nthreads = 256
nblocks = N // nthreads

initialize_array[nthreads, nblocks, kernel_stream](x)
```

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

### Copy data from device to host

Source: https://nvidia.github.io/numba-cuda/reference/memory.html

Copies data from a DeviceNDArray back to the host, either creating a new array or using an existing one.

```python
hary = d_ary.copy_to_host()
```

```python
ary = np.empty(shape=d_ary.shape, dtype=d_ary.dtype)
d_ary.copy_to_host(ary)
```

```python
hary = d_ary.copy_to_host(stream=stream)
```

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

### C ABI return value handling

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

PTX instructions showing how the result is directly placed in the return value when using the C ABI.

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

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

### Switch CUDA device

Source: https://nvidia.github.io/numba-cuda-mlir/latest/user/device-management.html

Selects a different device to create a new context, assuming multiple GPUs are available.

```python
cuda.select_device(1)  # assuming we have 2 GPUs
```

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

### bfloat16 Min/Max and Predicates

Source: https://nvidia.github.io/numba-cuda-mlir/latest/reference/types.html

Functions for finding min/max values and checking for special floating-point states.

```APIDOC
## numba_cuda_mlir.cuda.bf16.hmax(a, b)

### Description
Returns max(a, b). If either input is NaN, the other input is returned; if both are NaN, the canonical NaN is returned.

## numba_cuda_mlir.cuda.bf16.hisnan(a)

### Description
Returns True if a is a NaN, False otherwise.

## numba_cuda_mlir.cuda.bf16.hisinf(a)

### Description
Returns a nonzero integer if a is infinite, otherwise 0.
```

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

### mark_launch_config_sensitive()

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Mark the active compilation for this dispatcher as LC-S.

```APIDOC
## mark_launch_config_sensitive()

### Description
Mark the active compilation for this dispatcher as LC-S. This method must be called while a launch-triggered compilation is in progress.
```

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

### Generate input data

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

Prepare one-dimensional data on the device for the reduction kernel.

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

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

### numba.cuda.jit

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

JIT compile a Python function for CUDA GPUs. This decorator creates a CUDA dispatcher object that can be configured and launched.

```APIDOC
## numba.cuda.jit

### Description
JIT compile a Python function for CUDA GPUs. The decorator returns a Dispatcher object if a function is provided, or a function that accepts another function to compile.

### Signature
`numba.cuda.jit(func_or_sig=None, device=False, inline='never', forceinline=False, link=(), debug=None, opt=None, lineinfo=False, cache=False, launch_bounds=None, lto=None, shared_memory_carveout=None, **kws)`

### Parameters
- **func_or_sig** (function or signature) - Optional - A function to JIT compile, or signatures of a function to compile.
- **device** (bool) - Optional - Indicates whether this is a device function.
- **inline** (str) - Optional - Enables inlining at the Numba IR level when set to 'always'.
- **forceinline** (bool) - Optional - Enables inlining at the NVVM IR level.
- **link** (list) - Optional - A list of files containing PTX or CUDA C/C++ source to link.
- **debug** (bool) - Optional - If True, check for exceptions thrown when executing the kernel.
- **fastmath** (bool) - Optional - Enables fastmath optimizations.
- **max_registers** (int) - Optional - Request that the kernel is limited to using at most this number of registers per thread.
- **opt** (bool) - Optional - Whether to compile with optimization enabled.
- **lineinfo** (bool) - Optional - If True, generate a line mapping between source code and assembly code.
- **cache** (bool) - Optional - If True, enables the file-based cache for this function.
- **launch_bounds** (int or tuple) - Optional - Kernel launch bounds (max threads per block, min blocks per SM, max blocks per cluster).
- **lto** (bool) - Optional - Whether to enable LTO.
- **shared_memory_carveout** (str or int) - Optional - Controls the partitioning of shared memory and L1 cache.
```

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

### Access compile-time launch configuration

Source: https://nvidia.github.io/numba-cuda/reference/kernel.html

Use ensure_current_launch_config to retrieve the active launch configuration during compilation.

```python
from numba.cuda import launchconfig

cfg = launchconfig.ensure_current_launch_config()
print(cfg.griddim, cfg.blockdim, cfg.sharedmem, cfg.args)
```
Tokens:
Raw