### CUDA 11 Installation Example

Source: https://numba.readthedocs.io/en/stable/reference/utils.html

Instructions or code snippet related to installing Numba with CUDA 11 support. This might involve specific environment variables or package versions.

```bash
# Example: Ensure CUDA 11 toolkit is installed and Numba is configured correctly
# pip install numba
# export NUMBA_CUDA_LIBDIR=/path/to/cuda/lib64

```

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

### CUDA 12 Installation Example

Source: https://numba.readthedocs.io/en/stable/reference/utils.html

Instructions or code snippet related to installing Numba with CUDA 12 support. This might involve specific environment variables or package versions.

```bash
# Example: Ensure CUDA 12 toolkit is installed and Numba is configured correctly
# pip install numba
# export NUMBA_CUDA_LIBDIR=/path/to/cuda/lib64

```

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

### Example: Basic CUDA Ufunc

Source: https://numba.readthedocs.io/en/stable/release/0.58.1-notes.html

Demonstrates a basic CUDA universal function (ufunc) implementation in Numba. Ensure Numba and CUDA are installed.

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

x = np.arange(1000000).astype(np.float32)
y = x * 2
out = np.empty_like(x)

threadsperblock = 32
blockspergrid = (x.size + (threadsperblock - 1)) // threadsperblock

add_kernel[blockspergrid, threadsperblock](x, y, out)

print(out[:10])
```

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

### Example: C and Numba ABI Examples

Source: https://numba.readthedocs.io/en/stable/release/0.65.0-notes.html

Provides examples of how to compile Python functions using Numba for use with the C ABI and demonstrates Numba's ABI compatibility with C.

```python
from numba import njit, cfunc, types
import numpy as np

# Example function to be compiled with C ABI
@njit(types.int64(types.int64, types.int64), cache=True)
def add_c_abi(a, b):
    return a + b

# Example of a C function pointer type
add_ptr_type = types.int64(types.int64, types.int64)

# Example of using a C function pointer within Numba
@njit(cache=True)
def call_c_ptr(func_ptr, a, b):
    return func_ptr(a, b)

# Get the C function pointer for add_c_abi
c_add_ptr = add_c_abi.get_c_function_pointer()

# Call the C function pointer
result = call_c_ptr(c_add_ptr, 5, 7)
print(f"Result of calling C ABI function via pointer: {result}")

# Example demonstrating direct C ABI compatibility
# Assume 'external_c_add' is a function compiled with C ABI elsewhere
# @cfunc(add_ptr_type)
def external_c_add(a, b):
#     return a + b # This would be the C implementation
#     pass

# Numba can call functions compiled with the C ABI
# result_external = call_c_ptr(external_c_add, 10, 15)
# print(f"Result from external C ABI function: {result_external}")

print("C and Numba ABI examples demonstrated. Requires compilation and potential external linking.")
```

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

### Dividing Click Data into Sessions Example

Source: https://numba.readthedocs.io/en/stable/user/vectorize.html

An example illustrating how to group click data into sessions on the GPU.

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

@cuda.jit
def sessionize(timestamps, session_ids, session_timeout):
    i = cuda.grid(1)
    if i > 0:
        # Calculate time difference with the previous click
        time_diff = timestamps[i] - timestamps[i-1]
        # If the difference exceeds the timeout, start a new session
        if time_diff > session_timeout:
            session_ids[i] = session_ids[i-1] + 1
        else:
            session_ids[i] = session_ids[i-1]
    else:
        # First click always starts a new session
        session_ids[i] = 0

```

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

### Calling Device Functions Example

Source: https://numba.readthedocs.io/en/stable/user/vectorize.html

Example showing how to define and call a device function from a CUDA kernel.

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

@cuda.jit(device=True)
def device_multiply(a, b):
    return a * b

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

```

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

### Example: Device Selection

Source: https://numba.readthedocs.io/en/stable/cuda/faq.html

Shows how to select and manage CUDA devices programmatically.

```python
from numba import cuda

# Get the number of available CUDA devices
num_devices = cuda.device_count()
print(f"Number of CUDA devices found: {num_devices}")

