OpenACC与CUDA C交互(1)

目前大家对于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