简单 OpenACC 内核中向量子句的非法上下文
Posted
技术标签:
【中文标题】简单 OpenACC 内核中向量子句的非法上下文【英文标题】:Illegal context for vector clause in simple OpenACC kernel 【发布时间】:2021-09-13 01:13:58 【问题描述】:我正在尝试编译一个简单的 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)
的语法是合法的。
【问题讨论】:
【参考方案1】:在并行结构中使用“vector(value)”是非法的 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 a_stride, float * restrict c, 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
【讨论】:
非常感谢,马特。规范对此特别含糊:“当父计算构造是并行构造或孤立循环构造时,向量子句指定关联循环的迭代将在向量或 SIMD 模式下执行.. . 操作将使用为并行区域指定或选择的长度的向量执行”,因此也许可以使错误消息更清晰(即“不能为并行块内的单个循环指定向量长度”)?以上是关于简单 OpenACC 内核中向量子句的非法上下文的主要内容,如果未能解决你的问题,请参考以下文章