if num_devices > 0:
    # Select a specific device (e.g., the first one, index 0)
    cuda.select_device(0)
    print(f"Selected device: {cuda.current_device().name.decode()}")

    # You can also get information about the current device
    device = cuda.current_device()
    print(f"Device Compute Capability: {device.compute_capability}")

    # Perform operations on the selected device
    # ...

    # Reset the device if needed (releases resources)
    # cuda.close()
else:
    print("No CUDA devices available.")

```

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

### Clone Numba from GitHub

Source: https://numba.readthedocs.io/en/stable/user/installing.html

Clone the Numba repository from GitHub to build from source. Ensure llvmlite is installed first, following its installation guide.

```bash
$ git clone https://github.com/numba/numba.git

```

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

### 2D Kernel Instantiation Example

Source: https://numba.readthedocs.io/en/stable/_sources/cuda/kernels.rst.txt

Illustrates how to manually calculate the grid and block dimensions for launching a 2D kernel. Uses `math.ceil` to ensure all array elements are covered.

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

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

### Example: Context Management

Source: https://numba.readthedocs.io/en/stable/developer/mission.html

Shows how to manage CUDA contexts, including creating and destroying them.

```python
from numba import cuda

# Create a new context for device 0
with cuda.device(0):
    # Operations within this context will use device 0
    pass

# The context is automatically managed (created/destroyed) by the 'with' statement

```

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

### Example: Device Detection and Enquiry

Source: https://numba.readthedocs.io/en/stable/developer/mission.html

Demonstrates how to detect and query information about available CUDA devices.

```python
from numba import cuda

# Get the number of available devices
num_devices = cuda.device_count()

# Get the current device
current_device = cuda.get_current_device()

# Get properties of a specific device
properties = cuda.get_device_properties(0)
print(f"Device Name: {properties['name']}")

```

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

### Example CUDA Kernel with CuPy Arrays

Source: https://numba.readthedocs.io/en/stable/cuda/cuda_array_interface.html

This example demonstrates a typical usage pattern where a Numba CUDA kernel operates on CuPy arrays. Ensure CuPy is installed and Numba is configured for CUDA development.

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

```

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

### Import necessary libraries for CUDA vector addition

Source: https://numba.readthedocs.io/en/stable/_sources/cuda/examples.rst.txt

Imports required for the vector addition example. Ensure Numba and CUDA are installed.

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

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

### Example: Stream Management

Source: https://numba.readthedocs.io/en/stable/developer/mission.html

Demonstrates creating, using, and synchronizing CUDA streams.

```python
from numba import cuda

# Create a new stream
stream = cuda.stream()

# Launch a kernel on the stream
# kernel[blocks, threads, stream](args...)

# Synchronize the stream to ensure completion
stream.synchronize()

```

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

### Get LLVM Pass Timings

Source: https://numba.readthedocs.io/en/stable/developer/llvm_timings.html

Example demonstrating how to retrieve LLVM pass timings. This is useful for performance analysis of the Numba compilation process.

```python
from numba import njit

@njit(debug=True)
def foo(x):
    return x + 1

foo(1)
```

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

### Example: Basic CUDA Kernel for Device Functions

Source: https://numba.readthedocs.io/en/stable/cuda/faq.html

Illustrates defining and calling a simple device function from a CUDA kernel.

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

# Define a function that can be called from device code
@cuda.jit(device=True)
def square(x):
    return x * x

# Define a kernel that calls the device function
@cuda.jit
def device_function_kernel(input_array, output_array):
    idx = cuda.grid(1)
    if idx < input_array.shape[0]:
        # Call the device function
        output_array[idx] = square(input_array[idx])

# Example usage:
n = 1000
input_data = np.arange(n, dtype=np.float32)
output_data = np.empty_like(input_data)

input_device = cuda.to_device(input_data)
output_device = cuda.to_device(output_data)

threadsperblock = 128
blockspergrid = (n + (threadsperblock - 1)) // threadsperblock

device_function_kernel[blockspergrid, threadsperblock](input_device, output_device)

