Bump version to 0.1.5, add CUDA interop support, and enhance documentation

Author: offerrall

.gitignore

@@ -7,4 +7,5 @@ site/
 *.obj
 *.lib
 *.exp
-.vscode/
+.vscode/
+.claude/

README.md

@@ -1,4 +1,4 @@
-# PyImageCUDA 0.1.4
+# PyImageCUDA 0.1.5
 
 [![PyPI version](https://img.shields.io/pypi/v/pyimagecuda.svg)](https://pypi.org/project/pyimagecuda/)
 [![Build Status](https://github.com/offerrall/pyimagecuda/actions/workflows/build.yml/badge.svg)](https://github.com/offerrall/pyimagecuda/actions)
@@ -39,6 +39,7 @@ with Image(1024, 1024) as bg:
 * ✅ **Studio Quality:** 32-bit floating-point precision (float32) to prevent color banding.
 * ✅ **Advanced Memory Control:** Reuse GPU buffers across operations and resize without reallocation—critical for video processing and batch workflows.
 * ✅ **OpenGL Integration:** Direct GPU-to-GPU display for real-time preview widgets.
+* ✅ **CUDA Interop (Zero-Copy):** Native `__cuda_array_interface__` v3 support. Share GPU buffers with CuPy, PyTorch, Numba and the rest of the CUDA Python ecosystem without copies.
 * ✅ **API Simplicity:** Intuitive, Pythonic API designed for ease of use.
 
 ## Use Cases
@@ -82,6 +83,7 @@ pip install pyimagecuda
 * [Image & Memory](https://offerrall.github.io/pyimagecuda/image/) (Buffer management)
 * [IO](https://offerrall.github.io/pyimagecuda/io/) (Loading and Saving)
 * [OpenGL Integration](https://offerrall.github.io/pyimagecuda/opengl/) (Real-time preview, zero-copy display)
+* [CUDA Interop](https://offerrall.github.io/pyimagecuda/image/#cuda-interop-zero-copy) (Zero-copy with CuPy, PyTorch, Numba)
 
 ### Operations
 * [Fill](https://offerrall.github.io/pyimagecuda/fill/) (Solid colors, Gradients, Checkerboard, Grid, Stripes, Dots, Circle, Ngon, Noise, Perlin)

docs/image.md

@@ -154,6 +154,121 @@ print(f"Current: {img.width}×{img.height}")  # 1280×720
 
 ---
 
+## CUDA Interop (Zero-Copy)
+
+PyImageCUDA images implement the [`__cuda_array_interface__`](https://numba.readthedocs.io/en/stable/cuda/cuda_array_interface.html) v3 protocol. This allows any library in the CUDA Python ecosystem (CuPy, PyTorch, Numba, RAPIDS, JAX with the CUDA backend, etc.) to read **and write** the image's GPU memory **without any copies**.
+
+The memory layout exposed is `(height, width, 4)` interleaved RGBA, row-major contiguous:
+
+- `Image` → `float32` (typestr `<f4`)
+- `ImageU8` → `uint8` (typestr `|u1`)
+
+!!! warning "Memory ownership"
+    The CUDA buffer is owned by the `Image` / `ImageU8` instance. The image **must stay alive** while any external array views it, otherwise the view becomes a dangling pointer. Do **not** call `image.free()` while a CuPy / Torch tensor still references it.
+
+---
+
+### Raw device pointer
+
+Use the `cuda_ptr` property to retrieve the raw CUDA device pointer as a Python `int`. Useful for custom kernels or low-level interop.
+```python
+from pyimagecuda import Image
+
+img = Image(1920, 1080)
+print(hex(img.cuda_ptr))  # e.g. 0x7f1234500000
+```
+
+---
+
+### CuPy
+
+The `to_cupy()` helper wraps a pyimagecuda image as a zero-copy `cupy.ndarray`. CuPy is **not** a dependency of pyimagecuda; install it separately:
+```bash
+pip install cupy-cuda12x   # or cupy-cuda11x for CUDA 11
+```
+```python
+from pyimagecuda import Image, Fill, to_cupy, download
+import cupy as cp
+
+img = Image(512, 512)
+Fill.color(img, (1, 0, 0, 1))
+
+# Zero-copy view — shares the same GPU memory
+arr = to_cupy(img)
+assert arr.data.ptr == img.cuda_ptr
+assert arr.shape == (512, 512, 4)
+assert arr.dtype == cp.float32
+
+# Modify via CuPy → changes are visible in the pyimagecuda image
+arr[:, :, 1] = 1.0  # add green channel
+
+# Download to verify
+pixels = download(img)
+```
+
+Because `__cuda_array_interface__` is honored, `cp.asarray(img)` works too:
+```python
+arr = cp.asarray(img)  # zero-copy
+```
+
+---
+
+### Custom CUDA kernels (CuPy RawKernel)
+
+You can run your own CUDA kernels directly on pyimagecuda buffers:
+```python
+from pyimagecuda import Image, Fill
+import cupy as cp
+
+img = Image(512, 512)
+Fill.color(img, (0.2, 0.4, 0.8, 1.0))
+
+invert_rgb = cp.RawKernel(r'''
+extern "C" __global__
+void invert_rgb(float4* data, int n) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= n) return;
+    float4 p = data[i];
+    data[i] = make_float4(1.0f - p.x, 1.0f - p.y, 1.0f - p.z, p.w);
+}
+''', 'invert_rgb')
+
+arr = cp.asarray(img)
+n = img.width * img.height
+invert_rgb((n // 256 + 1,), (256,), (arr, n))
+```
+
+---
+
+### PyTorch
+
+PyTorch consumes `__cuda_array_interface__` via `torch.as_tensor`:
+```python
+import torch
+from pyimagecuda import Image, Fill
+
+img = Image(512, 512)
+Fill.color(img, (1, 0, 0, 1))
+
+tensor = torch.as_tensor(img, device='cuda')  # zero-copy
+# tensor.shape == torch.Size([512, 512, 4])
+# tensor.dtype == torch.float32
+```
+
+---
+
+### Numba CUDA
+
+```python
+from numba import cuda
+from pyimagecuda import Image
+
+img = Image(512, 512)
+device_array = cuda.as_cuda_array(img)  # zero-copy
+```
+
+---
+
 ## Best Practices
 
 ### For Simple Scripts

pyimagecuda/__init__.py

@@ -1,7 +1,7 @@
 import ctypes
 import sys
 
-__version__ = "0.1.4"
+__version__ = "0.1.5"
 
 def _check_nvidia_driver():
     try:
@@ -47,6 +47,7 @@ def _check_nvidia_driver():
         from .transform import Transform
         from .text import Text
         from .gl_interop import GLResource
+        from .cuda_interop import to_cupy
         from .pyimagecuda_internal import cuda_sync # type: ignore
     except ImportError as e:
         print(f"Warning: Error importing Python wrappers: {e}")

pyimagecuda/cuda_interop.py

@@ -0,0 +1,30 @@
+from .image import ImageBase, ImageU8
+
+
+def to_cupy(image: ImageBase):
+    """
+    Wraps a pyimagecuda Image or ImageU8 as a zero-copy CuPy ndarray.
+
+    The returned ndarray shares GPU memory with the source image. The image
+    must remain alive while the array is in use; CuPy holds a reference to
+    the image via UnownedMemory's owner argument.
+
+    Docs & Examples: https://offerrall.github.io/pyimagecuda/image/
+    """
+    try:
+        import cupy
+    except ImportError as e:
+        raise ImportError(
+            "CuPy is required for to_cupy(). "
+            "Install it with: pip install cupy-cuda12x."
+        ) from e
+
+    is_u8 = isinstance(image, ImageU8)
+    bytes_per_pixel = 4 if is_u8 else 16
+    size_bytes = image.width * image.height * bytes_per_pixel
+
+    unowned_mem = cupy.cuda.UnownedMemory(image.cuda_ptr, size_bytes, image)
+    memptr = cupy.cuda.MemoryPointer(unowned_mem, 0)
+
+    dtype = cupy.uint8 if is_u8 else cupy.float32
+    return cupy.ndarray((image.height, image.width, 4), dtype=dtype, memptr=memptr)

pyimagecuda/image.py

@@ -1,4 +1,4 @@
-from .pyimagecuda_internal import create_buffer_f32, free_buffer, create_buffer_u8  #type: ignore
+from .pyimagecuda_internal import create_buffer_f32, free_buffer, create_buffer_u8, get_buffer_ptr  # type: ignore
 
 class Buffer:
 
@@ -86,6 +86,21 @@ def get_max_capacity(self) -> int:
         """Returns the total pixel capacity of the buffer."""
         return self._buffer.capacity_pixels
 
+    @property
+    def cuda_ptr(self) -> int:
+        return get_buffer_ptr(self._buffer._handle)
+
+    @property
+    def __cuda_array_interface__(self) -> dict:
+        return {
+            'shape': (self._height, self._width, 4),
+            'typestr': self._cai_typestr,
+            'data': (self.cuda_ptr, False),
+            'version': 3,
+            'strides': None,
+            'stream': 1,
+        }
+
     def free(self) -> None:
         self._buffer.free()
 
@@ -100,18 +115,22 @@ def __repr__(self) -> str:
         return f"{self.__class__.__name__}({self.width}×{self.height})"
 
 class Image(ImageBase):
-    
+
+    _cai_typestr = '<f4'
+
     def __init__(self, width: int, height: int):
         """
         Creates a floating-point image with the given width and height.
-        
+
         Docs & Examples: https://offerrall.github.io/pyimagecuda/image/#image-float32-precision
         """
         super().__init__(width, height, is_u8=False)
 
 
 class ImageU8(ImageBase):
-    
+
+    _cai_typestr = '|u1'
+
     def __init__(self, width: int, height: int):
         """
         Creates an 8-bit unsigned integer image with the given width and height.

pyproject.toml

@@ -4,7 +4,7 @@ build-backend = "scikit_build_core.build"
 
 [project]
 name = "pyimagecuda"
-version = "0.1.4"
+version = "0.1.5"
 description = "GPU-accelerated image processing library for Python"
 readme = "README.md"
 requires-python = ">=3.10"

src/buffer/buffer.cu

@@ -338,14 +338,43 @@ PyObject* py_copy_buffer(PyObject* self, PyObject* args) {
 
 PyObject* py_cuda_sync(PyObject* self, PyObject* args) {
     cudaError_t err = cudaDeviceSynchronize();
-    
+
     if (err != cudaSuccess) {
-        PyErr_Format(PyExc_RuntimeError, "CUDA synchronize failed: %s", 
+        PyErr_Format(PyExc_RuntimeError, "CUDA synchronize failed: %s",
                     cudaGetErrorString(err));
         return NULL;
     }
-    
+
     Py_RETURN_NONE;
 }
 
+PyObject* py_get_buffer_ptr(PyObject* self, PyObject* args) {
+    PyObject* capsule;
+
+    if (!PyArg_ParseTuple(args, "O", &capsule)) {
+        return NULL;
+    }
+
+    if (!PyCapsule_CheckExact(capsule)) {
+        PyErr_SetString(PyExc_TypeError, "Expected a Buffer capsule");
+        return NULL;
+    }
+
+    const char* name = PyCapsule_GetName(capsule);
+    if (name == NULL ||
+        (strcmp(name, BUFFER_TYPE_FLOAT32) != 0 &&
+         strcmp(name, BUFFER_TYPE_UINT8) != 0)) {
+        PyErr_SetString(PyExc_TypeError, "Expected a Buffer capsule");
+        return NULL;
+    }
+
+    BufferContext* ctx = (BufferContext*)PyCapsule_GetPointer(capsule, name);
+    if (ctx == NULL || ctx->freed || ctx->ptr == NULL) {
+        PyErr_SetString(PyExc_ValueError, "Buffer has been freed");
+        return NULL;
+    }
+
+    return PyLong_FromVoidPtr(ctx->ptr);
+}
+
 }

src/include.h

@@ -44,6 +44,7 @@ PyObject* py_create_buffer_u8(PyObject* self, PyObject* args);
 PyObject* py_cuda_sync(PyObject* self, PyObject* args);
 PyObject* py_download_from_buffer(PyObject* self, PyObject* args);
 PyObject* py_free_buffer(PyObject* self, PyObject* args);
+PyObject* py_get_buffer_ptr(PyObject* self, PyObject* args);
 PyObject* py_upload_to_buffer(PyObject* self, PyObject* args);
 
 // convert.cu

src/python_wrapper.cpp

@@ -41,6 +41,7 @@ static PyMethodDef all_methods[] = {
     {"cuda_sync", py_cuda_sync, METH_VARARGS, "Cuda Sync"},
     {"download_from_buffer", py_download_from_buffer, METH_VARARGS, "Download From Buffer"},
     {"free_buffer", py_free_buffer, METH_VARARGS, "Free Buffer"},
+    {"get_buffer_ptr", py_get_buffer_ptr, METH_VARARGS, "Get raw CUDA device pointer as int"},
     {"upload_to_buffer", py_upload_to_buffer, METH_VARARGS, "Upload To Buffer"},
 
     // convert.cu