OpenACC 中的嵌套指令

Question

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

我的代码就是这么简单。当我尝试编译以下代码时，编译器 returns 这些消息

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(...);

}

我当然只尝试过 routine (seq,worker,gang) 而没有内部 parallel loop 指令。已编译，但未激活动态并行。

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

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

Answer 1

How am I supposed to use dynamic parallelism in OpenACC?

尽管嵌套区域（可能会使用动态并行性）是 a new feature in the OpenACC 2.0 specification，但我认为它尚未在 PGI 15.7 中实现。 PGI 15.7 代表 OpenACC 2.0 规范的部分实施。

此限制记录在 PGI 15.7 发行说明中，该发行说明应随 PGI 15.7 编译器 (pgirn157.pdf) 一起提供，位于第 2.7 节（这些发行说明当前可用 here）：

OpenACC 2.0 Missing Features

‣ The declare link directive for global data is not implemented.

‣ Nested parallelism (parallel and kernels constructs within a parallel or kernels region) is not implemented.

根据评论，对 #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 并行已在内（例程）循环执行。

此方法不依赖于动态并行性（相反，它 relies on a partitioning of parallelism between worker at the routine level and gang at the caller level）并且不会调用动态并行性。

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

不使用（也不支持）

Answer 2

尝试用#pragma acc loop 替换“#pragma acc parallel loop”

OpenACC 中的嵌套指令

Nested Directives in OpenACC

cuda

gpu

nvidia

openacc

dynamic-parallelism