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

Question

我是 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;
}

Answer 1

实际上，在 GPU 上见证并发内核行为有许多要求，这些要求在 SO cuda 标签上的其他问题中已涵盖，因此我不打算涵盖这些内容。

让我们假设您的内核有可能同时运行。

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

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

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

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

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

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

How can I make sure two kernels in two streams are sent to the GPU at the same time to run?

synchronization

cuda

cuda-streams