目前大家对于openacc的使用还是比较独立,也就是程序中要不就是使用openacc,要不就是使用cuda。但是结合openacc和cuda使用目前还是比较少。刚好看到几个好的案例,这就把例子搬过来。
/*openacc_c_main.c*/
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
extern void saxpy(int,float,float*,float*);
int main(int argc, char **argv)
{
float *x, *y, tmp;
int n = 1<<20, i;
x = (float*)malloc(n*sizeof(float));
y = (float*)malloc(n*sizeof(float));
#pragma acc data create(x[0:n]) copyout(y[0:n])
{
#pragma acc kernels
{
for( i = 0; i < n; i++)
{
x[i] = 1.0f;
y[i] = 0.0f;
}
}
#pragma acc host_data use_device(x,y)
{
saxpy(n, 2.0, x, y);
}
}
fprintf(stdout, "y[0] = %f\n",y[0]);
return 0;
}
/* saxpy_cuda.cu */
__global__
void saxpy_kernel(int n, float a, float *x, float *y)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if ( i < n )
y[i] += a * x[i];
}
extern "C" void saxpy(int n ,float a, float *x, float *y)
{
dim3 griddim, blockdim;
blockdim = dim3(128,1,1);
griddim = dim3(n/blockdim.x,1,1);
saxpy_kernel<<<griddim,blockdim>>>(n,a,x,y);
}
对于PGI编译,不能识别 .cu文件。所以对于这两个文件需要分开编译。编译测试平台:linux,k20c
对于saxpy_cuda.cu 文件:
nvcc -c saxpy_cuda.cu -arch=sm_35
对于openacc_c_main.c文件:
pgcc -c openacc_c_main.c -acc -Mcuda=cc3.5 -Minfo
最后把生成的openacc_c_main.o和saxpy_cuda.o文件一起链接
pgcc -o openacc_cuda openacc_c_main.o saxpy_cuda.o -Mcuda -acc