output_host = output_device.copy_to_host()
print(output_host[:10])

```

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

### Device Detection and Enquiry

Source: https://numba.readthedocs.io/en/stable/developer/numba-runtime.html

Code examples for detecting and enquiring about CUDA-enabled devices using Numba. This includes checking availability and getting device properties.

```python
from numba import cuda

# Check if CUDA is available on the system
is_cuda_available = cuda.is_available()
print(f"CUDA available: {is_cuda_available}")

if is_cuda_available:
    # Get the number of CUDA devices available
    device_count = cuda.device_count()
    print(f"Number of CUDA devices: {device_count}")

    # Get properties of each device
    for i in range(device_count):
        device = cuda.device(i)
        print(f"\n--- Device {i} ---")
        print(f"  Name: {device.name}")
        print(f"  Compute Capability: {device.compute_capability[0]}.{device.compute_capability[1]}")
        print(f"  Total Memory: {device.total_memory / (1024**3):.2f} GB")
        print(f"  Multiprocessor Count: {device.multiprocessor_count}")

    # Get the current device (usually the default device)
    current_device = cuda.get_current_device()
    print(f"\nCurrent device index: {current_device.id}")
    print(f"Current device name: {current_device.name}")

    # You can also select a different device
    # if device_count > 1:
    #     cuda.select_device(1) # Select the second device
    #     print(f"Switched to device: {cuda.get_current_device().name}")
else:
    print("No CUDA devices found or CUDA is not available.")

# Expected output (will vary based on your hardware):
# CUDA available: True
# Number of CUDA devices: 1
# 
# --- Device 0 ---
#   Name: NVIDIA GeForce RTX 3080
#   Compute Capability: 8.6
#   Total Memory: 10.00 GB
#   Multiprocessor Count: 68
# 
# Current device index: 0
# Current device name: NVIDIA GeForce RTX 3080
```

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

### Device Management Example

Source: https://numba.readthedocs.io/en/stable/cuda-reference/host.html

Shows how to select a specific CUDA device for use. This is essential when multiple GPUs are present and you need to target a particular one.

```python
from numba import cuda

# Select device 1
with cuda.device_api.get_device(1).context:
    # All subsequent CUDA operations will use device 1
    print(f"Using device: {cuda.get_current_device()}")
```

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

### Getting LLVM Pass Timings in Numba

Source: https://numba.readthedocs.io/en/stable/developer/custom_pipeline.html

Provides an example of how to retrieve LLVM pass timings within Numba. This is useful for performance analysis and debugging the compilation process.

```python
from numba import njit

@njit(pipeline_options={'print_llvm_passes': True})
def my_function(x):
    return x * 2
```

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

### Thread Indexing in CUDA Kernels

Source: https://numba.readthedocs.io/en/stable/developer/custom_pipeline.html

CUDA Kernel API example demonstrating how to get thread indices within a kernel. Essential for parallel data access and computation.

```python
from numba import cuda

@cuda.jit
def thread_indexing_example():
    # Get 1D thread index
    idx1d = cuda.grid(1)

    # Get 2D thread indices
    idx_x, idx_y = cuda.grid(2)

    # Get 3D thread indices
    idx_x, idx_y, idx_z = cuda.grid(3)
```

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

### Example: Using Grid Groups for Cooperative Launches

Source: https://numba.readthedocs.io/en/stable/cuda/faq.html

Demonstrates using Numba's cooperative groups API to manage cooperative kernel launches.

```python
from numba import cuda

# This is a conceptual example. Actual cooperative group usage requires specific hardware features
# and careful synchronization. Numba's support is evolving.

@cuda.jit
def cooperative_kernel(x):
    # Get the current thread's ID within the block
    tid = cuda.threadIdx.x
    # Get the current thread's ID within the grid
    grid_id = cuda.grid(0)

    # Use cooperative groups API (conceptual)
    # For example, to synchronize across threads within a grid group
    # This requires the kernel to be launched with cooperative=True
    # group = cuda.cg.grid_group()
    # group.sync()

    # Placeholder for operations that might use cooperative groups
    if tid == 0:
        print(f"Block {grid_id} starting cooperative work...")

