使用ctypes将杯形指针传递到CUDA内核

问题描述

template <typename T,typename C>
__global__
void cuda_ListArray_num(
  C *tonum,const T *fromstarts,const T *fromstops
) {
  int64_t block_id = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
  int64_t thread_id = block_id * blockDim.x + threadIdx.x;
  int64_t start = fromstarts[thread_id];
  int64_t stop = fromstops[thread_id];
  tonum[thread_id] = (C) (stop - start);
}

ERROR
awkward_ListArray32_num_64(
  int64_t* tonum,const int32_t* fromstarts,const int32_t* fromstops,int64_t length) {

  dim3 blocks_per_grid;
  dim3 threads_per_block;

  if (length > 1024) {
    blocks_per_grid = dim3(ceil((length) / 1024.0),1,1);
    threads_per_block = dim3(1024,1);
  } else {
    blocks_per_grid = dim3(1,1);
    threads_per_block = dim3(length,1);
  }

  cuda_ListArray_num<int32_t,int64_t><<<blocks_per_grid,threads_per_block>>>(
    tonum,fromstarts,fromstops);

  cudaDeviceSynchronize();

  return success();
}

我可以将其添加到.so文件中，然后使用ctypes从Python加载它。之后，我尝试从Python使用它，

这是在上面的代码块中返回的ERROR结构的Python等效项-

class Error(ctypes.Structure):
    _fields_ = [
        ("str",ctypes.POINTER(ctypes.c_char)),("identity",ctypes.c_int64),("attempt",("pass_through",ctypes.c_bool),]

这是我尝试从Python使用它的方式-

lib = ctypes.CDLL("cuda-kernels.so")

funcC = getattr(lib,'awkward_ListArray32_num_64')
funcC.restype = Error

tonum = cupy.array([123,123,123],dtype=cupy.in64)
tonumx = ctypes.cast(tonum.data.ptr,ctypes.POINTER(ctypes.c_int64))
fromstarts = cupy.array([2,2,1],dtype=cupy.int32)
fromstarts = ctypes.cast(fromstarts.data.ptr,ctypes.POINTER(ctypes.c_int32))
fromstops = cupy.array([3,4,5,3,6,11],dtype=cupy.int32)
fromstops = ctypes.cast(fromstops.data.ptr,ctypes.POINTER(ctypes.c_int32))
length = 3
funcC.argtypes = (ctypes.POINTER(ctypes.c_int64),ctypes.POINTER(ctypes.c_int32),ctypes.c_int64)
ret_pass = funcC(tonumx,fromstops,length)

但是当我打印tonum-

>>> tonum[:3]
array([0,0])

但值应为-[1,2]（基于cuda_ListArray_num的工作原理）

我可能做错了什么？我认为我如何将cupy指针传递到cuda内核可能是一个错误。

解决方法

您必须将python代码更改为

fromstarts = cupy.array([2,2,1,1],dtype=cupy.int32)
fromstarts_ctypes = ctypes.cast(fromstarts.data.ptr,ctypes.POINTER(ctypes.c_int32))
fromstops = cupy.array([4,4,5,3,6,11],dtype=cupy.int32)
fromstops_ctypes = ctypes.cast(fromstops.data.ptr,ctypes.POINTER(ctypes.c_int32))
length = 3
funcC.argtypes = (ctypes.POINTER(ctypes.c_int64),ctypes.POINTER(ctypes.c_int32),ctypes.c_int64)
ret_pass = funcC(tonumx,fromstarts_ctypes,fromstops_ctypes,length)

原因是CuPy数组是通过RAII管理的，因此，当您将fromstarts变量重新分配给另一个对象（ctypes指针）时，实际的数组将被销毁，其内存块将返回到CuPy的内存池中。此后，当您创建fromstops数组时，它将使用相同的内存块，覆盖fromstarts数组的内容，因为该数组不再存在，并共享相同的指针。>

然后，当您调用c代码时，fromstarts和fromstops实际上是相同的指针。您可以使用调试器或仅通过printf进行验证。

ctypes cuda cupy python