1

我正在尝试使用 OpenACC 的嵌套功能来激活我的 gpu 卡的动态并行性。我有 Tesla 40c，我的 OpenACC 编译器是 PGI 版本 15.7。

我的代码很简单。当我尝试编译以下代码时，编译器会返回这些消息

PGCC-S-0155-Illegal context for pragma: acc  parallel loop (test.cpp: 158)
PGCC/x86 Linux 15.7-0: compilation completed with severe errors

我的代码结构：

#pragma acc parallel loop
for( i = 0; i < N; i++ )
{
  // << computation >>

  int ss = A[tid].start;
  int ee = A[tid].end;

  #pragma acc parallel loop
  for(j = ss; j< ( ee + ss); j++)
  {
    // << computation >>
  }

我还尝试更改我的代码以使用例程指令。但我无法再次编译

#pragma acc routine workers
foo(...)
{

  #pragma acc parallel loop
  for(j = ss; j< ( ee + ss); j++)
  {
    // << computation >>
  }
}

#pragma acc parallel loop
for( i = 0; i < N; i++ )
{
  // << computation >>

  int ss = A[tid].start;
  int ee = A[tid].end;

  foo(...);

}

我当然只尝试过没有内部并行循环指令的例程（seq，worker，gang） 。它一直是编译器，但尚未激活动态并行性。

    37, Generating acc routine worker
         Generating Tesla code
         42, #pragma acc loop vector, worker /* threadIdx.x threadIdx.y */
         Loop is parallelizable

我应该如何在 OpenACC 中使用动态并行性？

4

2 回答 2

3

我应该如何在 OpenACC 中使用动态并行性？

尽管嵌套区域（可能会使用动态并行）是OpenACC 2.0 规范中的一项新功能，但我认为它尚未在 PGI 15.7 中实现。PGI 15.7 代表 OpenACC 2.0 规范的部分实现。

此限制记录在 PGI 15.7 发行说明中，该说明应随您的 PGI 15.7 编译器 (pgirn157.pdf) 在第 2.7 节中提供（这些发行说明目前可在此处获得）：

OpenACC 2.0 缺少的功能

‣ 未实现全局数据的声明链接指令。

‣ 未实现嵌套并行性（并行或内核区域内的并行和内核构造）。

根据评论，有一些关于 的问题#pragma acc routine worker，所以这里是一个完整的 PGI 15.7 示例：

$ cat t1.c
#include <stdio.h>
#include <stdlib.h>
#define D1 4096
#define D2 4096
#define OFFS 2

#pragma acc routine worker
void my_set(int *d, int len, int val){
  int i;
  for (i = 0; i < len; i++) d[i] += val+OFFS;
}

int main(){


  int i,*data;
  data = (int *)malloc(D1*D2*sizeof(int));
  for (i = 0; i < D1*D2; i++) data[i] = 1;

#pragma acc kernels copy(data[0:D1*D2])
  for (i = 0; i < D1; i++)
    my_set(data+(i*D2), D2, 1);

  printf("%d\n", data[0]);

  return 0;
}
$ pgcc -acc -ta=tesla -Minfo=accel t1.c -o t1
my_set:
      8, Generating acc routine worker
         Generating Tesla code
         10, #pragma acc loop vector, worker /* threadIdx.x threadIdx.y */
         Loop is parallelizable
main:
     20, Generating copy(data[:16777216])
     21, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         21, #pragma acc loop gang /* blockIdx.x */
$ ./t1
4
$

请注意，gang 并行性已在外部循环中执行，worker 并行性已在内部（例程）循环中执行。

这种方法不依赖于动态并行（相反，它依赖于例程级别的worker和调用者级别的gang之间的并行划分）并且不会调用动态并行。

PGI 15.7 目前不支持本机使用动态并行 (CDP) 。应该可以调用（即与）其他使用 OpenACC 代码中的 CDP 的函数（例如，CUDA 或库），但目前，在 PGI 15.7 中不使用（也不支持）本机

于 2015-08-12T15:11:14.300 回答
0

尝试用#pragma acc loop 替换“#pragma acc 并行循环”

于 2018-03-04T03:52:48.077 回答