# To launch this kernel cooperatively:
# threadsperblock = 128
# blockspergrid = 10
# cuda.launch_kernel(cooperative_kernel, blockspergrid, threadsperblock, cooperative=True, args=(...))

```

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

### Using parallel_diagnostics with level 4

Source: https://numba.readthedocs.io/en/stable/_sources/user/parallel.rst.txt

This example demonstrates how to use the `parallel_diagnostics` method with a verbosity level of 4 to get detailed information about Numba's parallelization transforms for a given function.

```python
@njit(parallel=True)
def test(x):
    n = x.shape[0]
    a = np.sin(x)
    b = np.cos(a * a)
    acc = 0
    for i in prange(n - 2):
        for j in prange(n - 1):
            acc += b[i] + b[j + 1]
    return acc

test(np.arange(10))

test.parallel_diagnostics(level=4)
```

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

### Example: Monte Carlo Integration (CUDA)

Source: https://numba.readthedocs.io/en/stable/release/0.65.0-notes.html

Demonstrates Monte Carlo integration on the GPU using Numba CUDA.

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

@cuda.jit
def monte_carlo_integration_kernel(rng_states, n_samples, result):
    tid = cuda.grid(1)
    if tid < result.shape[0]:
        # Use the thread-local RNG state
        r = cuda.random.xoroshiro128p_uniform_float64(rng_states, tid)
        # Example: integrate x^2 from 0 to 1
        result[tid] = r * r

def integrate_monte_carlo(n_samples_per_thread, n_threads):
    rng_states = cuda.random.create_xoroshiro128p_states(n_threads, seed=12345)
    results = np.zeros(n_threads, dtype=np.float64)

    threadsperblock = (128,)
    blockspergrid = (n_threads // threadsperblock[0] + 1,)

    monte_carlo_integration_kernel[blockspergrid, threadsperblock](rng_states, n_samples_per_thread, results)

    # Average the results from all threads
    final_result = np.mean(results)
    return final_result * n_samples_per_thread # Scale by total samples

# Example usage
# n_samples = 1000000
# n_threads = 100
# samples_per_thread = n_samples // n_threads
# integral = integrate_monte_carlo(samples_per_thread, n_threads)
# print(f"Approximate integral of x^2 from 0 to 1: {integral}")

print("Monte Carlo Integration example structure defined. Requires execution.")
```

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

### Configuring and Launching Kernels

Source: https://numba.readthedocs.io/en/stable/cuda-reference/kernel.html

Demonstrates the idiomatic way to configure and launch CUDA kernels using Numba's Dispatcher objects with subscripting for launch parameters.

```APIDOC
## Dispatcher Configuration and Launch

### Description
Configure and launch CUDA kernels using Numba's Dispatcher objects. The `griddim` and `blockdim` arguments specify the grid and thread block sizes, respectively. The `stream` parameter is an optional CUDA stream, and `sharedmem` specifies the dynamic shared memory size in bytes. The configuration object returned by subscripting can then be called with kernel arguments.

### Method
Subscripting the Dispatcher object followed by a call.

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

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```python
# Assuming 'func' is a function decorated with @cuda.jit
func[griddim, blockdim, stream, sharedmem](x, y, z)
```

### Response
#### Success Response (200)
Kernel execution.

#### Response Example
None
```

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

### Compiling code ahead of time (AOT) - Standalone example

Source: https://numba.readthedocs.io/en/stable/user/examples.html

Shows how to compile Numba functions ahead of time (AOT) into a shared library. This allows using Numba-compiled code without needing Numba installed at runtime.

```python
from numba import njit, types
from numba.pycc import CC

# Create a compilation context
cc = CC('my_module')

# Define a function to be compiled
@cc.export('add_floats', 'float64(float64, float64)')
@njit
def add_floats(x, y):
    return x + y

# Compile the module to a shared library
# This command would typically be run from the command line:
# python my_module.py
# cc.compile()

# After compilation, you can import and use it:
# import my_module_impl
# print(my_module_impl.add_floats(1.5, 2.5))
```

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

