OpenACC：从 OpenACC 并行循环调用 cuda __device__ 内核

Question

如果我在 hello.cu 文件中有简单的测试 cuda 内核：

extern "C" __device__ float radians( float f ){
    return f*3.14159265;
}

并在 mainacc.c 中测试 OpenACC 代码：

#include <stdio.h>
#include <stdlib.h>

#define N 10

#pragma acc routine seq
extern float radians( float );

int main() {

   int i;
   float *hptr, *dptr;
   hptr = (float *) calloc(N, sizeof(float));

   #pragma acc parallel loop copy(hptr[0:N])
   for(i=0; i<N; i++) {
       hptr[i] = radians(i*0.1f);
   }

   for( i=0; i< N; i++)
       printf("\n %dth value : %f", i, hptr[i]);
   return 0;
}

如果我尝试按如下方式编译此代码，则会出现链接时间错误：

nvcc hello.cu -c
cc -hacc -hlist=a mainacc.c hello.o
nvlink error   : Undefined reference to 'radians' in     '/tmp/pe_20271//app_cubin_20271.omainacc_1.o__sec.cubin'
cuda_link: nvlink fatal error

我尝试使用“--relocatable-device-code true”选项等的 nvcc，但没有成功。加载的模块是：

craype-accel-nvidia35
cudatoolkit/6.5
PrgEnv-cray/5.2.40

您能告诉我在 OpenACC 中使用 cuda device 内核的正确方法吗？

Answer 1

我已经能够使用 PGI 进行这种混合工作，但是我还不能生成一个可以使用 Cray 编译器的示例。这是一个适用于 PGI 的简单示例。

这是包含 CUDA 的文件。

// saxpy_cuda_device.cu
extern "C"
__device__
float saxpy_dev(float a, float x, float y)
{
  return a * x + y;
}

这是包含 OpenACC 的文件。

// openacc_cuda_device.cpp
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>

#pragma acc routine seq
extern "C" float saxpy_dev(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 parallel loop
    for( i = 0; i < n; i++ )
    {
      y[i] = saxpy_dev(2.0, x[i], y[i]);
    }
  }

  fprintf(stdout, "y[0] = %f\n",y[0]);
  return 0;
}

下面是编译命令。

$ make
nvcc  -rdc true -c saxpy_cuda_device.cu
pgc++ -fast -acc -ta=nvidia:rdc,cuda7.0 -c openacc_cuda_device.cpp
pgc++ -o openacc_cuda_device -fast -acc -ta=nvidia:rdc,cuda7.0  saxpy_cuda_device.o openacc_cuda_device.o -Mcuda

Answer 2

您可以使用 -Wc 命令行选项将生成的 ptx 文件添加到 CUDA 链接行。我已经打开了一个错误，以确保我们记录了如何执行此操作。

nvcc hello.cu -ptx -arch=sm_35
cc -hacc -hlist=a mainacc.c -Wc,hello.ptx

Answer 3

一个建议是提供子例程的主机和设备版本，然后使用“绑定”子句指示从计算区域调用哪个版本。这将允许您保持主机代码的可移植性。

例如：

% cat radians.cu
extern "C" __device__ float cuda_radians( float f ){
    return f*3.14159265;
}

extern "C" float radians( float f ){
    return f*3.14159265;
}
% cat test.c
#include <stdio.h>
#include <stdlib.h>

#define N 10

#pragma acc routine (radians) bind(cuda_radians) seq
extern float radians( float f);

int main() {

   int i;
   float *hptr, *dptr;
   hptr = (float *) calloc(N, sizeof(float));

   #pragma acc parallel loop copy(hptr[0:N])
   for(i=0; i<N; i++) {
       hptr[i] = radians(i*0.1f);
   }

   for( i=0; i< N; i++)
       printf("\n %dth value : %f", i, hptr[i]);
   return 0;
}
% nvcc -c radians.cu --relocatable-device-code true
% pgcc -acc -ta=tesla:cuda7.0 -Minfo=accel test.c radians.o -V15.7 -Mcuda
test.c:
main:
     15, Generating copy(hptr[:10])
         Accelerator kernel generated
         Generating Tesla code
         16, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
% a.out

 0th value : 0.000000
 1th value : 0.314159
 2th value : 0.628319
 3th value : 0.942478
 4th value : 1.256637
 5th value : 1.570796
 6th value : 1.884956
 7th value : 2.199115
 8th value : 2.513274
 9th value : 2.827434