如何在CuPy中使用C++传递的CUDAarray指针创建ndarray？

我尝试在C++中分配CUDAarray（即纹理内存）并将指针传递给CuPy，希望直接将其视为ndarray使用。目前已经实现了通过cudaMalloc()分配线性内存，再用cupy.cuda.UnownedMemory()共享给CuPy的方案，但传递cudaMallocArray()分配的CUDAarray指针时，会触发cudaErrorIllegalAddress错误。

查阅CuPy外部内存文档未发现明确限制，但ndarray文档提到需要指向内存头部的指针，而CUDAarray并非直接的内存起始地址，而是GPU内部内存的引用，因此考虑需要在C++中通过cudaMemcpy2DFromArray()将其转成线性内存后再传递。

不过CuPy源码中的TODO注释让我产生疑问：该注释是指当前已支持无需包装器，还是未来计划让ndarray支持CUDAarray时再修改？

我也尝试过实例化cupy.cuda.texture.CUDAarray()，但无法修改其只读的ptr属性，传递纹理对象也会遇到同样的非法内存访问问题。

想请教：

• 最优实现方式是什么？
• cudaMemcpy2DFromArray()是否是唯一可行的方案？
• 该方案是否会存在性能损耗？

cudaArray_t cuArray = NULL; // 全局变量，DLL加载期间保持存活，Python进程退出时释放

extern "C"
uintptr_t __declspec(dllexport) __stdcall getCudaArray()
{
    if (cuArray != NULL) {
        return reinterpret_cast<uintptr_t>(cuArray);
    }

    const int height = 1024;
    const int width = 1024;

    // 分配并初始化主机数据
    float* h_data = (float*)std::malloc(sizeof(float) * width * height);
    for (int i = 0; i < height * width; ++i)
        h_data[i] = i + 0.1;

    // 在设备内存中分配CUDA array
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaMallocArray(&cuArray, &channelDesc, width, height);

    // 设置源数据的行间距（字节为单位）
    const size_t srcpitch = width * sizeof(float);

    // 将主机数据复制到设备内存的CUDA array中
    cudaMemcpy2DToArray(cuArray, 0, 0, h_data, srcpitch, width * sizeof(float), height, cudaMemcpyHostToDevice);

    // 返回CUDAarray的指针
    return reinterpret_cast<uintptr_t>(cuArray);
}

extern "C"
uint64_t __declspec(dllexport) __stdcall clearCudaArray()
{
    cudaFreeArray(cuArray);
    cuArray = NULL;
    return 0;
}

import ShTexDll
import cupy as cp

ptr = ShTexDll.getCudaArray()

mem = cp.cuda.UnownedMemory(ptr, 4*1024*1024, owner=None, device_id=0)
memptr = cp.cuda.MemoryPointer(mem, offset=0)

arr = cp.ndarray((1024,1024), dtype=cp.float32, memptr=memptr, strides=(4,1024*4))

print(arr)

ShTexDll.clearCudaArray()

Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
  File "cupy\_core\core.pyx", line 1741, in cupy._core.core._ndarray_base.__str__
  File "cupy\_core\core.pyx", line 1854, in cupy._core.core._ndarray_base.get
  File "cupy\cuda\memory.pyx", line 586, in cupy.cuda.memory.MemoryPointer.copy_to_host_async
  File "cupy_backends\cuda\api\runtime.pyx", line 606, in cupy_backends.cuda.api.runtime.memcpyAsync
  File "cupy_backends\cuda\api\runtime.pyx", line 146, in cupy_backends.cuda.api.runtime.check_status
cupy_backends.cuda.api.runtime.CUDARuntimeError: cudaErrorIllegalAddress: an illegal memory access was encountered

为什么直接传递CUDAarray指针会报错？ #

CUDAarray是GPU纹理内存的抽象句柄，并非线性内存的直接地址，它的内部布局是硬件优化的（比如可能是tiled格式），无法像cudaMalloc()分配的线性内存那样直接通过指针进行随机访问。CuPy的ndarray依赖连续的线性内存地址空间，直接将CUDAarray的句柄当作内存指针传入，必然会触发非法内存访问错误。

最优实现方案 #

目前cudaMemcpy2DFromArray()是最直接的可行方案，但并非唯一选择，具体可以根据使用场景选择：

方案一：转成线性内存后共享（通用场景） #

• 实现方式：在C++侧通过cudaMemcpy2DFromArray()将CUDAarray的数据复制到cudaMalloc()分配的线性内存中，再将线性内存的指针传递给CuPy，用UnownedMemory包装成ndarray。
• 性能损耗：会产生一次设备内的数据拷贝，对于1024x1024的float数组来说，拷贝开销很小（约4MB数据，耗时微秒级），如果不是频繁拷贝，对整体性能影响可以忽略。
• 优点：兼容性最好，CuPy可以直接对线性内存进行任意操作（读写、计算等）。

方案二：在CuPy中直接使用纹理对象（计算场景） #

• 实现方式：如果核心需求是利用CUDAarray的纹理加速特性（比如插值、缓存优化），可以在C++侧创建纹理引用，然后通过CuPy的纹理API（cupy.cuda.texture）在核函数中直接访问CUDAarray，无需拷贝到线性内存。
• 步骤：
  1. 在C++侧创建并绑定纹理到CUDAarray（使用cudaBindTextureToArray）。
  2. 将纹理引用的句柄传递给Python，在CuPy的自定义核函数中通过CUDA纹理API读取数据。
• 优点：避免拷贝，充分利用纹理内存的硬件优化，适合需要纹理采样的计算场景。
• 缺点：需要编写CuPy自定义核函数，无法直接将CUDAarray转为普通ndarray进行通用操作。

方案三：使用CuPy的cudaArray包装（暂不可行） #

你尝试的cupy.cuda.texture.CUDAarray()是CuPy对CUDAarray的封装，但它的ptr属性是只读的，且目前CuPy并未提供直接从外部CUDAarray句柄创建该对象的API，因此无法通过这种方式直接共享外部CUDAarray。源码中的TODO注释属于未来规划，当前版本（截至CuPy v12.x）尚未支持直接将外部CUDAarray转为ndarray。

总结 #

• 如果需要将CUDAarray的数据作为普通ndarray进行任意操作，转成线性内存是唯一通用方案，性能损耗在大多数场景下可以接受。
• 如果仅需在计算中利用纹理特性，直接使用纹理对象访问是最优选择，无需拷贝，性能最优。

内容的提问来源于stack exchange，提问作者Nicholas Masso