### Standalone AOT Compilation Example

Source: https://numba.readthedocs.io/en/stable/_sources/user/pycc.rst.txt

This snippet demonstrates how to use Numba's `CC` class for Ahead-of-Time compilation. It defines functions with explicit signatures and compiles them into a standalone extension module named 'my_module'. Ensure `setuptools` is installed for compilation.

```python
from numba.pycc import CC

cc = CC('my_module')
# Uncomment the following line to print out the compilation steps
#cc.verbose = True

@cc.export('multf', 'f8(f8, f8)')
@cc.export('multi', 'i4(i4, i4)')
def mult(a, b):
    return a * b

@cc.export('square', 'f8(f8)')
def square(a):
    return a ** 2

if __name__ == "__main__":
    cc.compile()
```

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

### Complete Example: Linking and Calling Functions

Source: https://numba.readthedocs.io/en/stable/developer/numba-runtime.html

A complete example demonstrating how to link a C library and call its functions from Numba. This involves CFFI for loading the library and Numba for wrapping the calls.

```python
from numba import njit, cgтуре, cffi_support
import numpy as np
from cffi import FFI

# --- Step 1: Define and compile a simple C library ---
ffi = FFI()
ffi.cdef("int multiply(int x, int y);\n         void add_to_array(int* arr, int size, int value);")
ffi.set_source("_mymathlib", """ 
    int multiply(int x, int y) {
        return x * y;
    }
    void add_to_array(int* arr, int size, int value) {
        for (int i = 0; i < size; ++i) {
            arr[i] += value;
        }
    }
""", 
    libraries=[],
    library_dirs=[],
    include_dirs=[])

# Compile the C code into a shared library (requires a C compiler)
try:
    ffi.compile(verbose=True)
    import _mymathlib
    lib = _mymathlib.lib
    cffi_support.register_module(_mymathlib)
    print("C library compiled and loaded successfully.")
    c_lib_available = True
except Exception as e:
    print(f"Failed to compile or load C library: {e}. Skipping C function calls.")
    c_lib_available = False

# --- Step 2: Wrap C functions using Numba ---

# Wrap the 'multiply' function
@njit(cgтуре.int64(cgтуре.int64, cgтуре.int64), cache=True)
def numba_multiply(x, y):
    # This Numba function calls the C function 'multiply' from the loaded library
    # Numba uses CFFI integration to find and call the C function.
    if c_lib_available:
        return lib.multiply(x, y)
    else:
        return x * y # Fallback if C lib is not available

# Wrap the 'add_to_array' function
# Signature: void(pointer_to_int, int_size, int_value)
c_add_array_sig = cgтуре.void(cgтуре.int64ptr, cgтуре.int64, cgтуре.int64)

@njit(c_add_array_sig, cache=True)
def numba_add_to_array(arr_ptr, size, value):
    if c_lib_available:
        lib.add_to_array(arr_ptr, size, value)
    else:
        # Simulate the operation if C lib is not available
        arr = np.frombuffer(arr_ptr, dtype=np.int64, count=size)
        arr += value

# --- Step 3: Use the Numba-wrapped C functions ---

if c_lib_available:
    # Example 1: Using multiply
    num1 = np.int64(10)
    num2 = np.int64(5)
    result_mult = numba_multiply(num1, num2)
    print(f"Numba multiply({num1}, {num2}) = {result_mult}")
    # Expected output: Numba multiply(10, 5) = 50

    # Example 2: Using add_to_array
    data_array = np.arange(10, dtype=np.int64)
    print(f"Original array: {data_array}")
    
    array_ptr = data_array.ctypes.data
    array_size = len(data_array)
    add_value = np.int64(3)
    
    numba_add_to_array(array_ptr, array_size, add_value)
    print(f"Array after adding {add_value}: {data_array}")
    # Expected output: Array after adding 3: [ 3  4  5  6  7  8  9 10 11 12]
else:
    print("Skipping function calls as C library is not available.")

print("Complete example finished.")
```

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

### Example: Calling Device Functions in CUDA ufuncs (Version 2)

