Блог

У меня есть ядро 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, 1);
  } else {
    blocks_per_grid = dim3(1, 1, 1);
    threads_per_block = dim3(length, 1, 1);
  }

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

  cudaDeviceSynchronize();

  return success();
}
  

Я могу добавить это в .so файл и загрузить его из Python с помощью ctypes . После этого я пытаюсь использовать его из Python,

Это эквивалент Python для ERROR структуры, возвращаемой в приведенном выше блоке кода —

 class Error(ctypes.Structure):
    _fields_ = [
        ("str", ctypes.POINTER(ctypes.c_char)),
        ("identity", ctypes.c_int64),
        ("attempt", ctypes.c_int64),
        ("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, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123, 123], dtype=cupy.in64)
tonumx = ctypes.cast(tonum.data.ptr, ctypes.POINTER(ctypes.c_int64))
fromstarts = cupy.array([2, 0, 2, 1, 2, 1, 2, 1, 2, 1, 2, 1], dtype=cupy.int32)
fromstarts = ctypes.cast(fromstarts.data.ptr, ctypes.POINTER(ctypes.c_int32))
fromstops = cupy.array([3, 2, 4, 5, 3, 4, 2, 5, 3, 4, 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.POINTER(ctypes.c_int32), ctypes.c_int64)
ret_pass = funcC(tonumx, fromstarts, fromstops, length)
  

но когда я печатаю tonum —

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

но значения должны быть — [1, 2, 2] (в зависимости от того, как cuda_ListArray_num работает)

Что я мог делать неправильно? Я думаю, что, возможно, я совершаю ошибку в том, как я передаю cupy указатель в ядро cuda.

Вы должны изменить код python на

 fromstarts = cupy.array([2, 0, 2, 1, 2, 1, 2, 1, 2, 1, 2, 1], dtype=cupy.int32)
fromstarts_ctypes = ctypes.cast(fromstarts.data.ptr, ctypes.POINTER(ctypes.c_int32))
fromstops = cupy.array([4, 2, 4, 5, 3, 4, 2, 5, 3, 4, 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.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.