From 6cced6f137117a025380ec3625eb7e2d9381918d Mon Sep 17 00:00:00 2001 From: Max Date: Tue, 6 Oct 2026 13:15:09 +0200 Subject: [PATCH] Add C-contiguous array views CArray1D-4D --- CHANGELOG.md | 5 + docs/source/kernels/arguments.md | 34 +++++ src/cunumpy/_cuda_kernel.py | 110 ++++++++++++---- src/cunumpy/_emulation.py | 15 ++- .../cuda/include/cunumpy/array_view.cuh | 116 ++++++++++++++++- tests/unit/test_cuda_kernel.py | 123 ++++++++++++++++++ tests/unit/test_emulation.py | 31 +++++ 7 files changed, 400 insertions(+), 34 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 94dd615..7d0caef 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -17,6 +17,11 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 The former function names are removed without compatibility aliases. ### Added +- C-contiguous array views `CArray1D` to `CArray4D` in + `cunumpy/array_view.cuh`. They hold a pointer and shape only, so `a(i, j)` is + `data[i * shape[1] + j]`. As kernel parameters or struct fields, they reject + non-contiguous arrays and never copy them. `CudaStruct.from_signature` and + `from_pyccel_class` take `contiguous=True` or field names to generate them. - Device implementation selection via `kernels.set_device_kernel_implementation`, `get_device_kernel_implementation`, `use_device_kernel_implementation`, and `CUNUMPY_DEVICE_KERNEL_IMPLEMENTATION`. Accept `"cuda"` or automatic selection diff --git a/docs/source/kernels/arguments.md b/docs/source/kernels/arguments.md index 5e4e80f..87cf523 100644 --- a/docs/source/kernels/arguments.md +++ b/docs/source/kernels/arguments.md @@ -325,6 +325,40 @@ extern "C" __global__ void push(MarkerArgs m, double dt) { } ``` +### Contiguous views: `CArray2D` + +`Array2D` takes any view, so its strides are only known at run time and a +kernel cannot tell which index is the fast one. When an array is always +C-contiguous (a marker array, a grid), declare it as `CArray1D` to +`CArray4D` instead. The view holds a pointer and the shape, no strides, and +`m.markers(ip, 0)` is `data[ip * shape[1] + 0]`: the last index is always the +fast one, as in the row-major memory the host code uses. + +```python +MarkerArgs = xp.cuda.CudaStruct.from_pyccel_class( + "my_sim/kernel_arguments/pusher_args_kernels.py", + "MarkerArguments", + "MarkerArgs", + contiguous=["markers"], # or True for every array field +) +``` + +```c +struct MarkerArgs { + CArray2D markers; + long long n_markers; + Array1D valid; +}; +``` + +* A non-contiguous array (e.g. `markers[:, 0:3]`) raises `TypeError` at + packing or launch. It is **not** copied with `ascontiguousarray`, since + the kernel would write into the copy and the result would be lost. + Use `Array2D` for arguments that are sometimes views. +* A `CArrayND` converts to an `ArrayND`, so device helper functions + written for strided views also take it. +* `contiguous=` works the same in `CudaStruct.from_signature`. + ### Write the struct to a header Kernels in `.cu` files include the struct from a header. Generate it from the diff --git a/src/cunumpy/_cuda_kernel.py b/src/cunumpy/_cuda_kernel.py index c3f931f..2d0eb56 100644 --- a/src/cunumpy/_cuda_kernel.py +++ b/src/cunumpy/_cuda_kernel.py @@ -21,7 +21,9 @@ (from the shipped header ``cunumpy/array_view.cuh``, see :func:`cuda_include_dir`) takes a 2D CuPy array, contiguous or not, and receives its pointer, shape and strides, so that kernels index ``a(i, j)`` - like the pyccel kernels they are ported from; + like the pyccel kernels they are ported from; ``CArray2D`` is the + C-contiguous variant (shape only, ``a(i, j)`` is ``data[i * shape[1] + j]``), + which takes C-contiguous arrays only and raises for other views; * C++ function templates are instantiated with ``template_args``, and generated kernels (one source per variant) are compiled once per variant by :class:`CudaKernelVariants`. @@ -187,7 +189,11 @@ class CudaParameter(NamedTuple): The struct type, for a struct passed by value. view_ndim : int | None The number of dimensions, for an array view (``Array1D`` to - ``Array4D``, see :func:`cuda_include_dir`) passed by value. + ``Array4D`` or ``CArray1D`` to ``CArray4D``, see + :func:`cuda_include_dir`) passed by value. + contiguous : bool + Whether the array view is C-contiguous (``CArray2D``): packed + without strides, and only C-contiguous arrays are accepted. """ name: str @@ -196,6 +202,7 @@ class CudaParameter(NamedTuple): pointer: bool struct: CudaStruct | None = None view_ndim: int | None = None + contiguous: bool = False # C types (after removing qualifiers) and their NumPy dtypes. ``long`` is 64 bit, @@ -259,30 +266,30 @@ class CudaParameter(NamedTuple): _QUALIFIERS = {"const", "volatile", "__restrict__", "__restrict", "restrict"} _COMPLEX = re.compile(r"(?:(?:thrust|cuda::std)::)?complex\s*<\s*(float|double)\s*>") -# Array1D to Array4D (cunumpy/array_view.cuh), T a scalar type of _CTYPES -_VIEW = re.compile(r"\bArray([1234])D\s*<((?:[^<>]|complex<[^<>]*>)+?)>") +# Array1D to Array4D and CArray1D to CArray4D (cunumpy/array_view.cuh), +# T a scalar type of _CTYPES +_VIEW = re.compile(r"\b(C?)Array([1234])D\s*<((?:[^<>]|complex<[^<>]*>)+?)>") _TOKEN = re.compile( - r"Array[1234]D<[^<>]*(?:<[^<>]*>[^<>]*)?>|complex<(?:float|double)>" + r"C?Array[1234]D<[^<>]*(?:<[^<>]*>[^<>]*)?>|complex<(?:float|double)>" r"|[A-Za-z_]\w*|\*|\[\s*\]", ) def _normalize_view(match: re.Match) -> str: """``Array2D< const double >`` -> ``Array2D``.""" - words = [w for w in match.group(2).split() if w not in _QUALIFIERS] - return f"Array{match.group(1)}D<{' '.join(words)}>" - - -def _view_dtype(ndim: int) -> np.dtype: - """The structured dtype with the C layout of ``ArrayD``.""" - return np.dtype( - [ - ("data", np.uint64), - ("shape", np.int64, (ndim,)), - ("strides", np.int64, (ndim,)), - ], - align=True, - ) + words = [w for w in match.group(3).split() if w not in _QUALIFIERS] + return f"{match.group(1)}Array{match.group(2)}D<{' '.join(words)}>" + + +def _view_dtype(ndim: int, contiguous: bool = False) -> np.dtype: + """The structured dtype with the C layout of ``ArrayD``. + + ``CArrayD`` (`contiguous`) has no strides. + """ + fields = [("data", np.uint64), ("shape", np.int64, (ndim,))] + if not contiguous: + fields.append(("strides", np.int64, (ndim,))) + return np.dtype(fields, align=True) def ctype_of(dtype: Any) -> str: @@ -515,14 +522,15 @@ def _parse_parameter( return CudaParameter(name, ctype, struct.dtype, False, struct) view = _VIEW.fullmatch(ctype) if view is not None: - element = view.group(2) + element = view.group(3) if pointers or element not in _CTYPES: raise ValueError( f"cannot check the kernel parameter {text.strip()!r}: array views " f"take a scalar element type and are passed by value", ) - ndim = int(view.group(1)) - return CudaParameter(name, ctype, np.dtype(_CTYPES[element]), False, None, ndim) + ndim = int(view.group(2)) + dtype = np.dtype(_CTYPES[element]) + return CudaParameter(name, ctype, dtype, False, None, ndim, bool(view.group(1))) if ctype == "void" and pointers == 1: return CudaParameter(name, ctype, None, True) if pointers > 1 or ctype not in _CTYPES: @@ -714,10 +722,13 @@ def _view_checker(param: CudaParameter, index: int) -> Callable[[Any], Any]: """Checker for an array view: packs (pointer, shape, strides) of a device array. Strides are converted from bytes to elements; the array need not be - contiguous. + contiguous. A contiguous view (``CArray2D``) packs (pointer, shape) and + takes C-contiguous arrays only: it is never copied, since what the kernel + writes into a copy would be lost. """ ndim = param.view_ndim - dtype = _view_dtype(ndim) + contiguous = param.contiguous + dtype = _view_dtype(ndim, contiguous) def check(value: Any) -> Any: _check_device_array(param, index, value) @@ -725,6 +736,19 @@ def check(value: Any) -> Any: raise TypeError( f"{_describe(param, index)} must be a {ndim}D array, got {value.ndim}D", ) + if contiguous: + flags = getattr(value, "flags", None) + if flags is not None and not flags.c_contiguous: + raise TypeError( + f"{_describe(param, index)} must be C-contiguous: got a view " + f"with strides {tuple(value.strides)}; declare the parameter " + f"as Array{ndim}D to take strided views, or pass " + f"cupy.ascontiguousarray(...) and copy the result back", + ) + packed = np.zeros((), dtype=dtype) + packed["data"] = value.data.ptr + packed["shape"] = value.shape + return packed[()] itemsize = value.dtype.itemsize strides = [s // itemsize for s in value.strides] if any(s * itemsize != stride for s, stride in zip(strides, value.strides)): @@ -835,7 +859,7 @@ def _field_dtype(field: CudaParameter) -> np.dtype: if field.pointer: return np.dtype(np.uint64) if field.view_ndim is not None: - return _view_dtype(field.view_ndim) + return _view_dtype(field.view_ndim, field.contiguous) return field.dtype @@ -891,6 +915,27 @@ def _pyccel_ctype(annotation: Any, scalars: Mapping[str, str], what: str) -> str return f"Array{ndim}D<{ctype}>" +def _apply_contiguous( + fields: list[tuple[str, str]], + contiguous: bool | Iterable[str], + what: str, +) -> list[tuple[str, str]]: + """Make array fields ``CArrayD``: all of them, or those named.""" + if contiguous is True: + names = {f for f, ctype in fields if ctype.startswith("Array")} + elif contiguous is False: + return fields + else: + names = {contiguous} if isinstance(contiguous, str) else set(contiguous) + arrays = {f for f, ctype in fields if ctype.startswith("Array")} + if names - arrays: + raise ValueError( + f"{what}: contiguous names {sorted(names - arrays)} that are not " + f"array fields (array fields: {sorted(arrays)})", + ) + return [(f, f"C{ctype}" if f in names else ctype) for f, ctype in fields] + + def _header_guard(name: str) -> str: guard = re.sub(r"\W", "_", name).upper().strip("_") return guard if re.match(r"[A-Z_]", guard) else f"_{guard}" @@ -1084,6 +1129,7 @@ def from_signature( *, int_type: str = "long long", scalar_names: Mapping[str, str] | None = None, + contiguous: bool | Iterable[str] = False, ) -> CudaStruct: """Build the struct from the annotated parameters of a Python function. @@ -1109,6 +1155,10 @@ def from_signature( scalar_names : Mapping[str, str] | None Additional (or changed) mappings from annotation scalar names to C types, e.g. ``{"float": "float"}`` for single precision. + contiguous : bool | Iterable[str] + Array fields that become C-contiguous views (``CArray2D`` + instead of ``Array2D``): ``True`` for all of them, or their + names. Packing such a field raises for a non-contiguous array. Raises ------ @@ -1141,7 +1191,8 @@ def from_signature( continue what = f"parameter {param.name!r} of {getattr(func, '__qualname__', func)}" fields.append((param.name, _pyccel_ctype(param.annotation, scalars, what))) - return cls(name, fields) + what = f"{getattr(func, '__qualname__', func)}" + return cls(name, _apply_contiguous(fields, contiguous, what)) @classmethod def from_pyccel_class( @@ -1154,6 +1205,7 @@ def from_pyccel_class( scalar_names: Mapping[str, str] | None = None, exclude: Sequence[str] = (), attribute_names: bool = True, + contiguous: bool | Iterable[str] = False, ) -> CudaStruct: """Build the struct from the ``__init__`` of a class in a Python source file. @@ -1177,8 +1229,9 @@ def from_pyccel_class( Name of the class in the source. name : str | None Name of the struct type in C; `class_name` by default. - int_type, scalar_names - As for :meth:`from_signature`. + int_type, scalar_names, contiguous + As for :meth:`from_signature` (`contiguous` takes field names, i.e. + attribute names when `attribute_names` is set). exclude : Sequence[str] Parameter or attribute names that do not become fields, e.g. scratch arrays the host class allocates for itself. @@ -1217,6 +1270,7 @@ def from_pyccel_class( raise ValueError(f"{what} has no type annotation") annotation = _annotation_text(arg.annotation) fields.append((field, _pyccel_ctype(annotation, scalars, what))) + fields = _apply_contiguous(fields, contiguous, f"{class_name} in {where}") return cls(class_name if name is None else name, fields) def __repr__(self) -> str: diff --git a/src/cunumpy/_emulation.py b/src/cunumpy/_emulation.py index 499b888..d3e2de4 100644 --- a/src/cunumpy/_emulation.py +++ b/src/cunumpy/_emulation.py @@ -343,12 +343,18 @@ def emulate_cuda_kernel( f"argument {i} ({param.name}) must be a {param.view_ndim}D " f"array, got {value.ndim}D", ) + if param.contiguous and not value.flags.c_contiguous: + # as on the GPU: a copy would drop what the kernel writes + raise TypeError( + f"argument {i} ({param.name}) must be C-contiguous for " + f"{param.ctype}", + ) buffer = np.ascontiguousarray(value) path = tmp_path / f"{name}.bin" buffer.tofile(path) ctype = param.ctype if param.dtype is not None else "unsigned char" element = ( - re.match(r"Array\dD<(.*)>", ctype).group(1) + re.match(r"C?Array\dD<(.*)>", ctype).group(1) if param.view_ndim is not None else ctype ) @@ -362,10 +368,11 @@ def emulate_cuda_kernel( strides = ", ".join( f"{s // buffer.itemsize}LL" for s in buffer.strides ) + members = f"{{{shape}}}" + if not param.contiguous: + members += f", {{{strides}}}" globals_.append(f"static {ctype} {name}_view;") - inits.append( - f" {name}_view = {ctype}{{{name}, {{{shape}}}, {{{strides}}}}};", - ) + inits.append(f" {name}_view = {ctype}{{{name}, {members}}};") call_args.append(f"{name}_view") else: call_args.append(name) diff --git a/src/cunumpy/cuda/include/cunumpy/array_view.cuh b/src/cunumpy/cuda/include/cunumpy/array_view.cuh index c1ebe6a..d93667f 100644 --- a/src/cunumpy/cuda/include/cunumpy/array_view.cuh +++ b/src/cunumpy/cuda/include/cunumpy/array_view.cuh @@ -1,4 +1,4 @@ -// Strided array views for CUDA kernels, passed by value from Python. +// Array views for CUDA kernels, passed by value from Python. // // Array1D to Array4D describe a (possibly non-contiguous) // device array the way NumPy/CuPy do: a data pointer, a shape and strides. @@ -20,6 +20,19 @@ // the end of this file check the size. Member functions do not change the // layout. // +// CArray1D to CArray4D are C-contiguous (row-major) views: a data +// pointer and a shape, no strides. `a(i, j)` is `data[i * shape[1] + j]`, so +// the last index is always the fast one and the compiler knows it has unit +// stride. A parameter or field of these types only takes C-contiguous arrays; +// cunumpy raises for a non-contiguous view instead of copying it (a copy would +// silently drop what the kernel writes). Their layout, for ndim = 1 to 4: +// +// T* data; // 8 bytes +// long long shape[ndim]; // ndim * 8 bytes +// +// (sizeof == 8 * (1 + ndim)). A CArrayND converts to an ArrayND, so it +// can be passed to device functions written for strided views. +// // Bounds checks: compile with -DCUNUMPY_BOUNDS_CHECK to check every index // against the shape (an out-of-bounds index prints a message and traps the // kernel, which CuPy reports as a CUDA error). Without the macro, indexing is @@ -117,12 +130,111 @@ struct Array4D { } }; -// The layout the Python side packs: pointer, shape, strides, 8-byte aligned. +// C-contiguous views: shape only, the strides follow from it. + +template +struct CArray1D { + T* data; + long long shape[1]; + + __device__ __forceinline__ T& operator()(long long i) const { + CUNUMPY_CHECK_INDEX(i, 0, shape[0]); + return data[i]; + } + + // Number of elements. + __device__ __forceinline__ long long size() const { return shape[0]; } + + __host__ __device__ operator Array1D() const { + return Array1D{data, {shape[0]}, {1}}; + } +}; + +template +struct CArray2D { + T* data; + long long shape[2]; + + __device__ __forceinline__ T& operator()(long long i, long long j) const { + CUNUMPY_CHECK_INDEX(i, 0, shape[0]); + CUNUMPY_CHECK_INDEX(j, 1, shape[1]); + return data[i * shape[1] + j]; + } + + // Number of elements. + __device__ __forceinline__ long long size() const { + return shape[0] * shape[1]; + } + + __host__ __device__ operator Array2D() const { + return Array2D{data, {shape[0], shape[1]}, {shape[1], 1}}; + } +}; + +template +struct CArray3D { + T* data; + long long shape[3]; + + __device__ __forceinline__ T& operator()(long long i, long long j, + long long k) const { + CUNUMPY_CHECK_INDEX(i, 0, shape[0]); + CUNUMPY_CHECK_INDEX(j, 1, shape[1]); + CUNUMPY_CHECK_INDEX(k, 2, shape[2]); + return data[(i * shape[1] + j) * shape[2] + k]; + } + + // Number of elements. + __device__ __forceinline__ long long size() const { + return shape[0] * shape[1] * shape[2]; + } + + __host__ __device__ operator Array3D() const { + return Array3D{data, + {shape[0], shape[1], shape[2]}, + {shape[1] * shape[2], shape[2], 1}}; + } +}; + +template +struct CArray4D { + T* data; + long long shape[4]; + + __device__ __forceinline__ T& operator()(long long i, long long j, + long long k, long long l) const { + CUNUMPY_CHECK_INDEX(i, 0, shape[0]); + CUNUMPY_CHECK_INDEX(j, 1, shape[1]); + CUNUMPY_CHECK_INDEX(k, 2, shape[2]); + CUNUMPY_CHECK_INDEX(l, 3, shape[3]); + return data[((i * shape[1] + j) * shape[2] + k) * shape[3] + l]; + } + + // Number of elements. + __device__ __forceinline__ long long size() const { + return shape[0] * shape[1] * shape[2] * shape[3]; + } + + __host__ __device__ operator Array4D() const { + return Array4D{data, + {shape[0], shape[1], shape[2], shape[3]}, + {shape[1] * shape[2] * shape[3], shape[2] * shape[3], + shape[3], 1}}; + } +}; + +// The layouts the Python side packs: pointer, shape (and strides), 8-byte +// aligned. static_assert(sizeof(Array1D) == 24, "unexpected Array1D layout"); static_assert(sizeof(Array2D) == 40, "unexpected Array2D layout"); static_assert(sizeof(Array3D) == 56, "unexpected Array3D layout"); static_assert(sizeof(Array4D) == 72, "unexpected Array4D layout"); static_assert(sizeof(Array1D) == 24, "unexpected Array1D layout"); static_assert(alignof(Array2D) == 8, "unexpected Array2D alignment"); +static_assert(sizeof(CArray1D) == 16, "unexpected CArray1D layout"); +static_assert(sizeof(CArray2D) == 24, "unexpected CArray2D layout"); +static_assert(sizeof(CArray3D) == 32, "unexpected CArray3D layout"); +static_assert(sizeof(CArray4D) == 40, "unexpected CArray4D layout"); +static_assert(alignof(CArray2D) == 8, "unexpected CArray2D alignment"); #endif // CUNUMPY_ARRAY_VIEW_CUH diff --git a/tests/unit/test_cuda_kernel.py b/tests/unit/test_cuda_kernel.py index 1309007..582164e 100644 --- a/tests/unit/test_cuda_kernel.py +++ b/tests/unit/test_cuda_kernel.py @@ -756,6 +756,129 @@ def test_view_parameter_errors(): ) +SCALE_COLUMN_CONTIGUOUS = r""" +#include "cunumpy/array_view.cuh" +#include +extern "C" __global__ +void scale_column(CArray2D a, long long column, double factor) { + CUNUMPY_THREAD_1D(i, a.shape[0]); + a(i, column) *= factor; +} +""" + + +def test_parse_contiguous_view_parameters(): + source = ( + "__global__ void f(CArray1D a, const CArray3D< float > b, " + "Array2D c, CArray4D> d) {}" + ) + params = parse_cuda_signature(source, "f") + assert [(p.ctype, p.view_ndim, p.contiguous, p.dtype) for p in params] == [ + ("CArray1D", 1, True, np.dtype(np.float64)), + ("CArray3D", 3, True, np.dtype(np.float32)), + ("Array2D", 2, False, np.dtype(np.float64)), + ("CArray4D>", 4, True, np.dtype(np.complex128)), + ] + with pytest.raises(ValueError, match="array views"): + parse_cuda_signature("__global__ void f(CArray2D* a) {}", "f") + + +def test_contiguous_view_packs_pointer_and_shape_only(): + kernel = CudaKernel(SCALE_COLUMN_CONTIGUOUS, "scale_column") + a = FakeDeviceArray(np.float64, ptr=0xF00, shape=(4, 3)) + packed, _, _ = kernel.prepare_args(a, 1, 2.0) + assert packed.dtype.names == ("data", "shape") + assert packed.dtype.itemsize == 8 + 2 * 8 # data, shape[2] + assert packed["data"] == 0xF00 + assert packed["shape"].tolist() == [4, 3] + + +def test_contiguous_view_rejects_strided_arrays_instead_of_copying(): + kernel = CudaKernel(SCALE_COLUMN_CONTIGUOUS, "scale_column") + every_second_row = FakeDeviceArray(np.float64, shape=(4, 3), strides=(48, 8)) + with pytest.raises(TypeError, match=r"CArray2D a\) must be C-contiguous"): + kernel.prepare_args(every_second_row, 1, 2.0) + with pytest.raises(TypeError, match="must be a 2D array, got 1D"): + kernel.prepare_args(FakeDeviceArray(np.float64, shape=(3,)), 1, 2.0) + + +def test_struct_with_contiguous_view_fields(): + struct = CudaStruct( + "Markers", + [("markers", "CArray2D"), ("valid", "Array1D"), ("n", "int")], + ) + assert [f.contiguous for f in struct.fields] == [True, False, False] + assert struct.has_views + assert " CArray2D markers;" in struct.declaration + # 24 (CArray2D) + 24 (Array1D) + 4 (int) + 4 padding + assert struct.dtype.itemsize == 56 + value = struct( + markers=FakeDeviceArray(np.float64, ptr=0xA0, shape=(10, 7)), + valid=FakeDeviceArray(np.bool_, shape=(10,)), + n=10, + ) + assert value.packed["markers"]["shape"].tolist() == [10, 7] + assert value.packed["markers"].dtype.names == ("data", "shape") + with pytest.raises(TypeError, match="must be C-contiguous"): + struct( + markers=FakeDeviceArray(np.float64, shape=(10, 3), strides=(56, 8)), + valid=FakeDeviceArray(np.bool_, shape=(10,)), + n=10, + ) + + +def test_from_signature_contiguous(): + class Args: + def __init__(self, markers: "float[:, :]", valid: "bool[:]", n: "int"): ... + + every = CudaStruct.from_signature(Args.__init__, "A", contiguous=True) + assert [f.ctype for f in every.fields] == [ + "CArray2D", + "CArray1D", + "long long", + ] + some = CudaStruct.from_signature(Args.__init__, "A", contiguous=["markers"]) + assert [f.ctype for f in some.fields][:2] == ["CArray2D", "Array1D"] + with pytest.raises(ValueError, match=r"\['n'\] that are not array fields"): + CudaStruct.from_signature(Args.__init__, "A", contiguous=["n"]) + + source = ( + "class MarkerArguments:\n" + " def __init__(self, mks: 'float[:, :]', n: 'int'):\n" + " self.markers = mks\n" + ) + struct = CudaStruct.from_pyccel_class( + source, "MarkerArguments", contiguous=["markers"] + ) + assert [(f.name, f.ctype) for f in struct.fields] == [ + ("markers", "CArray2D"), + ("n", "long long"), + ] + + +def test_scale_column_of_contiguous_view_on_gpu(): + _skip_without_cupy() + import cupy as cp + + a = cp.arange(12, dtype=cp.float64).reshape(4, 3) + expected = a.get() + kernel = CudaKernel(SCALE_COLUMN_CONTIGUOUS, "scale_column", block_size=2) + kernel(a, 1, 10.0, n_threads=4) + expected[:, 1] *= 10.0 + assert np.array_equal(a.get(), expected) + with pytest.raises(TypeError, match="must be C-contiguous"): + kernel(a[::2], 1, 10.0, n_threads=2) + + +def test_contiguous_view_layout_on_gpu(): + _skip_without_cupy() + struct = CudaStruct( + "ContiguousLayout", + [("a", "CArray3D"), ("n", "int"), ("b", "Array2D")], + ) + struct.verify_layout() + + MARKERS = CudaStruct( "Markers", [("markers", "Array2D"), ("valid", "Array1D"), ("n", "int")], diff --git a/tests/unit/test_emulation.py b/tests/unit/test_emulation.py index 1da8a7d..cd16a75 100644 --- a/tests/unit/test_emulation.py +++ b/tests/unit/test_emulation.py @@ -69,6 +69,37 @@ def test_strided_view_written_back_into_the_callers_array(): np.testing.assert_array_equal(markers, expected) +CONTIGUOUS = r""" +#include "cunumpy/array_view.cuh" +#include +__device__ double sum_row(Array2D a, long long i) { + double total = 0.0; + for (long long j = 0; j < a.shape[1]; ++j) total += a(i, j); + return total; +} +extern "C" __global__ +void row_sums(CArray2D a, CArray3D b, CArray1D out) { + CUNUMPY_GRID_STRIDE_1D(i, a.shape[0]) { + out(i) = sum_row(a, i) + b(i, 1, 2); // CArray2D converts to Array2D + } +} +""" + + +def test_contiguous_views(): + a = np.arange(12.0).reshape(4, 3) + b = np.arange(4 * 2 * 3.0).reshape(4, 2, 3) + out = np.zeros(4) + kernel = CudaKernel(CONTIGUOUS, "row_sums") + emulate_cuda_kernel(kernel, a, b, out, grid=1, block=2) + np.testing.assert_array_equal(out, a.sum(axis=1) + b[:, 1, 2]) + + with pytest.raises(TypeError, match="must be C-contiguous"): + emulate_cuda_kernel( + kernel, a[:, :2].copy()[::2], b[::2], out[:2], grid=1, block=2 + ) + + TEMPLATE = r""" template __global__ void power(T* x, int n) {