Source: https://numba.readthedocs.io/en/stable/developer/hashing.html

Demonstrates calling device functions within a CUDA ufunc, possibly with different device functions or ufunc logic.

```python
from numba import cuda, int32

@cuda.jit('int32(int32)', device=True)
def square_int(x):
    return x * x

@cuda.vectorize('int32(int32)')
def square_ufunc(x):
    return square_int(x)

input_array = cuda.to_device(np.arange(10, dtype=np.int32))
output_array = square_ufunc(input_array)

print(output_array.copy_to_host())
```

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

### Set and Get Parallel Chunk Size

Source: https://numba.readthedocs.io/en/stable/_sources/user/parallel.rst.txt

Demonstrates setting a specific chunk size (4) outside a parallel region, which is reset to 0 inside the `prange` loop. If the `prange` is not executed in parallel, the chunk size inside would remain 4. This example uses `set_parallel_chunksize` and `get_parallel_chunksize`.

```python
from numba import prange, set_parallel_chunksize, get_parallel_chunksize


def func1():
    set_parallel_chunksize(4)
    print(f"Outside prange: {get_parallel_chunksize()}")
    for i in prange(10):
        # The chunk size is reset to 0 upon entering a parallel region
        # and restored upon exiting.
        if i == 0:
            print(f"Inside prange: {get_parallel_chunksize()}")
    print(f"After prange: {get_parallel_chunksize()}")

def func2():
    set_parallel_chunksize(8)
    print(f"Outside prange: {get_parallel_chunksize()}")
    for i in prange(10):
        if i == 0:
            print(f"Inside prange: {get_parallel_chunksize()}")
    print(f"After prange: {get_parallel_chunksize()}")


func1()
func2()
```

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

### Install numba-cuda with conda

Source: https://numba.readthedocs.io/en/stable/_sources/cuda/overview.rst.txt

Install the numba-cuda package from the conda-forge channel using conda. This is an alternative installation method for users who prefer conda.

```bash
conda install conda-forge::numba-cuda
```

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

### gdb Session Example

Source: https://numba.readthedocs.io/en/stable/user/troubleshoot.html

This example demonstrates a typical gdb session after hitting a breakpoint set by `numba.gdb()`. It shows how to step through the code, inspect local variables, and set new breakpoints.

```bash
$ NUMBA_OPT=0 NUMBA_EXTEND_VARIABLE_LIFETIMES=1 python demo_gdb.py
...
Breakpoint 1, 0x00007fb75238d830 in numba_gdb_breakpoint () from numba/_helperlib.cpython-39-x86_64-linux-gnu.so
(gdb) s
Single stepping until exit from function numba_gdb_breakpoint,
which has no line number information.
0x00007fb75233e1cf in numba::misc::gdb_hook::hook_gdb::_3clocals_3e::impl_242[abi:c8tJTIeFCjyCbUFRqqOAK_2f6h0phxApMogijRBAA_3d](StarArgTuple) ()
(gdb) s
Single stepping until exit from function _ZN5numba4misc8gdb_hook8hook_gdb12_3clocals_3e8impl_242B44c8tJTIeFCjyCbUFRqqOAK_2f6h0phxApMogijRBAA_3dE12StarArgTuple,
which has no line number information.
__main__::foo_241[abi:c8tJTC_2fWgEeGLSgydRTQUgiqKEZ6gEoDvQJmaQIA](long long) (a=123) at demo_gdb.py:7
7           c = a * 2.34
(gdb) l
2
3       @njit(debug=True)
4       def foo(a):
5           b = a + 1
6           gdb() # instruct Numba to attach gdb at this location and pause execution
7           c = a * 2.34
8           d = (a, b, c)
9           print(a, b, c, d)
10
11      r= foo(123)
(gdb) p a
$1 = 123
(gdb) p b
$2 = 124
(gdb) p c
$3 = 0
(gdb) b 9
Breakpoint 2 at 0x7fb73d1f7287: file demo_gdb.py, line 9.
(gdb) c
Continuing.

Breakpoint 2, __main__::foo_241[abi:c8tJTC_2fWgEeGLSgydRTQUgiqKEZ6gEoDvQJmaQIA](long long) (a=123) at demo_gdb.py:9
9           print(a, b, c, d)
(gdb) info locals
b = 124
c = 287.81999999999999
d = {f0 = 123, f1 = 124, f2 = 287.81999999999999}

```

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

