如何适当地结合 OpenMP 和 CUDA答案

【问题标题】：How to combine OpenMP with CUDA appropriately如何适当地结合 OpenMP 和 CUDA
【发布时间】：2012-09-18 20:15:40
【问题描述】：

我最近尝试将我的 CUDA 代码与 OpenMP 结合，但出现了一些问题。我的 CUDA-OpenMP 代码写成

    int main (void)
    {
       cudaGetDeviceCount(&gpuNum);

       //declare variables
       float *data_h; *data_d[gpuNum];
       data_h = (float*)malloc(Mem);
       ....
       ....
       omp_set_num_threads(gpuNum); 
       #pragma omp parallel
       {
         int cpuid = omp_get_thread_num();
         cudaSetDevice(cpuid);

         cudaMalloc((void**)&data_d[cpuid], Mem );

         cudaMemcpy( data_d[cpuid], data_h, Mem, cudaMemcpyHostToDevice);
         kernel<<< gpu_block, gpu_thread >>>();
         cudaMemcpy( data_h, data_d[cpuid], Mem, cudaMemcpyDeviceToHost);
       }
       printf("end of parallel\n");
       //post process
     }

问题是有时我运行这段代码时，一切都很顺利，但有时它会停止并且不会打印出“并行结束”这句话。这种情况是随机发生的，谁能告诉我这可能是什么原因？

【问题讨论】：

我不确定原因，但我的猜测是每个线程都将最终结果复制到“data_h”（内核完成后）。它可能涉及一些锁定。可以尝试为每个线程分配内存并复制最终结果。检查该程序是否挂起。
请检查所有 API 调用的返回值（请参阅stackoverflow.com/tags/cuda/info 了解提问技巧）。

标签： cuda openmp

【解决方案1】：

我想提供一些失败的可能性：

假设在并行区域中，当前两行执行时，活动线程与另一行切换，

#pragma omp parallel{
  int cpuid = omp_get_thread_num();
  cudaSetDevice(cpuid);
  ...
}

然后另一个线程将调用设置设备函数并更改所选设备。

虽然 memcopy 操作会阻塞内核调用，但不会。因此，如果在内核调用之后切换线程，当一个内核调用未完成时，将发出另一个内核调用。要执行并发内核，您需要使用“streams”。来看看

CUDA concurrent kernel execution with multiple kernels per stream

【讨论】：

这个推测不太正确。 cudaSetDevice() 是主机线程安全的。运行时 api 手册：“可以在任何时间从任何主机线程、任何设备进行此调用。”
据我所知，这里有一个竞争条件。一个线程将设备设置为 0，而另一个线程将其设置为 1，依此类推..
多个主机线程调用 cudaSetDevice 不会导致竞争条件。如果您不相信运行时 API 手册，请查看对应的驱动程序 API。它明确表示 cuCtxCreate、cuCtxPopCurrent 和 cuCtxPushCurrent 与调用 CPU 线程相关联，这意味着与上下文相关的变量是线程本地的。 OP 的问题更有可能是内核调用失败引起的。

【解决方案2】：

如果您想获得最佳性能，我建议您不要使用 OpenMP 来运行 CUDA。根据我的经验，我知道通过 OpenMP 创建线程并非没有成本，而且需要一些时间。在创建线程期间，您将能够在许多设备上运行更多的 gpu 内核。

如@phoad 所述，如果您的数据集是独立的，您可以使用流。你可以在网上找到很多例子。

另一种可能性是重新设计内核。一个内核可以完成相同的工作并返回数组作为结果。

【讨论】：