3

単精度と倍精度の両方のカードで実行できるようにしたいので、float と double 用にテンプレート化された CUDA アプリケーションを作成しようとしています。アプリケーションは、動的に割り当てられたグローバル、動的に割り当てられた共有、および定数メモリと静的グローバル メモリを使用します。

動的に割り当てられたグローバルおよび共有メモリ変数をテンプレート化する例を見てきました。そして、定数メモリは静的であるため、この投稿で述べられているように、テンプレート化は通常不可能であることを認識しています: cuda でテンプレート化された定数変数を定義します。

この絶え間ないメモリの問題に対する回避策を見つけることができませんでした。この問題に遭遇したのは私が初めてではないと確信しているため、驚くべきことです。現時点では、定数メモリを使用する場合、同じアプリケーションの 2 つのコピー (1 つは double 用、もう 1 つは float 用) を作成する必要があることに直面しているようです。そうでないことを願っています。

回避策として、テンプレート化され、定数メモリ変数宣言以外のすべてを実装する (仮想?) 基本クラスを作成することを検討しています。次に、主に定数変数の宣言を処理するだけの、ベースから継承する 2 つのクラス (float 用、double 用) を作成したいと思います。私の質問は、この戦略が機能するかどうか、または明らかな欠陥があるかどうかです。設計を実装する前に、それが機能しないことを確認するためだけに尋ねると思いました。この戦略が機能しない場合、少なくとも問題を軽減する実績のある他の戦略はありますか? それとも、float 用と double 用の 2 つのアプリケーションのコピーを作成するだけでよいのでしょうか?

4

1 に答える 1

2

この回答は歴史的な関心事、または CUDA ツールキット バージョン 6.5 以前を使用しているユーザー向けであることに注意してください。CUDA 7.0 以降、 のみをサポートする CUDA デバイスはサポートされていないfloatため、nvccCUDA コンパイラは、以下で説明する、自動的に に降格する機能を保持しなくなりdoubleましたfloat。

とのみに関心があると述べてfloatおり、をサポートしていないデバイスでdoubleのみ関心があると述べているため、これを処理するためにdouble から float へのコンパイラの自動降格を利用できるようです。.floatdoublenvcc

__constant__メモリを使用した例を次に示します。

$ cat t264.cu
#include <stdio.h>

#define DSIZE 64

__constant__ double my_const_data[DSIZE];

__global__ void my_kernel(double *data){
  data[1] = my_const_data[0];
  data[0] = sqrt(my_const_data[0]);
}

int main(){
  double my_data[DSIZE], h_data[DSIZE], *d_data;
  my_data[0] = 256.0;
  cudaMemcpyToSymbol(my_const_data, my_data, sizeof(double)*DSIZE);
  printf("hello\n");
  cudaMalloc((void **)&d_data, sizeof(double)*DSIZE);
  my_kernel<<<1,1>>>(d_data);
  cudaMemcpy(&h_data, d_data, sizeof(double)*DSIZE, cudaMemcpyDeviceToHost);
  printf("data = %lf\n", h_data[1]);
  printf("sqrt = %lf\n", h_data[0]);
  return 0;
}

$ nvcc -o t264 t264.cu
ptxas /tmp/tmpxft_00003228_00000000-5_t264.ptx, line 62; warning : Double is not supported. Demoting to float
$ ./t264
hello
data = 256.000000
sqrt = 16.000000
$
于 2013-10-21T14:21:04.170 に答える