cuda streams:在流执行后没有调用回调

Question

我遇到了cuda流回调函数的问题.问题是它在内核运行后不会执行.我将cuda流封装在C++类中,如下所示:

class Stream
{
public:
    Stream();
    void run(float *input, int points);
    ~Stream();
    static void CUDART_CB callback(cudaStream_t stream,
                                   cudaError_t status,
                                   void * user_data);

private:
    void callback_function();

    cudaStream_t s;

};

现在,我按照以下建议设置了回调函数,如下所示:

void CUDART_CB StreamWorkflow::callback(cudaStream_t stream,
                                    cudaError_t status,
                                    void * user_data)
{
    Stream* thiz = (Stream *)(user_data);
    thiz->callback_function();
}

构造函数设置流并附加回调

Stream::Stream()
{
    checkCudaErrors(cudaStreamCreate(&s));
    checkCudaErrors(cudaStreamAddCallback(s, Stream::callback, this, 0));
}

现在run方法调用了这个流上的cuda内核,我知道这个执行正常.它实现如下:

void Stream::run(float *f, int p)
{
    dim3 block(16, 16);
    dim3 grid((int)ceil(double(p) / 256.0));
    my_kernel(f, p, grid, block, s);
    checkCudaErrors(cudaDeviceSynchronize());
    getLastCudaError("kernel launch failed");
}

我的问题是只在构造函数执行时调用回调.因此,只要创建了对象,就会通过已注册的callback()执行callback_function().内核由主机多次执行,并且在完成后它永远不会执行回调函数.我可以看到内核已成功执行但回调永远不会到来.

每次内核运行时我是否必须注册回调(所以在执行run函数之前注册这个回调)或者我在这里做错了什么？

Answer 1

我想你可能对回调如何运作感到困惑.通常,当您向CUDA流发出CUDA操作时,在发出到该流的所有先前CUDA活动已完成时(以及在开始向该流发布的任何后续活动之前)执行该CUDA操作.

回调也不例外.如果要在内核执行后执行回调,则必须在发出内核后向该流发出回调.add回调函数的定义不是:

"任何时候内核在此流中完成,运行此回调"

add回调函数 的定义是:

"当CUDA流完成所有CUDA活动到目前为止,然后执行回调"

或者直接引用文档:

完成前一个流操作之前调用的函数

因此这个构造函数对我没有意义:

Stream::Stream()
{
    checkCudaErrors(cudaStreamCreate(&s));
    checkCudaErrors(cudaStreamAddCallback(s, Stream::callback, this, 0));
}

这个构造函数说:

"创建一个CUDA流"

"当此流的所有先前发布的活动完成后,运行此回调"

但是,当然,您尚未向该流发布任何活动,因此回调在流创建后立即(并且仅一次)运行.

"所以在执行run函数之前注册这个回调"

不,如果这是你想要的:

问题是它在内核运行后不会执行

然后在你的内核启动后移动你的add callback函数,我认为你更有可能看到你想要的活动.