控制从 OpenACC 计算区域进入 CUDA 函数的线程答案

【问题标题】：Controlling the threads that go into CUDA function from OpenACC compute region控制从 OpenACC 计算区域进入 CUDA 函数的线程
【发布时间】：2016-02-11 15:40:23
【问题描述】：

我正在从 OpenACC 计算区域调用一个 cuda 函数，并且我想指定应该进入 cuda 函数的线程数，但似乎我不知道如何控制它。

%main.cpp
..

#pragma acc routine vector
extern "C" void CUDA_KERNEL_FUNCTION(double *B, int ldb,const double *A, int lda);
..
#pragma acc parallel loop independent collapse(3) gang vector(128)
  for(int i0 = 0; i0 < size0 - 31; i0+= 32)
     for(int i1 = 0; i1 < size1 - 31; i1+= 32)
        for(int i2 = 0; i2 < size2; i2+= 1)
          CUDA_KERNEL_FUNCTION(B, ldb, A, lda);

..
..


%cuda_code.cu

extern "C" __device__ void CUDA_KERNEL_FUNCTION(double *B, int ldb,const double *A, int lda)
{

    Num_Threads_gpu = blockDim.x * blockDim.y* blockDim.z;

    //Num_Threads_gpu is always 32 
}

编译很好。但是无论我使用什么向量长度，进入cuda函数的线程数总是32。有什么方法可以指定吗？

我使用“cuda/7.0.28”和“pgi/15.10”

谢谢

【问题讨论】：

您的Num_Threads_gpu 正在计算块中的线程数，而不是“进入函数”的线程总数。
但是我已经指定了向量（128）。所以只有一个 128 个线程的向量应该进入，对吗？ ——

标签： cuda gpu nvcc openacc pgi

【解决方案1】：

尝试将 vector(128) 更改为 vector_length(128)。我认为 PGI 15.10 支持这两种语法，但以防万一……

如果这不起作用，您能否使用-Minfo=accel 发布编译器输出，以便我们可以看到编译器在做什么？

【讨论】：

太棒了！使用parallel loop 时，调整帮派/工人/向量的正确方法是在区域顶部使用num_gangs、num_workers 和vector_length。在 OpenACC 2.5 中，我们还可以像您展示的那样进行操作，但这仅在一个月前得到批准，因此编译器也可能需要一些时间才能采用该语法。