NVCC-コンパイル時間のアーチに応じて異なるブロックサイズ

Question

最高のパフォーマンスを示すカーネルがあります 違う KeplerおよびFermiハードウェアで実行するときのブロックサイズ。コンパイル時に、現在のアーキテクチャのターゲットを確認し、 THREADS_PER_BLOCK マクロからi）カーネルを起動します。 ii）必要なブロックの数を決定します。 iii）カーネルに共有メモリサイズを静的に設定します。

以下は私がやろうとしていることを示しています。 GK104ハードウェアをターゲットにしていると仮定して、使用してください nvcc -arch=sm_30. 。これはまだ生じます THREADS_PER_BLOCK = 256 以来 __CUDA_ARCH__ ホストコードコンピレーションの場合は定義されていません。 （例えば、私は理解しています この答え, 、なぜこのように機能しないのか。）

#if __CUDA_ARCH__ >= 300
#define THREADS_PER_BLOCK 512
#else
#define THREADS_PER_BLOCK 256
#endif

__global__ void some_kernel(int* a, int* b) {
    __shared__ sm_data[THREADS_PER_BLOCK];
    // Do something.
}

int main(void) {
    // Initialize data.
    // Calculate blocks based on THREADS_PER_BLOCK, problem size and some max.
    some_kernel<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b)
    return 0;
}

実行時にデバイスのプロパティを確認し、動的共有メモリを使用することができますが、これをコンパイル時にハードコーディングできるかどうかを知りたいのです。 THREADS_PER_BLOCK それに基づいて。 NB：このコードのユーザーは自分でコンパイルします。 ほとんど 確かに1つのアーキテクチャにとって、これは不合理なオプションではありません。通過することに照らして不必要に思えます -arch= 国旗。

Answer 1

nvcc コンパイラはローカルで利用可能なGPUを検出しません。デフォルトでは常にSM 1.0をターゲットにします。それ以外の場合は、異なるシステムを構築するときに、非常に混乱する動作を導入する可能性があります。

使用可能なデバイスをコンパイルするには、ユーザーにSMバージョンの指定を依頼するか、ビルド時間中に検出コードを実行する必要があります。ハードウェア検出コードをランタイムに配置し、必要に応じてカーネルの起動を構成する方が簡単だと確信しています。