### CUDA 12 Installation

Source: https://numba.readthedocs.io/en/stable/developer/numba-runtime.html

Instructions for installing Numba with support for CUDA 12. This typically involves ensuring the correct CUDA toolkit and driver versions are installed.

```bash
# Ensure you have the NVIDIA driver compatible with CUDA 12 installed.
# Download and install the CUDA Toolkit 12.x from the NVIDIA website.
# Make sure the CUDA environment variables (like CUDA_HOME, PATH) are set correctly.

# Install Numba using pip, ensuring it can find the CUDA toolkit.
# It's often recommended to install within a virtual environment.

pip install --upgrade numba

# If Numba fails to find CUDA, you might need to specify the CUDA path during installation
# or ensure it's in your system's PATH.
# Example: export PATH=/usr/local/cuda-12.0/bin:$PATH
# Example: export LD_LIBRARY_PATH=/usr/local/cuda-12.0/lib64:$LD_LIBRARY_PATH

# Verify the installation by running a simple Numba CUDA example.
# python -c "from numba import cuda; print(cuda.is_available())"

```

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

### Example: Basic CUDA Kernel

Source: https://numba.readthedocs.io/en/stable/release/0.59.1-notes.html

A simple CUDA kernel for vector addition. Demonstrates basic kernel definition and invocation.

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

x = np.arange(1000000).astype(np.float32)
y = x * 2
out = np.empty_like(x)

threadsperblock = 32
blockspergrid = (x.shape[0] + (threadsperblock - 1)) // threadsperblock

add_kernel[blockspergrid, threadsperblock](x, y, out)

cuda.synchronize()

print(out[:10])
```

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

### Install CUDA 11 support with Conda

Source: https://numba.readthedocs.io/en/stable/_sources/user/installing.rst.txt

Install the necessary CUDA 11 toolkit components using conda-forge. Ensure you have compatible NVIDIA drivers installed separately.

```bash
$ conda install -c conda-forge cudatoolkit "cuda-version>=11.2,<12.0"
```

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

### GDB Session Example

Source: https://numba.readthedocs.io/en/stable/_sources/user/troubleshoot.rst.txt

This example demonstrates a typical gdb session after hitting a breakpoint set by numba.gdb(). It shows stepping through code, listing source lines, inspecting variables, setting new breakpoints, and continuing execution.

```none
$ NUMBA_OPT=0 NUMBA_EXTEND_VARIABLE_LIFETIMES=1 python demo_gdb.py
...
Breakpoint 1, 0x00007fb75238d830 in numba_gdb_breakpoint () from numba/_helperlib.cpython-39-x86_64-linux-gnu.so
(gdb) s
Single stepping until exit from function numba_gdb_breakpoint,
which has no line number information.
0x00007fb75233e1cf in numba::misc::gdb_hook::hook_gdb::_3clocals_3e::impl_242[abi:c8tJTIeFCjyCbUFRqqOAK_2f6h0phxApMogijRBAA_3d](StarArgTuple) ()
(gdb) s
Single stepping until exit from function _ZN5numba4misc8gdb_hook8hook_gdb12_3clocals_3e8impl_242B44c8tJTIeFCjyCbUFRqqOAK_2f6h0phxApMogijRBAA_3dE12StarArgTuple,
which has no line number information.
__main__::foo_241[abi:c8tJTC_2fWgEeGLSgydRTQUgiqKEZ6gEoDvQJmaQIA](long long) (a=123) at demo_gdb.py:7
7           c = a * 2.34
(gdb) l
2
3       @njit(debug=True)
4       def foo(a):
5           b = a + 1
6           gdb() # instruct Numba to attach gdb at this location and pause execution
7           c = a * 2.34
8           d = (a, b, c)
9           print(a, b, c, d)
10
11      r= foo(123)
(gdb) p a
$1 = 123
(gdb) p b
$2 = 124
(gdb) p c
$3 = 0
(gdb) b 9
Breakpoint 2 at 0x7fb73d1f7287: file demo_gdb.py, line 9.
(gdb) c
Continuing.

