简单 OpenACC 内核中向量子句的非法上下文

如何解决简单 OpenACC 内核中向量子句的非法上下文

我正在尝试编译一个简单的 OpenACC 基准：

void foo(const float * restrict a,int a_stride,float * restrict c,int c_stride) {
#pragma acc parallel copyin(a[0:a_stride*256]) copyout(c[0:c_stride*256])
#pragma acc loop vector(128)
  {
    for (int i = 0; i < 256; ++i) {
      float sum = 0;
      for (int j = 0; j < 256; ++j) {
        sum += *(a + a_stride * i + j);
      }
      *(c + c_stride * i) = sum;
    }
  }
}

使用 Nvidia HPC SDK 21.5 并遇到错误

$ nvc++ -S tmp.cc -Wall -Wextra -O2 -acc -acclibs -Minfo=all -g -gpu=cc80
NVC++-S-0155-Illegal context for gang(num:) or worker(num:) or vector(length:)  (tmp.cc: 7)
NVC++/x86-64 Linux 21.5-0: compilation completed with severe errors

知道是什么原因造成的吗？据我所知，vector(128) 的语法是合法的。

解决方法

将“向量（值）”与并行构造一起使用是非法的 OpenACC 语法。您需要在并行指令上使用“vector_length”子句来定义向量长度。原因是因为“并行”定义了要卸载的单个计算区域，因此该区域中的所有向量循环都需要具有相同的向量长度。

您只能将“vector(value)”与“kernels”构造一起使用，因为编译器随后可以将该区域拆分为多个内核，每个内核具有不同的向量长度。

选项 1：

% cat test.c
void foo(const float * restrict a,int a_stride,float * restrict c,int c_stride) {
#pragma acc parallel vector_length(128) copyin(a[0:a_stride*256]) copyout(c[0:c_stride*256])
#pragma acc loop vector
  {
    for (int i = 0; i < 256; ++i) {
      float sum = 0;
      for (int j = 0; j < 256; ++j) {
        sum += *(a + a_stride * i + j);
      }
      *(c + c_stride * i) = sum;
    }
  }
}
% nvc -acc -c test.c -Minfo=accel
foo:
      4,Generating copyout(c[:c_stride*256]) [if not already present]
         Generating copyin(a[:a_stride*256]) [if not already present]
         Generating Tesla code
          5,#pragma acc loop vector(128) /* threadIdx.x */
          7,#pragma acc loop seq
      5,Loop is parallelizable
      7,Loop is parallelizable

选项 2：

% cat test.c
void foo(const float * restrict a,int c_stride) {
#pragma acc kernels copyin(a[0:a_stride*256]) copyout(c[0:c_stride*256])
#pragma acc loop independent vector(128)
  {
    for (int i = 0; i < 256; ++i) {
      float sum = 0;
      for (int j = 0; j < 256; ++j) {
        sum += *(a + a_stride * i + j);
      }
      *(c + c_stride * i) = sum;
    }
  }
}
% nvc -acc -c test.c -Minfo=accel
foo:
      4,Generating copyout(c[:c_stride*256]) [if not already present]
         Generating copyin(a[:a_stride*256]) [if not already present]
      5,Loop is parallelizable
         Generating Tesla code
          5,#pragma acc loop gang,vector(128) /* blockIdx.x threadIdx.x */
          7,#pragma acc loop seq
      7,Loop is parallelizable