如何确保两个流中的两个内核同时发送到 GPU 运行？

Posted 2023-03-23

【中文标题】如何确保两个流中的两个内核同时发送到 GPU 运行？

【英文标题】：How can I make sure two kernels in two streams are sent to the GPU at the same time to run?

【发布时间】：2022-01-08 02:36:37

【问题描述】：

我是 CUDA 的初学者。我正在使用 NVIDIA Geforce GTX 1070 和 CUDA 工具包 11.3 和 ubuntu 18.04。 如下代码所示，我使用两个 CPU 线程将两个内核以两个流的形式发送到一个 GPU。我希望这两个内核同时发送到 GPU。 有没有办法做到这一点？

或者至少比我做的更好。

提前谢谢你。

我的代码：

//Headers
pthread_cond_t cond;
pthread_mutex_t cond_mutex;
unsigned int waiting;
cudaStream_t streamZero, streamOne;  

//Kernel zero defined here
__global__ void kernelZero()...

//Kernel one defined here
__global__ void kernelOne()...

//This function is defined to synchronize two threads when sending kernels to the GPU.
void threadsSynchronize(void) 
    pthread_mutex_lock(&cond_mutex);
    if (++waiting == 2) 
        pthread_cond_broadcast(&cond);
     else 
        while (waiting != 2)
            pthread_cond_wait(&cond, &cond_mutex);
    
    pthread_mutex_unlock(&cond_mutex);



void *threadZero(void *_) 
    // ...
    threadsSynchronize();
    kernelZero<<<blocksPerGridZero, threadsPerBlockZero, 0, streamZero>>>();
    cudaStreamSynchronize(streamZero);
    // ...
    return NULL;



void *threadOne(void *_) 
    // ...
    threadsSynchronize();
    kernelOne<<<blocksPerGridOne, threadsPerBlockOne, 0, streamOne>>>();
    cudaStreamSynchronize(streamOne);
    // ...
    return NULL;



int main(void) 
    pthread_t zero, one;
    cudaStreamCreate(&streamZero);
    cudaStreamCreate(&streamOne); 
    // ...
    pthread_create(&zero, NULL, threadZero, NULL);
    pthread_create(&one, NULL, threadOne, NULL);
    // ...
    pthread_join(zero, NULL);
    pthread_join(one, NULL);
    cudaStreamDestroy(streamZero);  
    cudaStreamDestroy(streamOne);  
    return 0;

【参考方案1】：

实际上，在 GPU 上观察并发内核行为有许多要求，这些要求在 SO cuda 标签上的其他问题中有所涉及，因此我不打算介绍这些内容。

假设您的内核可以同时运行。

在这种情况下，无论你是否使用线程，你都不会做得比这更好：

cudaStream_t s1, s2;
cudaStreaCreate(&s1);
cudaStreamCreate(&s2);
kernel1<<<...,s1>>>(...);
kernel2<<<...,s2>>>(...);

如果您的内核有一个“长”的持续时间（比内核启动开销长得多，大约 5-50 微秒），那么它们似乎“几乎”在同一时间启动。通过切换到线程，您不会做得比这更好。据我所知，其原因尚未公布，所以我会简单地说，我自己的观察表明，内核启动到同一个 GPU 是由 CUDA 运行时序列化的，不知何故。你可以在各种论坛上找到这方面的轶事证据，如果你不相信我也没关系。对于我熟悉的 CPU 线程机制，也没有理由假设 CPU 线程同步执行。因此，没有理由假设一个线程系统会导致两个不同线程中的内核启动甚至被主机线程在同一时刻到达。

使用cudaLaunchKernel 进行内核启动可能会做得更好，而不是使用三字形启动语法：<<<...>>>，但确实没有文档支持这种说法。 YMMV。

请记住，GPU 作为吞吐量处理器正在尽其所能。没有明确的机制来确保同时启动内核，也不清楚您为什么需要它。