c - 使用 OpenACC 并行化嵌套循环

Question

我编写了一个涉及四个嵌套 for 循环的串行方法 - 我想使用 OpenACC 并行化此方法(这是我第一次尝试使用它，我对所有指令都不是很熟悉)。

我尝试了以下操作，但看到以下错误:调用 cuStreamSynchronize 返回错误 700:内核执行期间的非法地址

我在下面粘贴了我的方法的简化伪代码版本，我非常感谢您帮助找出并行化这四个嵌套循环结构的最佳方法。

// a, b, and c are input arguments to this method
#pragma acc parallel
for(int j = 0; j < a; j++){
    for(int i = 0; i < b; i++){

        // computing mins and maxs based on formulas with i, j, a, b, and c
        int minX = ...
        int maxX = ...
        int minY = ...
        int maxY = ...

        double count = (maxX - minX + 1)*(maxY - minY + 1);
        int sum1 = 0;
        int sum2 = 0;
        int sum3 = 0;

        #pragma acc loop
        for (int y = minY; y < maxY; y++) {
          for (int x = minX; x < maxX; x++) {
            #pragma acc routine(function_call_name) seq
            sum1 += // some function call;
            sum2 += // some function call;
            sum3 += // some function call;
          }
        }
        int result1 = (int)(sum1/count);
        int result2 = (int)(sum2/count);
        int result3 = (int)(sum3/count);

        #pragma acc routine(function_call_name) seq
        // calling some function call to store result1, result2, result3 in the output
    }
}

Answer 1

“非法地址”意味着您的程序正在访问 GPU 上的错误地址。通常，这是由于越界访问、访问设备上的主机地址、使用具有动态数据成员的聚合数据结构而不是“附加”成员(即在父结构中设置设备指针)引起的。不太常见的情况是堆或堆栈溢出。

您如何管理数据？代码中其他地方的数据区域？

如果使用 PGI，请首先尝试针对多核 CPU (-ta=多核)，这样您就无需担心数据移动。一旦并行区域开始工作，您就可以返回使用 GPU 并处理数据移动。我建议您首先使用 CUDA 统一内存 (-ta=tesla: Managed)，以便 CUDA 驱动程序为您处理数据移动(仅限动态数据)。然后，一旦该方法起作用，请尝试添加数据区域以手动管理数据。

我看到的其他东西:

并行构造需要在外部循环上使用循环指令。

#pragma acc parallel loop
for(int j = 0; j < a; j++){
    for(int i = 0; i < b; i++){

您可以考虑根据循环行程计数折叠循环:

#pragma acc parallel loop collapse(2)
for(int j = 0; j < a; j++){
    for(int i = 0; i < b; i++){

此外，“例程”指令应装饰例程的原型(prototype)或定义，但不应在计算区域中使用。

如果您在设备例程中使用任何全局变量，请务必将其放入“声明”指令中，以便在设备上创建数据的全局副本。

如果您使用 PGI，请向编译器添加“-Minfo=accel”选项。这将为您提供有关编译器如何并行化代码的编译器反馈消息。

如果您不使用数据指令，编译器将需要隐式复制数据。这些消息将告诉您正在复制哪些数组以及复制的大小。

如果您在理解反馈消息时遇到困难，请发布您的编译输出，我将帮助您完成它们。