2

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

1 回答 1

2

实际上，在 GPU 上见证并发内核行为有许多要求，这些要求在 SO 标签上的其他问题中有所涉及cuda，因此我不打算讨论这个问题。

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

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

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

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

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

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

于 2021-12-01T14:23:45.023 回答