Breakpoint 2, __main__::foo_241[abi:c8tJTC_2fWgEeGLSgydRTQUgiqKEZ6gEoDvQJmaQIA](long long) (a=123) at demo_gdb.py:9
9           print(a, b, c, d)
(gdb) info locals
b = 124
c = 287.81999999999999
d = {f0 = 123, f1 = 124, f2 = 287.81999999999999}
```

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

### Install Numba using Pip

Source: https://numba.readthedocs.io/en/stable/_sources/user/installing.rst.txt

Install Numba using pip. This method downloads pre-built wheels for Windows, Mac, and Linux. LLVM is bundled and does not need to be installed separately.

```bash
$ pip install numba
```

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

### Example: Invoking a CUDA Kernel with Block Size Calculation

Source: https://numba.readthedocs.io/en/stable/cuda/cudapysupported.html

Demonstrates how to launch a CUDA kernel, calculating the appropriate grid and block dimensions based on the input array size. This ensures efficient parallel execution.

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

# Assume vector_add kernel is defined as above

size = 1000000
x = np.arange(size, dtype=np.float32)
y = np.arange(size, dtype=np.float32)

out = np.zeros_like(x)

threads_per_block = 128
blocks_per_grid = (size + (threads_per_block - 1)) // threads_per_block

vector_add[blocks_per_grid, threads_per_block](x, y, out)

cuda.synchronize()

print(out[:10])

```

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

### Example: Using NRT API from C

Source: https://numba.readthedocs.io/en/stable/developer/numba-runtime.html

Demonstrates how to use the NRT_api_functions struct from C code to manage memory. It shows allocating memory, wrapping it with MemInfo, and managing references.

```c
#include <stdio.h>
#include "numba/core/runtime/nrt_external.h"

void my_dtor(void *ptr) {
    free(ptr);
}

NRT_MemInfo* my_allocate(NRT_api_functions *nrt) {
    /* heap allocate some memory */
    void * data = malloc(10);
    /* wrap the allocated memory; yield a new reference */
    NRT_MemInfo *mi = nrt->manage_memory(data, my_dtor);
    /* acquire reference */
    nrt->acquire(mi);
    /* release reference */
    nrt->release(mi);
    return mi;
}


```

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

### CUDA 11 Installation Notes

Source: https://numba.readthedocs.io/en/stable/release/0.58.0-notes.html

Provides specific instructions or considerations for installing Numba with CUDA 11 support. This may include required driver versions or specific installation steps.

```text
Ensure you have the correct NVIDIA driver version installed for CUDA 11 compatibility.
```

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

### CUDA 12 Installation Notes

Source: https://numba.readthedocs.io/en/stable/release/0.58.0-notes.html

Provides specific instructions or considerations for installing Numba with CUDA 12 support. This may include required driver versions or specific installation steps.

```text
Ensure you have the correct NVIDIA driver version installed for CUDA 12 compatibility.
```

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

### Example: CUDA Minor Version Compatibility (MVC) Support

Source: https://numba.readthedocs.io/en/stable/developer/mission.html

Enabling CUDA Minor Version Compatibility (MVC) support for Numba.

```bash
export NUMBA_CUDA_ENABLE_MVC=1

```

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

### Example: Basic CUDA ufunc (Version 2)

Source: https://numba.readthedocs.io/en/stable/developer/hashing.html

A basic example of a CUDA universal function (ufunc) using Numba, similar to the first example but potentially with different type signatures or usage.

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

@cuda.vectorize('int32(int32, int32)')
def multiply_int_ufunc(x, y):
    return x * y

x = np.arange(10, dtype=np.int32)
y = np.arange(10, dtype=np.int32)

result = multiply_int_ufunc(x, y)
print(result)
```