[FEA] Make `cuda.core.utils.StridedMemoryView` recognized by numba-cuda

[leofang]

Closes #122

[leofang]

xref: NVIDIA/cccl#4072 (comment)

[cpcloud]

Is there an example somewhere that would illustrate what it would mean for numba-cuda to recognize StridedMemoryView? Where is/would it be used?

[leofang]

This is what I currently have in mind:

1. Explicit construction & reuse of views

# Using CuPy as example here, but it should work for any array objects supporting either DLPack or CAI
import cupy as cp
from cuda.core import utils, Device
from numba import cuda

dev = Device().set_current()
arr = cp.zeros(10, dtype=cp.float32)
stream = dev.create_stream()

@utils.args_viewable_as_strided_memory((0,))
def my_workload(arr, stream):
    view = arr.view(stream.handle)
    assert isinstance(view, utils.StridedMemoryView)

    @cuda.jit
    def my_kernel(arr):
        tid = cuda.grid(1)
        if tid < len(arr):
            view[tid] += arr.dtype(42.)

    grid = ...
    block = ...
    for _ in range(100):
        # view can be accessed as if it is an array argument 
        my_kernel[grid, block, 0, stream](view)

my_workload(arr, stream)
stream.sync()
assert cp.all(arr == 420.)

2. Implicit construction for one-off use cases. I am not a big fan of this, but I can see why people love it.

import cupy as cp
from cuda.core import Device
from numba import cuda

dev = Device().set_current()
arr = cp.zeros(10, dtype=cp.float32)

@cuda.jit
def my_kernel(arr):
    tid = cuda.grid(1)
    if tid < len(arr):
        arr[tid] += arr.dtype(42.)

grid = ...
block = ...
# A StridedMemoryView is implicitly constructed inside numba
my_kernel[grid, block, 0, stream](arr)

Basically, this is saying numba-cuda need not implement any support for DLPack (#122, #531). Just leave the job to cuda.core.

[leofang]

Basically, this is saying numba-cuda need not implement any support for DLPack (#122, #531). Just leave the job to cuda.core.

I'd really like to emphasize this design philosophy. It enables us to build reusable & performant components, that can be used by numba-cuda and beyond. For example, we already knew how to offer fast paths for building a StridedMemoryView for CuPy and PyTorch tensors (NVIDIA/cuda-python#180). It'd be really nice if we could do this optimization work in one place and accelerate everyone.