Fermi Compatibility Guide日本語訳
ソフトウェア要件
このFermi Compatibility Guide for CUDA Applicationsは、CUDAアプリケーションがFerimiアーキテクチャをベースにしたGPUで動作することを開発者が確認することを助けるアプリケーションノートです。このガイドは、すでにCUDA C/C++でのプログラミングを良く知っており、ソフトウェアアプリケーションがFermiと互換性があるかを確認したい開発者への案内を提供することを意図しています。
重要事項:
Fermiアーキテクチャの導入より前、すべてのNVIDIAのTelsaブランドの商品は、Teslaアーキテクチャをベースとしていました。この文書に対して、"Tesla" という単語はGPUアーキテクチャのみを差し、いかなる特定のNVIDIA商品をも差しません。以下では、Teslaは、compute capability 1.xのデバイスを差し、Fermiは、compute capabitily 2.0のデバイスを差します。
NVIDIAのCUDA Cコンパイラnvccは、アーキテクチャ特有のCUBINファイルとそれぞれのカーネルの上位互換のあるPTX版のいずれも生成するのに使うことができます。
カーネルのPTX版をすでに含むアプリケーションでは、FermiベースのGPUで現状のままで動くはずです。しかしながら、CUBINファイル経由で特定のGPUアーキテクチャのみをサポートするアプリケーションは、Fermiと将来のGPUへジャストインタイム(JIT)でコンパイルされるカーネルのPTX版を提供するか、カーネルのFermi特定のCUBIN版を含むよう改定するかのいずれかを行う必要があります。このため、アプリケーションがリリースされた後に導入されたCUDAアーキテクチャへの上位互換性を保証するために、すべてのアプリケーションは、カーネルのPTX版の起動をサポートすることが推奨されます。
それぞれのCUBINファイルは、特定のcompute capabilityバージョンをターゲットとし、同じメジャーバージョン番号のCUDAアーキテクチャとのみ上位互換性があります:たとえば、 compute capability 1.0をターゲットとするCUBINファイルは、すべてのcompute capability 1.x (Tesla)デバイスをサポートしますが、compute capability 2.0 (Fermi)デバイスをサポートしません。
CUDA Toolkit バージョン2.1から2.3を使いビルドされたCUDAアプリケーションは、カーネルのPTX版が含まれるようビルドされた限りは、Fermiと互換性があります。NVIDIAドライババージョン 195.xxもしくはそれ以降は、アプリケーションがPTX JITコードパスを使うのを可能にします。PTX JUTがあなたのアプリケーションで動作しているかをテストするために、以下を行うことができます
CUDA Toolkit バージョン3.0以降を使いビルドされたCUDAアプリケーションは、FermiネイティブのCUBINフォーマット(1.3参照)もしくはPTXフォーマット(1.2.1.1参照)のいずれかもしくは両方でカーネルを含むようビルドされた限りFermiと互換性があります。
warp内でスレッドがシェアードもしくはグローバルメモリ経由でお互いに値を通信する必要があるとき、一般的な最適化は、値をメモリに書いた後で__syncthreads()を省略することです(5.4.3とCUDA C Programming GuideのB.2.5参照)。この場合、warp内でスレッドの実行の同期性のために、__syncthreads()は省略することができます
この最適化のよくあるアプリケーションは、並列リダクションでであり、それは、それぞれの出力がすべての入力に依存する問題のセットをカバーする一般的なデータ並列演算です(つまり、大きなグループの数の和を求めたり、大きなセットの値の中で値nの実体を数えるなどです)。そのようなアプリケーションは、たびたびこの節の最後の例(それは、GPU Computing SDKからのreductionサンプルからの単純化した抜粋です)に似たコードを採用します。
warp内でシェアードもしくはグローバルメモリを使いスレッド間の値を渡す際、もしあなたのカーネルがこの種の最適化を実装するなら、そのメモリへのポインタが、コンパイラが値(以下の例では smem[tid])をレジスタで保持するのではなく、それぞれのステップ後に中間値をメモリへ完全に書くことを強制するvolatile修飾子で宣言する(元文書では赤で示されている)ことは必要不可欠です。
volatile修飾子を省略するコードは、強化されたコンパイラ最適化のため、Fermi上では正しく動作しません。以下の例では、volatile修飾子が、単純にレジスタに割り当て、シェアードメモリへの書き込みを取り除く最適化ではなく、すべての割り当ての後にsmem[tid]をシェアードメモリに書き戻さなければならないことをコンパイラに伝えています。
__device__ void reduce(float *g_idata, float *g_odata)
{
unsigned int tid = threadIdx.x;
extern __shared__ float sdata[];
sdata[tid] = g_idata[...]; // 初期値を割り当てる
__syncthreads();
// シェアードメモリでリダクションをする
// この例はブロックサイズが256であることを仮定している;
// 完全で一般的な実装はGPU Computing SDKの
// "reduction"サンプルを参照のこと
if (tid < 128) {sdata[tid]+=sdata[tid+128];} __syncthreads();
if (tid < 64) {sdata[tid]+=sdata[tid+ 64];} __syncthreads();
if (tid < 32) {
// それぞれのwarpで32スレッドが、
// お互い横並びで実行するため、
// volatileで宣言されたポインタ経由でデータにアクセスする限り、
// 以下の行のそれぞれの後に__syncthreads()は必要ない
volatile float *smem = sdata;
smem[tid] += smem[tid + 32];
smem[tid] += smem[tid + 16];
smem[tid] += smem[tid + 8];
smem[tid] += smem[tid + 4];
smem[tid] += smem[tid + 2];
smem[tid] += smem[tid + 1];
}
// このブロックに対する結果をグローバルメモリに書き込む
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
CUDA Toolkit 2.1、2.2と2.3に含まれるコンパイラは、TeslaアーキテクチャにネイティブのCUBINファイルを生成しますが、FermiアーキテクチャにネイティブなCUBINファイルを生成することができません(これはCUDA Toolkit 3.0以降を必要とします)。CUDA Toolkitの2.xバージョンを使うとき、Fermiと将来のアーキテクチャのサポートを可能にするために、コンパイラは、それぞれのカーネルの PTXバージョンを生成することができます。標準では、PTXバージョンは、実行ファイルに含まれ、ジャストインタイム(JIT)コンパイル経由で Fermiデバイス上で実行が可能です。
CUDA Toolkitのバージョン3.0から、nvccはFermiアーキテクチャにネイティブなCUBINファイルも作ることが可能です。CUDA Toolkit 3.0以降を利用する際、将来のGPUアーキテクチャのためのPTXバージョンだけでなく、すべてのリリースされたGPUアーキテクチャに対する CUBINファイルをnvccが生成することを確実にするために、いかに示すようにnvccコマンドラインで適切な"-arch=sm_xx"パラメータを指定してください。
CUDAアプリケーションが、カーネルを起動するとき、CUDAランライムライブラリ(CUDART)はシステムでのそれぞれのGPUのcompute capabilityを決定し、最も一致するカーネルのCUBINまたはPTXバージョンを探すためにこの情報を使います。アプリケーションを起動する GPUのアーキテクチャをサポートするCUBINファイルがもし利用可能なら、それを使います。それ以外はCUDAランタイムは、PTXを読み込み、 GPU上で起動するためにPTXからCUBINへJITコンパイルします。
以下はTeslaデバイスでネイティブに動きPTX経由でFermiデバイスで動く、cuda_kernel.cuをビルドするコンパイラ設定です。ネイティブコードを提供する主な利点は、エンドユーザに対しPTXへコンパイルされたCUDAカーネルをPTX JITするのにかかる時間を節約する点です。しかしながら、CUDAドライバは、PTX JITの結果として生成されたネイティブのISAをキャッシュするため、これはほぼ一回のコストです。CUDAランタイムは現在のGPUのアーキテクチャを確認し、CUDAカーネルの最も有効なバージョンを明示的に呼ぶため、追加の起動毎のオーバヘッドがまだあるでしょう。
Windows:
nvcc.exe -ccbin "C:\vs2008\VC\bin" -I"C:\CUDA\include" -Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT" –arch=sm_10 --compile -o "Release\cuda_kernel.cu.obj" "cuda_kernel.cu"
Mac/Linux:
/usr/local/cuda/bin/nvcc -arch=sm_10 --compiler-options -fno-strict-aliasing -I. -I/usr/local/cuda/include -DUNIX -O2 -o release/cuda_kernel.cu.o -c cuda_kernel.cu
注意:nvccのコマンドラインオプション"-arch=sm_xx"は以下のより明示的な-gencodeコマンドラインオプションの省略された等価内容です。
–gencode=arch=compute_xx,code=sm_xx
–gencode=arch=compute_xx,code=compute_xx
-gencodeオプションは、以下に示すよう複数のターゲットアーキテクチャに対するCUBINやPTXコードをコンパイルなら-archの代わりに使用しなければなりません。
代わりに、CUDA Toolkitのバージョン 3.0では、コンパイラは以下に示すようにTeslaデバイスとFermiデバイス両方でネイティブに動くcuda_kernel.cuをビルドすることができます。この例は、また上位互換性のあるPTXコードでもビルドします。
Windows:
nvcc.exe -ccbin "C:\vs2008\VC\bin" -I"C:\CUDA\include"
-Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT"
-gencode=arch=compute_10,code=sm_10
-gencode=arch=compute_10,code=compute_10
-gencode=arch=compute_20,code=sm_20
-gencode=arch=compute_20,code=compute_20
--compile -o "Release\cuda_kernel.cu.obj" "cuda_kernel.cu"
Mac/Linux:
/usr/local/cuda/bin/nvcc
-gencode=arch=compute_10,code=sm_10
-gencode=arch=compute_10,code=compute_10
-gencode=arch=compute_20,code=sm_20
-gencode=arch=compute_20,code=compute_20
--compiler-options -fno-strict-aliasing -I.
-I/usr/local/cuda/include -DUNIX
-O2 -o release/cuda_kernel.cu.o -c cuda_kernel.cu
-gencodeへの"code=sm_10"引数(特定のcompute capabilityに対するCUBINファイルを生成します)と"code=compute_10"引数(そのcompute capabilityに対するPTXを生成します)のコマンドラインの区別に注意してください。
どんなステップをFermiをサポートするためにとる必要がありますか?
答え:いくつかの選択肢があります。
* CUDAカーネルファイルをPTXにコンパイルする方法。CUBINファイルは、CUDA Toolkit 2.1から2.3のコンパイルを使い生成することができますが、それらのCUBINファイルは、Teslaデバイスとのみ互換性があり、Fermiデバイスにはありません。
以下のCUDAドライバAPIがPTXカーネルを起動する方法を示すGPU Computing SDKコード例を参照してください。
CUDAソースファイルからPTX出力を作るために以下のコンパイラ設定を使用してください。
Windows:
nvcc.exe -ccbin "C:\vs2008\VC\bin" -I"C:\CUDA\include"
-Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT"
-ptx
–o "cuda_kernel.ptx" "cuda_kernel.cu"
Mac/Linux:
/usr/local/cuda/bin/nvcc
-ptx
--compiler-options -fno-strict-aliasing -I.
-I/usr/local/cuda/include -DUNIX -O2
-o cuda_kernel.ptx cuda_kernel.cu
*CUDAカーネルをCUBINとPTX出力ファイルの両方へコンパイルする方法。nvccはどちらのタイプのそれぞれの生成された出力のために一度呼ばれなければならないため、これはコンパイル時に明示的に指定されなければいけません。
実行時に、アプリケーションは、以下のCUDAドライバAPI関数で現在のGPUのcompute capabilityを明示的に検査します。この関数の使い方の詳細な例についてはGPU Computing SDK内のdeviceQueryDrvコード例を参照してください。
cuDeviceComputeCapability(&major, &minor, dev)
この関数で返却されたメジャーバージョンとマイナーバージョンに基づいて、アプリケーションは、それぞれのカーネルの適切なCUBINもしくはPTXバージョンを選択することができます。
CUDA ドライバAPIを使いPTXへコンパイルされたカーネルを読み込むために、以下の例のようなコードを使うことができます。 cuModuleLoadDataExの呼び出しは、PTXソースファイルをJITコンパイルします(開発者がカーネルを正しくコンパイルするために気を使う必要のある少しのJITオプションがあることに気をつけてください)。GPU Computing SDKの例matrixMulDrvとsimple TextureDrvはこの手順をさらに説明します。
CUmodule cuModule;
CUfunction cuFunction = 0;
string ptx_source;
// ヘルパー関数がPTXソースをstringに読み込む
findModulePath ("matrixMul_kernel.ptx",
module_path, argv, ptx_source));
// 我々はPTXJITコンパイルをパラメータで指定する
const unsigned int jitNumOptions = 3;
CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
void **jitOptVals = new void*[jitNumOptions];
// コンパイルログバッファのサイズをセットアップする
jitOptions[0] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
int jitLogBufferSize = 1024;
jitOptVals[0] = (void *)jitLogBufferSize;
// コンパイルログバッファへのポインタをセットアップする
jitOptions[1] = CU_JIT_INFO_LOG_BUFFER;
char *jitLogBuffer = new char[jitLogBufferSize];
jitOptVals[1] = jitLogBuffer;
// レジスタの最大数へのポインタをセットアップする
jitOptions[2] = CU_JIT_MAX_REGISTERS;
int jitRegCount = 32;
jitOptVals[2] = (void *)jitRegCount;
// モジュールの読み込みはPTXがJITであることを強制する
status = cuModuleLoadDataEx(&cuModule, ptx_source.c_str(),
jitNumOptions, jitOptions,
(void **)jitOptVals);
printf("> PTX JIT log:\n%s\n", jitLogBuffer);
(本文の訳はこれで完了)