前のシリーズをまだ見ていない人のために：CUDAを使用すると、ビデオカードnVidia（8xxx以降）の特別な拡張機能を使用して、C ++で記述されたプログラムを作成および実行できます。 適切なタスクでは、従来のCPUと比較して$の大幅なパフォーマンスの優位性が達成されます。
GTX295で1秒間に1兆回以上のオペレーションを達成できます。
NB：この記事は簡単な紹介であり、CUDAの下でのプログラミングのすべてのニュアンスを1つの記事でカバーすることはほとんど不可能です:-)
鉄について
CUDAは、8400GS以降のビデオカードで実行されます。 ビデオカードごとに機能が異なります。 一般に、ビデオカードにたとえば128 SP（ストリーミングプロセッサ）がある場合、これは8つのSIMD MP（マルチプロセッサ）があり、それぞれが16の操作を同時に行うことを意味します。 1つのMPには、16kbの共有メモリ、8192個の4バイトレジスタがあります（GTX2xxシリーズカードにはさらに多くの値があります）。 すべてのMPに共通の64kbの定数もあり、キャッシュされます。キャッシュに入らない場合は、かなり大きな遅延（400〜600クロックサイクル）があります。 グローバルビデオカードメモリがあり、アクセスはキャッシュされません。テクスチャ（キャッシュ、2Dサンプル用に最適化されたキャッシュ）。 複数のビデオカードを使用するには、まず木材でSLIを無効にする必要があります。次に、ストリームで各ストリームを実行し、cudaSetDevice（）を呼び出します。
どこから始めますか？
CUDAでのプログラミング方法を学ぶ最速の方法は、SDKからいくつかの例を取り上げ、実行し、動作中に変更することです（実際、これはBarsWFを書いたときにやったことです）:-)
開始するには、 http：//www.nvidia.com/object/cuda_get.htmlにアクセスして、目的のビットサイズのオペレーティングシステム用のSDKとツールキットをダウンロードします。 （残念ながら、たとえば、32ビットSDKと64ビットツールキットは干渉できません）。 ビデオカードドライバーを最新バージョンに更新すると便利です（CUDAは急速に開発されているため、プログラムのユーザーとユーザーの両方に最新のfireを用意しておくと便利です）。 。
たとえば、Mandelbrot SDKの例をご覧ください。 最も重要なのは.cuファイルです。カスタムビルドルールに注意してください。
$(CUDA_BIN_PATH)\nvcc.exe -ccbin "$(VCInstallDir)bin" -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT -I"$(CUDA_INC_PATH)" -I./ -I../../common/inc -I"$(DXSDK_DIR)\Include" -o $(ConfigurationName)\Mandelbrot_sm10.obj Mandelbrot_sm10.cu
すべてのプロジェクトで使用できます。「../../ common / inc」の代わりに、絶対パス（または環境変数）を指定できます。
nvccは素晴らしい、ひどいCUDAコンパイラです。 出力では、ビデオカード用にコンパイルされたプログラムが既に含まれているオブジェクトファイルを生成します。
Mandelbrot_kernel.hのインターフェイスの説明に注意してください。ここでは、メインC ++プログラムから呼び出すカーネルを自分の手で説明する必要があります（通常はそれほど多くないので、これは怖くないです）。
SDKサンプルを実行した後、CUDAプログラムが通常のプログラムとどのように異なるかを検討できます。
注意：-keepオプションを追加すると、コンパイル後に多くの興味深い中間ファイルを見つけることができます。
関数定義
次の「修飾子」の前には、.cuファイル内の関数があります。
__device__-これは、関数がビデオカードでのみ実行されることを意味します。 通常のプロセッサ（ホスト）で実行されているプログラムから呼び出すことはできません。
__global__-この関数は、コンピューティングカーネルの始まりです。 ビデオカードで実行されますが、ホストからのみ起動されます。
__host__-ホストからのみ実行され、実行されます（つまり、通常のC ++関数）。 たとえば、関数の前に__host__と__device__を指定すると、2つのバージョンの関数がコンパイルされます（すべての組み合わせが有効であるとは限りません）。
データ定義
__device__-変数がビデオカードのグローバルメモリ（つまり512-1024MB以上）にあることを意味します。 ビデオカードでのコンピューティングの標準により非常に遅いメモリ（中央処理装置のメモリよりも数倍高速ですが）をできるだけ使用しないことをお勧めします。 このメモリには、異なるコンピューティングコアの呼び出し間でデータが保存されます。 ここのデータは、ホストを使用して読み書きできます
cudaMemcpy(device_variable, host_variable, size, cudaMemcpyHostToDevice); //cudaMemcpyDeviceToHost -
__constant__-定数メモリに変数を設定します。 定数の値はcudaMemcpyToSymbol関数でロードする必要があることに注意してください。 定数はすべてのスレッドから利用でき、速度はレジスタに匹敵します（キャッシュに入るとき）。
__shared__-スレッドブロックの共有メモリに変数を設定します（つまり、値はすべてで共有されます）。 ここでは注意してアプローチする必要があります-コンパイラはここでのアクセスを積極的に最適化します（volatile修飾子で絞ることができます）。競合状態になる可能性があるため、__ syncthreads（）を使用する必要があります。 データが確実に記録されるようにします。 共有メモリはバンクに分割され、2つのスレッドが同時に1つのバンクにアクセスしようとすると、バンクの競合が発生し、速度が低下します。
カーネル（__device__）で定義したすべてのローカル変数は、最高のアクセス速度のレジスタにあります。
スレッドが何に取り組むべきかをどのように学習するか
CUDAの主なアイデアは、問題を解決するために何千ものスレッドを開始するということです。したがって、ここでさらに説明することを恐れないでください:-)
200x200の画像に対して何らかの操作を行う必要があるとしましょう。 画像は10x10のピースに分割され、そのようなピースのピクセルごとに上流から開始します。 次のようになります。
dim3スレッド（10、10）; // 4次サイズ、10 * 10
dim3 grid（20、20）; //画像全体を覆うために必要な正方形の数
your_kernel <<< grid、threads >>>（image、200,200）; //この行は40'000スレッドを開始します（同時にではなく、約200-2000スレッドが同時に機能します）。
AMDのBrook +とは異なり、CUDAでどのスレッドをどのデータで処理するかを即座に決定します。すべてが異なります。渡されたカーネルパラメーターはすべてのスレッドで同じであり、スレッドはこれを行うためにデータを受信する必要があり、スレッドは計算する必要があります画像がある場所。 これには、マジック変数blockDim、blockIdxが役立ちます。
const int ix = blockDim.x * blockIdx.x + threadIdx.x;
const int iy = blockDim.y * blockIdx.y + threadIdx.y;
ixおよびiyでは、画像配列からソースデータを取得し、作業の結果を記録できる座標。
最適化
プログラムを非常に遅くしないようにする方法についてのいくつかの言葉（CPUよりも低速で実行するプログラムの作成は、10倍の速度で作業するよりもはるかに簡単です:-)）
- __global__メモリをできるだけ少なく使用します。
- __shared__メモリを使用する場合は、バンクの競合を避けてください（ただし、多くのタスクは共有メモリなしで解決できます）。
- 異なるスレッドが異なるパスを通るコード内の可能な限り少ないブランチ。 このようなコードは並行して実行されません。
- できるだけ少ないメモリを使用してください。 使用するメモリが少ないほど、コンパイラとハードウェアがカーネルをより積極的に実行できるようになります（たとえば、100個のスレッドを使用し、100個のレジスタを使用して1つのMPで同時に実行し、レイテンシを大幅に削減します）
動かない？
まず第一に、SDK（NVIDIA_CUDA_Programming_Guide、CudaReferenceManual、ptx_isa）と共にドキュメントを読む必要があります。その後、 公式フォーラムで質問できます。 ロシア語では、例えばフォーラムで私に尋ねることができます。そこで私は答えます:-)また、多くの人々がgpgpu.ruに住んでいます 。
この紹介が、ビデオカードのプログラミングを試してみる人々に役立つことを願っています。 問題/質問がある場合-私はお手伝いさせていただきます。 さて、私たちの目の前には、Brook +とSIMD x86の紹介があります
オリジナルはこちらhttp://3.14.by/ru/read/cuda-crash-course-vvedenie