OpenACC中的嵌套指令

Posted 2021-04-21

我正在尝试使用OpenACC的嵌套功能来激活我的gpu卡的动态并行性。我有特斯拉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中使用动态并行？

答案

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

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

这一限制记录在PGI 15.7发行说明中，该说明应随附于第2.7节中的PGI 15.7编译器（pgirn157.pdf）（这些发行说明目前可用here）：

OpenACC 2.0缺少功能

‣未实现全局数据的declare link指令。

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

基于这些评论，对#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
", 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中没有使用（并且不受支持）。

另一答案

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