Global Variables and Captured Values

Numba CUDA kernels and device functions can reference global variables defined at module scope. This section describes how these values are captured and the implications for your code.

Capture as constants

By default, global variables referenced in kernels are captured as constants at compilation time. This applies to scalars and host arrays (e.g. NumPy arrays).

The following example demonstrates this behavior. Both TAX_RATE and PRICES are captured when the kernel is first compiled. Because they are embedded as constants, modifications to these variables after compilation have no effect—the second kernel call still uses the original values:

Demonstrating constant capture of global variables

import numpy as np
from numba import cuda

TAX_RATE = 0.08
PRICES = np.array([10.0, 25.0, 5.0, 15.0, 30.0], dtype=np.float64)

@cuda.jit
def compute_totals(quantities, totals):
    i = cuda.grid(1)
    if i < totals.size:
        totals[i] = quantities[i] * PRICES[i] * (1 + TAX_RATE)

d_quantities = cuda.to_device(
    np.array([1, 2, 3, 4, 5], dtype=np.float64)
)
d_totals = cuda.device_array(5, dtype=np.float64)

# First kernel call - compiles and captures values
compute_totals[1, 32](d_quantities, d_totals)
print("Value of d_totals:", d_totals.copy_to_host())

# These modifications have no effect on subsequent kernel calls
TAX_RATE = 0.10  # noqa: F841
PRICES[:] = [20.0, 50.0, 10.0, 30.0, 60.0]

# Second kernel call still uses the original values
compute_totals[1, 32](d_quantities, d_totals)
print("Value of d_totals:", d_totals.copy_to_host())

Running the above code prints:

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

Note that both outputs are identical—the modifications to TAX_RATE and PRICES after the first kernel call have no effect.

This behaviour is useful for small amounts of truly constant data like configuration values, lookup tables, or mathematical constants. For larger arrays, consider using device arrays instead.

Device array capture

Device arrays are an exception to the constant capture rule. When a kernel references a global device array—any object implementing __cuda_array_interface__, such as CuPy arrays or Numba device arrays—the device pointer is captured rather than the data. No copy occurs, and modifications to the array are visible to subsequent kernel calls.

The following example demonstrates this behavior. The global PRICES device array is mutated after the first kernel call, and the second kernel call sees the updated values:

Demonstrating device array capture by pointer

import numpy as np
from numba import cuda

# Global device array - pointer is captured, not data
PRICES = cuda.to_device(
    np.array([10.0, 25.0, 5.0, 15.0, 30.0], dtype=np.float32)
)

@cuda.jit
def compute_totals(quantities, totals):
    i = cuda.grid(1)
    if i < totals.size:
        totals[i] = quantities[i] * PRICES[i]

d_quantities = cuda.to_device(
    np.array([1.0, 1.0, 1.0, 1.0, 1.0], dtype=np.float32)
)
d_totals = cuda.device_array(5, dtype=np.float32)

# First kernel call
compute_totals[1, 32](d_quantities, d_totals)
print(d_totals.copy_to_host())  # [10. 25.  5. 15. 30.]

# Mutate the device array in-place
PRICES.copy_to_device(
    np.array([20.0, 50.0, 10.0, 30.0, 60.0], dtype=np.float32)
)

# Second kernel call sees the updated values
compute_totals[1, 32](d_quantities, d_totals)
print(d_totals.copy_to_host())  # [20. 50. 10. 30. 60.]

Running the above code prints:

[10. 25.  5. 15. 30.]
[20. 50. 10. 30. 60.]

Note that the outputs are different—the mutation to PRICES after the first kernel call is visible to the second call, unlike with host arrays.

This makes device arrays suitable for global state that needs to be updated between kernel calls without recompilation.

Note

Kernels and device functions that capture global device arrays cannot use cache=True. Because the device pointer is embedded in the compiled code, caching would serialize an invalid pointer. Attempting to cache such a kernel will raise a PicklingError. See On-disk Kernel Caching for more information on kernel caching.