0

現在、ループとさまざまな CUBLAS 呼び出しを含むホスト関数があります。CC 3.5 デバイスにアクセスできるようになったので、動的並列処理を使用して、より効率的な単一のカーネルを作成できます。ただし、CC < 3.5 デバイスの古い機能を引き続きサポートしたいと考えています。いくつかの gencode を使用して、同じバイナリで複数のデバイスをサポートするようになりました。

-gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35

両方のアーキテクチャをサポートする単一のバイナリを作成し続けたいと思っていますが、ホスト コードでこれを切り替える方法が思い浮かびません。NVCC は、ホストの知る限り、コンパイルされたコード イメージを生成することはできません。

CC < 3.5 用にビルドしているユーザーは 3.5 の機能を使用してカーネルをビルドできないため、これは良くありません (そしてひどく醜いことでもあります)。

cudaGetDevice (&current_device);
cudaGetDeviceProperties (&current_device_properties, current_device);
if (current_device_properties.major < 3 && ... etc) {
  ...
}
else ...

__CUDACC__ または __CUDA_ARCH__ もここでは役に立ちません。

私の推測では、これは不可能であり、別のバイナリのコンパイルを開始し、プリプロセッサでアーキテクチャを切り替える必要があります。しかし、誰かが何かを考えることができれば、素晴らしいです。

4

1 に答える 1

1

それはあなたの目標が何であるかによって異なります。ここでは、2 つの異なるケースについて質問しているようです。

まず、ユーザーがCC 3.5 をサポートしない nvcc でコードをコンパイルする可能性があると思われる場合は、 CUDA_ARCHでプリプロセッサ チェックを使用して計算機能をテストし、サポートされていないコードをコンパイルしようとするのを防ぐ必要があります。

第 2 に、CC 3.5 とそれよりも低い機能の両方の実装を一緒に含めるようにコードをコンパイルする場合は、正しいホスト実装を選択するために既に述べたように、cudaGetDeviceProperties チェックを使用する必要があります。

これらの両方が同時に必要な場合は、このような実装を使用する必要があります。

cudaGetDevice (&current_device);
cudaGetDeviceProperties (&cdp, current_device);
if (cdp.major < 3 || (cdp.major >= 3 &&  cdp.minor < 5)) {
  //loop and CUBLAS
}else {
  kernel35<<<>>>();
}

同様に、カーネルは__CUDA_ARCH__ >= 350.

#if (__CUDA_ARCH__ >= 350)
__global__ void kernel35()
{
  ...
}
#else
__global__ void kernel35()
{
  //fake stub kernel to allow non 35 compatible nvcc to build the code
}
#endif

また、新しいカーネルがより効率的であることをテストしたと思いますが、反復回数が事前にわかっている場合、動的並列処理はほとんどの場合、CPU から正しく起動するよりも遅くなります。私のテストでは 40% もの差があったため、この切り替えを Kepler GPU に行う前に、パフォーマンスを徹底的にテストすることをお勧めします。

編集：より互換性があり、より安全なオプションは、2番目の部分をこのように表現することだと思います.

__global void kernel35(){
  #if (__CUDA_ARCH__ >=350 )
  ...
  #else
  //stub
  #endif
}
于 2013-06-14T15:16:06.293 に答える