在 OpenACC 中修改循环变量（索引）

Question

我有一种情况，我需要多次重复循环的特定迭代。所以，在那个特定的迭代中，我将索引减少了一步，这样循环索引的下一个增量就没有区别了。

这种方法是我必须实施的方法，适用于多线程 OpenMP 代码。但是，它不适用于 OpenACC（对于 multicore 和 tesla 目标）。我收到以下错误：

Floating point exception (core dumped)

以下是这两种情况的代码：

#include <stdio.h>
#include <omp.h>
#include <unistd.h>

int main() {

    int x = 52;
    int count = 5;
    int i;

    omp_set_num_threads(6);

    #pragma omp parallel for
    for(i=0;i<100;i++) {
        if(i == x) {
            printf("%d\n", i);
            i--;
            count--;
            if(count == 0)
                x = 10000;
        }
    }


    int gpu_count = 0;
    count = 5;
    x = 52;

    #pragma acc parallel loop independent
    for(i=0;i<1000000;i++) {
        if(i == x) {
            #pragma acc atomic
            gpu_count++;

            i--;
            count--;
            if(count == 0)
                x = 2000000;
        }
    }
    printf("gpu_count: %d\n", gpu_count);


    return 0;
}

对于 OpenMP，我得到了正确的输出：

52
52
52
52
52

但是，对于 OpenACC，我得到了上述错误。

如果我注释第 35 行 (i--;)， 代码将正确执行并输出重复迭代次数（即 1）。

注意：我将 PGI 16.5 与 Geforce GTX 970 和 CUDA 7.5 一起使用。

我用 PGI 编译器编译如下：

pgcc -mp -acc -ta=multicore -g f1.c

所以，我的问题是：为什么我会看到这样的行为？我不能在 OpenACC 中更改循环索引变量吗？

Answer 1

您的 OpenMP 版本有误。您依赖于块大小大于“计数”的静态计划。如果您增加 OMP 线程的数量以使块大小小于计数，或者如果您更改调度以交错块（即“调度（静态，1）”），那么您将得到错误的答案。 “x”和“count”也有竞争条件。

请注意，OpenACC 调度更像 OpenMP “static,1”，因此向量可以跨工作程序（也称为 CUDA 扭曲）访问连续的内存块。所以你的算法在这里也行不通。

此外，通过使用“独立”子句（使用“并行循环”时隐含），您向编译器断言此循环不包含依赖项，或者用户将通过“原子”指令处理它们.但是，更改循环主体内的循环索引变量将创建循环依赖项，因为循环索引的值取决于上一次迭代是否更改了它的值。

编辑：下面是一个示例，它是您的代码的可并行版本。

% cat test2.c
#include <stdio.h>
#include <omp.h>
#include <unistd.h>

int main() {

    int x = 52;
    int count = 5;
    int i;
    int mycnt;

    #pragma omp parallel for schedule(static,1) private(mycnt)
    for(i=0;i<100;i++) {
        if(i == x) {
          mycnt = count;
          while(mycnt > 0) {
            printf("%d\n", i);
            mycnt--;
          }
        }
    }

#ifdef _OPENACC
    int gpu_count = 0;
    #pragma acc parallel loop reduction(+:gpu_count)
    for(i=0;i<1000000;i++) {
        if(i == x) {
          mycnt = count;
          while(mycnt > 0) {
            gpu_count++;
            mycnt--;
          }
        }
    }
    printf("gpu_count: %d\n", gpu_count);
#endif
    return 0;
}

% pgcc -fast -mp -acc test2.c  -Minfo=mp,acc
main:
     13, Parallel region activated
         Parallel loop activated with static cyclic schedule
     24, Barrier
         Parallel region terminated
     25, Accelerator kernel generated
         Generating Tesla code
         25, Generating reduction(+:gpu_count)
         26, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
         29, #pragma acc loop seq
     29, Loop carried scalar dependence for gpu_count at line 30
% a.out
52
52
52
52
52
gpu_count: 5