c - 将值分配给指针引用的结构的动态分配的结构成员数组时出现openacc错误
问题描述
我试图将 openacc 与指向包含动态分配成员的结构的指针结合起来。下面的代码失败了
线程失败:1 调用 cuStreamSynchronize 返回错误 700:内核执行期间地址非法
使用 nvc 编译时(“x86-64 Linux -tp haswell 上的 nvc 20.9-0 LLVM 64 位目标”)。据我所知,我正在遵循例如 OpenACC“入门”指南中建议的方法。但不知何故,指针可能不会粘在设备上(?)。有谁知道这里出了什么问题?
#include <stdlib.h>
#include <stdio.h>
typedef struct grid
{
int N;
double *X;
} grid;
void allocate(grid* g, int N)
{
g->N = N;
g->X = (double*) malloc(sizeof(double) * g->N);
#pragma acc enter data create(g[0:1])
#pragma acc enter data create(g->X[0:N])
}
void release(grid* g)
{
#pragma acc exit data delete(g->X[0:g->N])
#pragma acc exit data delete(g[0:1])
free(g->X);
}
void fill(grid * g)
{
int i;
#pragma acc parallel loop
for (i = 0; i < g->N; i++)
{
g->X[i] = 42; // the cuprit, commenting this removes the error too
}
}
int main()
{
grid g;
allocate(&g, 10);
fill(&g);
release(&g);
return 0;
}```
解决方案
从编译器反馈消息中,您会看到如下内容:
fill:
32, Accelerator restriction: size of the GPU copy of g is unknown
Generating Tesla code
32, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
32, Generating implicit copyin(g) [if not already present]
37, Generating update self(g->X[:g->N])
问题是编译器无法隐式复制具有动态数据成员的聚合类型,因此您需要添加“present(g)”以指示 g 已经是设备。
此外,您需要复制以在设备上获取 N 的值,并且无需在退出数据删除指令中包含数组形状。例如:
% cat test.c
#include <stdlib.h>
#include <stdio.h>
typedef struct grid
{
int N;
double *X;
} grid;
void allocate(grid* g, int N)
{
g->N = N;
g->X = (double*) malloc(sizeof(double) * g->N);
#pragma acc enter data copyin(g[0:1])
#pragma acc enter data create(g->X[0:N])
}
void release(grid* g)
{
#pragma acc exit data delete(g->X)
#pragma acc exit data delete(g)
free(g->X);
}
void fill(grid * g)
{
int i;
#pragma acc parallel loop present(g)
for (i = 0; i < g->N; i++)
{
g->X[i] = 42; // the cuprit, commenting this removes the error too
}
#pragma acc update self(g->X[:g->N])
for (i = 0; i < 4; i++)
{
printf("%d : %f \n",i,g->X[i]);
}
}
int main()
{
grid g;
allocate(&g, 10);
fill(&g);
release(&g);
return 0;
}
% nvc -acc test.c -Minfo=accel -V20.9 ; a.out
allocate:
17, Generating enter data copyin(g[:1])
Generating enter data create(g->X[:N])
release:
24, Generating exit data delete(g[:1],g->X[:1])
fill:
32, Generating present(g[:1])
Generating Tesla code
32, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
37, Generating update self(g->X[:g->N])
0 : 42.000000
1 : 42.000000
2 : 42.000000
3 : 42.000000
推荐阅读
- asp.net - Asp.Net Razor 自定义样式单选列表框
- java - 如何在多个类中抽象出具有不同返回类型的相同方法
- python - Python单例设计模式 - 静态方法与类方法方法
- react-native - 我是否必须制作单独的 MyComponent.android.js 文件才能使用 SafeArea
- android - Android studio 3.4.2 R8 混淆器不混淆类名,只混淆里面的java代码
- tex - 为什么我的 sage 命令的结果变成了问号?
- javascript - Office.js 如何将字符样式应用于选定的单词
- c# - 注销用户 OneDrive API
- reactjs - 如何删除 Chartjs 中的鬼线?
- java - 有没有办法在javascript中使用jspdf和autotable导出所有html表格分页数据