如何调用具有双指针作为参数的内核？

Posted 2023-02-16

【中文标题】如何调用具有双指针作为参数的内核？

【英文标题】：How to invoke a kernel which has a double pointer as argument?

【发布时间】：2021-06-17 21:34:20

【问题描述】：

我正在尝试使用 CuPy 以双指针作为参数调用自定义内核。

给定一个大小为 n 的 d 维点的数据集，我为每个维度分配一个 n 个浮点数的数组。 然后我将所有指向这些数组的指针放入另一个数组中。这就是我处理双指针的原因。

我的自定义内核具有如下签名： __global__ void myKernel(float **dataset, int n, int d, int* output, ...)

在本机 cuda 代码（通过 nvcc 编译）中，我的代码按预期工作。

我的 python 尝试调用内核看起来像：

import cupy as cp
from sklearn.datasets import make_blobs

X, y = make_blobs(n_samples=1_000, n_features=2, cluster_std=.5, center_box=(- 10.0, 10.0), shuffle=True, random_state=None)

with open('path', 'r') as f:
  code = f.read()

output = cp.empty(len(X), dtype=cp.int32)

kern = cp.RawKernel(code, 'myKernel')

d1 = cp.array(X[:, 0])
d2 = cp.array(X[:, 1])

blocks = ...
grid = ...
args = (cp.array([d1.data.ptr, d2.data.ptr]), 2, len(X), output, ...)
shared_mem = ...
kern((grid, 1, 1), (blocks, 1, 1), args=args, shared_mem=shared_mem)
cp.cuda.Stream.null.synchronize()

您可能会发现，我的问题在args = ... 行。 代码没有报错，但是输出向量中的数据肯定是错的。

我是否正确传递了数据？有没有更好的方法？

【参考方案1】：

您能尝试创建一个最小复制器吗？我认为传递双指针参数没有问题。

import cupy as cp
import sys

code = '''
extern "C" __global__ void myKernel(float **dataset, int n, int* output) 
for (int i = 0; i < n; ++i) 
    printf("[");
    for (int j = 0; j < 10; ++j) 
        printf("%f, ", dataset[i][j]);
    
    printf("]\\n");


'''

kern = cp.RawKernel(code, 'myKernel')

d0 = cp.arange(10, dtype=cp.float32)
d1 = cp.arange(10, dtype=cp.float32) * 2

dataset = cp.array([d0.data.ptr, d1.data.ptr])
output = cp.empty(10, dtype=cp.int32)

def run():
    print('--expected--')
    print(d0)
    print(d1)
    print('--actual--')
    sys.stdout.flush()
    kern((1, 1, 1), (1, 1, 1), args=(dataset, len(dataset), output))
    cp.cuda.Stream.null.synchronize()
    print('--end--')
    print()

run()

输出：

--expected--
[0. 1. 2. 3. 4. 5. 6. 7. 8. 9.]
[ 0.  2.  4.  6.  8. 10. 12. 14. 16. 18.]
--actual--
[0.000000, 1.000000, 2.000000, 3.000000, 4.000000, 5.000000, 6.000000, 7.000000, 8.000000, 9.000000, ]
[0.000000, 2.000000, 4.000000, 6.000000, 8.000000, 10.000000, 12.000000, 14.000000, 16.000000, 18.000000, ]
--end--

【讨论】：

感谢您的回答！我尝试在 Tesla T4 上的 Google Colab 上运行您的代码，并得到以下输出：--expected-- [0. 1. 2. 3. 4. 5. 6. 7. 8. 9.] [ 0. 2. 4. 6. 8. 10. 12. 14. 16. 18.] --actual-- --end--您使用什么设置来运行您的代码？

printf 在 CUDA 内核中无法被笔记本（Colab/Jupyter/etc）捕获。请尝试在本地运行。