vector_length 和 num_workers 如何在 OpenACC 例程中工作？

Question

当使用 OpenACC“#pragma acc routine worker”-routine 时，它包含向量（和 worker）级并行的多个循环，vector_length 和 num_workers 是如何工作的？

我尝试了一些代码（见下文）并偶然发现了一些东西：

设置这些循环的向量长度让我非常困惑。在比较运行次时，在外部 parallel 区域使用 vector_length(#) 子句似乎很奇怪。当我将矢量长度增加到巨大的数字时，例如4096，运行时间居然变小了。以我的理解，当矢量循环中只有 10 次迭代时，大量线程应该处于休眠状态。我是不是做错了什么？
我注意到输出奇怪地取决于 foo() 中的工人数量。如果它是 16 或更小，则输出是“正确的”。如果它是 32 甚至更大，则工作例程内的循环会以某种方式执行两次。我在这里错过了什么？

有人可以帮助我了解 OpenACC routine 子句吗？非常感谢。

示例代码如下：

#include <iostream>
#include <chrono>

class A{
public:
    int out;
    int* some_array;
    A(){
        some_array = new int[1000*100*10];
        for(int i = 0; i < 1000*100*10; ++i){
            some_array[i] = 1;
        }
        #pragma acc enter data copyin(this, some_array[0:1000*100*10])
    };
    
    ~A(){ 
        #pragma acc exit data delete(some_array, this)
        delete [] some_array;
    }
    
    #pragma acc routine worker
    void some_worker(int i){
        int private_out = 10;
        #pragma acc loop vector reduction(+: private_out)
        for(int j=0; j < 10; ++j){
            //do some stuff
            private_out -= some_array[j];
        }
        #pragma acc loop reduction(+: private_out) worker
        for(int j=0; j < 100; ++j){
            #pragma acc loop reduction(+: private_out) vector
            for(int k=0; k < 10; ++k){
                //do some other stuff
                private_out += some_array[k+j*10+i*10*100];
            }
        }
        #pragma acc atomic update
        out += private_out;
    }
    
    void foo(){
        #pragma acc data present(this, some_array[0:1000*100*10]) pcreate(out)
        {
            #pragma acc serial
            out=0;
            //#######################################################
            //# setting num_workers and vector_length produce weird #
            //# results and runtimes                                #
            //#######################################################
            #pragma acc parallel loop gang num_workers(64) vector_length(4096)
            for(int i=0; i < 1000; ++i){
                some_worker(i);
            }
            #pragma acc update host(out)
        }
    }
};

int main() {
    using namespace std::chrono;
    A a;
    auto start = high_resolution_clock::now();
    a.foo();
    auto stop = high_resolution_clock::now();
    std::cout << a.out << std::endl
              << "took " << duration_cast<microseconds>(stop - start).count() << "ms" << std::endl;
    //output for num_workers(16) vector_length(4096)
    //1000000
    //took 844ms
    //
    //output for num_workers(16) vector_length(2)
    //1000000
    //took 1145ms
    //
    //output for num_workers(32) vector_length(2)
    //1990000
    //took 1480ms
    //
    //output for num_workers(64) vector_length(1)
    //1990000
    //took 502ms
    //
    //output for num_workers(64) vector_length(4096)
    //1000000
    //took 853ms
    return 0;
}

机器规格：nvc++ 21.3-0 with OpenACC 2.7, Tesla K20c with cc35, NVIDIA-driver 470.103.01 with CUDA 11.4

编辑：

2 的附加信息：

我只是在 worker 中使用了一些 printfs 来查看中间结果。我将它们放在循环之间的隐式障碍中。我可以看到 private_out 的值最初是 10

到-10而不是循环之间的0和
到1990而不是1000。

在我看来这两个循环都被执行了两次。

为方便起见更多结果

给这个例子增加一些奇怪的地方：代码不能为 num_workers/vector_length 的某些组合编译。例如，将 num_workers 留在 64 并将 vector_length 设置为 2、4、8、16 甚至32（增加线程数超过 1024 个限制）。它给出了错误信息

ptxas error   : Entry function '_ZN1A14foo_298_gpu__1Ev' with max regcount of 32 calls function '_ZN1A11some_workerEi' with regcount of 41

然而，只需如上所述插入 printfs，它突然编译正常但运行s 进入运行time 错误：“调用 cuLaunchKernel 返回错误 1：无效价值”。

但最奇怪的是，它编译运行对于 64/64 没问题，但 returns 结果不正确。下面是使用 NV_ACC_TIME=1 设置的输出，但请注意，所有编译和运行ning 配置的输出几乎完全相同，block: [1x#-######] 部分除外。

Accelerator Kernel Timing data
/path/to/src/main.cpp
  _ZN1AC1Ev  NVIDIA  devicenum=0
    time(us): 665
    265: data region reached 1 time
        265: data copyin transfers: 3
             device time(us): total=665 max=650 min=4 avg=221
/path/to/src/main.cpp
  _ZN1AD1Ev  NVIDIA  devicenum=0
    time(us): 8
    269: data region reached 1 time
        269: data copyin transfers: 1
             device time(us): total=8 max=8 min=8 avg=8
/path/to/src/main.cpp
  _ZN1A3fooEv  NVIDIA  devicenum=0
    time(us): 1,243
    296: data region reached 2 times
    298: compute region reached 2 times
        298: kernel launched 2 times
            grid: [1-1000]  block: [1-32x1-24]
             device time(us): total=1,230 max=1,225 min=5 avg=615
            elapsed time(us): total=1,556 max=1,242 min=314 avg=778
    304: update directive reached 1 time
        304: data copyout transfers: 1
             device time(us): total=13 max=13 min=13 avg=13

Answer 1

worker 和 vector 的精确映射将取决于目标设备和实现。特别是当使用针对 NVIDIA GPU 的 NVHPC 时，“gang”映射到 CUDA 块，“worker”映射线程块的 y 维度，“vector”映射到 x-dimension。鉴于目标的限制，“num_workers”或“vector_length”中使用的值可能会减少。 CUDA 块最多可以包含 1024 个线程，因此“4096”值将减少到硬件允许的值。其次，为了支持设备例程中的向量缩减，最大 vector_length 可以是 32。换句话说，由于这些限制，您的“4096”值实际上是“32”。

注意查看设备上的最大线程块大小，运行“nvaccelinfo”实用程序并查找“每个块的最大线程数”和“最大块尺寸”字段。此外，设置环境变量“NV_ACC_TIME=1”将使运行time 产生一些基本的分析信息，包括运行期间使用的实际块数和线程块大小。

In my understanding, a huge amount of threads should lie dormant when there are only as many as 10 iterations in the vector loop.

CUDA 线程被分组为 32 个线程的“warp”，其中 warp 的所有线程同时执行相同的指令（也称为 SIMT 或单指令多线程）。因此，尽管只有 10 个线程在做有用的工作，但其余 12 个线程并未处于休眠状态。此外，它们仍然占用寄存器等资源，因此为循环次数较少的循环添加过多线程实际上可能会损害性能。

在这种情况下，将矢量长度设置为 1 最有可能是最好的情况，因为扭曲现在可以由 y-dimension 个线程组成。将其设置为 2，将在 x-dimension 中导致完整的 32 个线程扭曲，但只有 2 个线程在做有用的工作。

至于为什么有些组合会给出错误的结果，我没有调查。很少使用日常工作人员，尤其是在减少时，因此我们可能会遇到某种类型的代码生成问题，例如 off-by 在这些不规则的计划大小下，减少中的一个错误。我稍后会调查并确定是否需要提交问题报告。

对于 #2，您如何确定它获得运行两次？这只是基于运行时间吗？

vector_length 和 num_workers 如何在 OpenACC 例程中工作？

How does vector_length and num_workers work in an OpenACC routine?

c++

openacc

编辑：

2 的附加信息：

为方便起见更多结果