1

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デバイスカーネルを使用する正しい方法を教えていただけますか?

4

3 に答える 3

2

私はこの種のミキシングを 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
于 2015-07-31T17:20:10.817 に答える
1

-Wc コマンド ライン オプションを使用して、生成された ptx ファイルを CUDA リンク ラインに追加できます。これを行う方法を文書化するためにバグを開きました。

nvcc hello.cu -ptx -arch=sm_35
cc -hacc -hlist=a mainacc.c -Wc,hello.ptx
于 2015-08-13T14:42:39.937 に答える
1

1 つの提案は、サブルーチンのホスト バージョンとデバイス バージョンの両方を提供し、"bind" 句を使用して、計算領域から呼び出すバージョンを示すことです。これにより、ホスト コードとの移植性を維持できます。

例えば：

% 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
于 2015-08-20T18:50:48.267 に答える