始めましょう。
GPUコンピューティングモデル：
GPU計算モデルをより詳細に検討してください。
- GPUコアの上位レベルは、グリッドまたは次元N1 * N2 * N3のグリッドにグループ化されたブロックで構成されます。 これは次のように表すことができます。
図 1.コンピューティングデバイスのGPU。
ブロックのグリッドの次元はcudaGetDeviceProperties関数を使用して見つけることができます;結果の構造体では、maxGridSizeフィールドがこれを担当します。 たとえば、GeForce 9600M GSでは、ブロックグリッドの次元は65535 * 65535 * 1です。つまり、ブロックグリッドは2次元です（受信データはCompute Capability v.1.1を満たします）。
- すべてのブロックはスレッドで構成され、スレッドは計算の直接実行者です。 ブロック内のスレッドは3次元配列の形式で形成され（図2）、その次元はcudaGetDeviceProperties関数を使用して見つけることもできます; maxThreadsDimフィールドがこれを担当します。
図 2. GPUユニットデバイス。
GPUを使用する場合、必要なサイズのグリッドを使用して、タスクのニーズに合わせてブロックを構成できます。
CUDAおよびC言語：
CUDAテクノロジー自体（nvcc.exeコンパイラー）は、C言語の追加の拡張機能をいくつか導入します。これらは、GPUのコードを記述するために必要です。
- 関数がどのように、どこから実行されるかを示す関数指定子。
- 使用されるGPUメモリのタイプを示す変数修飾子。
- GPUカーネル起動修飾子。
- GPUコアでコードを実行するときにスレッド、ブロック、およびその他のパラメーターを識別するための組み込み変数。
- 追加のタイプの変数。
前述のように、関数指定子は、関数を呼び出す方法と場所を決定します。 合計で、CUDAには次の3つの修飾子があります。
- __host__- CPUから実行され、CPUから呼び出されます（原則として、省略することができます）。
- __global__- GPUで実行され、CPUから呼び出されます。
- __device__- GPUで実行され、GPUから呼び出されます。
カーネル開始修飾子は、GPUで計算するときに割り当てるブロック、スレッド、およびメモリの数を記述するために使用されます。 カーネルの起動構文は次のとおりです。
myKernelFunc <<< gridSize、blockSize、sharedMemSize、cudaStream >>>（float * param1、float * param2）、ここで
- gridSize-計算に割り当てられたブロックのグリッドの次元（dim3）、
- blockSize-計算に割り当てられたブロックサイズ（dim3）、
- sharedMemSize-カーネルの起動時に割り当てられる追加メモリのサイズ、
- cudaStream-呼び出しが行われるストリームを定義するcudaStream_t変数。
そしてもちろん、myKernelFunc自体はカーネル関数（__global__指定子）です。 sharedMemSizeやcudaStreamなど、一部の変数はカーネルを呼び出すときに省略できます。
また、組み込み変数に言及する価値があります。
- gridDim - dim3タイプのグリッドディメンション。 現在のカーネル呼び出し中に割り当てられたグリッドのサイズを知ることができます。
- blockDim - dim3タイプのブロック次元。 現在のカーネル呼び出し中に割り当てられたブロックのサイズを確認できます。
- blockIdx -GPUでの計算における現在のブロックのインデックスは、uint3型です。
- threadIdx -GPUでの計算における現在のスレッドのインデックス。タイプはuint3です。
- warpSize-ワープのサイズ。int型です（自分で使用しようとしませんでした）。
ちなみに、gridDimとblockDimは、GPUカーネルの起動時に渡すまさに変数ですが、カーネルでのみ読み取ることができます。
追加のタイプの変数とその修飾子は、メモリを操作する例で直接考慮されます。
CUDAホストAPI：
CUDAをコンピューティングに直接使用する前に、いわゆるCUDAホストAPIに慣れる必要があります。これは、CPUとGPUの間のリンクです。 CUDAホストAPIは、CUDAユーザーモードドライバーへのアクセスを提供するCUDAドライバーAPIと呼ばれる低レベルAPIと、CUDAランタイムAPIに分けられます。 私の例では、CUDAランタイムAPIを使用します。
CUDAランタイムAPIには、次の機能グループが含まれています。
- デバイス管理 -一般的なGPU制御（GPU機能に関する情報の受信、SLIモードでの作業中のGPU間の切り替えなど）のための機能が含まれます。
- スレッド管理 -スレッド管理。
- ストリーム管理 - ストリーム管理。
- イベント管理 - イベントを作成および管理する機能。
- 実行制御 -CUDAカーネルの起動および実行機能。
- メモリ管理 -GPUメモリ管理機能。
- テクスチャ参照マネージャ -CUDAを介してテクスチャオブジェクトを操作します。
- OpenGL相互運用性 -OpenGL APIと対話するための機能。
- Direct3D 9の相互運用性 -Direct3D 9 APIと対話するための機能。
- Direct3D 10相互運用性 -Direct3D 10 APIと対話するための機能。
- エラー処理 -エラー処理関数。
GPUの動作を理解する：
前述のように、スレッドは計算を直接実行します。 では、スレッド間の計算の並列化はどのように行われますか？ 単一ブロックの操作を検討してください。
チャレンジ。 次元N要素の2つのベクトルの合計を計算する必要があります。
ブロックの最大サイズは512 * 512 * 64スレッドです。 ベクトルは1次元であるため、これまでのところ、ブロックのx次元の使用に制限します。つまり、ブロックからスレッドのストリップを1つだけ使用します（図3）。
図 3.使用済みブロックからのスレッドのストリップ。
ブロック512のx次元、つまり、一度にベクトルを加算できることに注意してください。ベクトルの長さはN <= 512要素です。 他の場合、より大規模な計算では、より多くのブロックと多次元配列を使用できます。 また、1つのブロックで512 * 512 * 64 = 16777216のスレッドを使用できると思った人もいるかもしれませんが、これはそうではありません。一般に、この製品は512（少なくとも私のグラフィックカード）。
プログラム自体で、次の手順を実行する必要があります。
- 計算用のデータを取得します。
- このデータをGPUメモリにコピーします。
- カーネル関数を介してGPUで計算を実行します。
- 計算されたデータをGPUメモリからRAMにコピーします。
- 結果を表示します。
- 使用済みのリソースを解放します。
コードの記述に直接進みます。
まず最初に、ベクトルの追加を実行するカーネル関数を作成します。
//
__global__ void addVector( float * left, float * right, float * result)
{
// id .
int idx = threadIdx.x;
// .
result[idx] = left[idx] + right[idx];
}
* This source code was highlighted with Source Code Highlighter .
したがって、カーネルの起動時に並列化が自動的に実行されます。 この関数は、組み込み変数threadIdxとそのフィールドxも使用します。これにより、ベクトル要素の計算とブロック内のスレッド間の対応を設定できます。 ベクトルの各要素の計算は、個別のスレッドで行います。
プログラムのポイント1および2を担当するコードを記述します。
#define SIZE 512
__host__ int main()
{
//
float * vec1 = new float [SIZE];
float * vec2 = new float [SIZE];
float * vec3 = new float [SIZE];
//
for ( int i = 0; i < SIZE; i++)
{
vec1[i] = i;
vec2[i] = i;
}
//
float * devVec1;
float * devVec2;
float * devVec3;
//
cudaMalloc(( void **)&devVec1, sizeof ( float ) * SIZE);
cudaMalloc(( void **)&devVec2, sizeof ( float ) * SIZE);
cudaMalloc(( void **)&devVec3, sizeof ( float ) * SIZE);
//
cudaMemcpy(devVec1, vec1, sizeof ( float ) * SIZE, cudaMemcpyHostToDevice);
cudaMemcpy(devVec2, vec2, sizeof ( float ) * SIZE, cudaMemcpyHostToDevice);
…
}
* This source code was highlighted with Source Code Highlighter .
ビデオカードにメモリを割り当てるには、次のプロトタイプを持つcudaMalloc関数を使用します。
cudaError_t cudaMalloc（void ** devPtr、size_t count）、ここで
- devPtr-割り当てられたメモリのアドレスへのポインタ、
- count-割り当てられたメモリのサイズ（バイト単位）。
戻り値：
- cudaSuccess-メモリの割り当てが成功したとき
- cudaErrorMemoryAllocation-メモリ割り当てエラー
データをビデオカードのメモリにコピーするには、次のプロトタイプを持つcudaMemcpyを使用します。
cudaError_t cudaMemcpy（void * dst、const void * src、size_t count、enum cudaMemcpyKind kind）、ここで
- dst-コピー先のアドレスを含むポインター、
- src-コピー元アドレスを含むポインター、
- count-コピーされたリソースのサイズ（バイト単位）
- cudaMemcpyKind-コピーの方向を示す列挙（cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyHostToHost、cudaMemcpyDeviceToDeviceなど）。
戻り値：
- cudaSuccess-コピーが成功したとき
- cudaErrorInvalidValue-無効な引数パラメーター（たとえば、コピーサイズが負）
- cudaErrorInvalidDevicePointer-グラフィックスカードの無効なメモリポインター
- cudaErrorInvalidMemcpyDirection-間違った方向（たとえば、コピーのソースと宛先が混同されている）
次に、GPUでの計算のためのカーネルの直接呼び出しに目を向けます。
…
dim3 gridSize = dim3(1, 1, 1); //
dim3 blockSize = dim3(SIZE, 1, 1); //
//
addVector<<<gridSize, blockSize>>>(devVec1, devVec2, devVec3);
…
* This source code was highlighted with Source Code Highlighter .
この場合、グリッドとブロックのサイズを決定する必要はありません。ブロックとブロックで1つの測定値のみを使用するため、上記のコードを記述できます。
addVector<<<1, SIZE>>>(devVec1, devVec2, devVec3);
* This source code was highlighted with Source Code Highlighter .
ここで、計算結果をビデオメモリからホストメモリにコピーするだけです。 しかし、カーネル関数には特異性があります-非同期実行、つまり、カーネルの呼び出し後にコードの次のセクションが動作を開始した場合、これはGPUが計算を実行したことを意味しません。 特定のカーネル機能を完了するには、イベントなどの同期ツールを使用する必要があります。 したがって、結果をホストにコピーする前に、イベントを介してGPUスレッドを同期します。
カーネルを呼び出した後のコード：
//
addVector<<<blocks, threads>>>(devVec1, devVec2, devVec3);
// event'
cudaEvent_t syncEvent;
cudaEventCreate(&syncEvent); // event
cudaEventRecord(syncEvent, 0); // event
cudaEventSynchronize(syncEvent); // event
//
cudaMemcpy(vec3, devVec3, sizeof ( float ) * SIZE, cudaMemcpyDeviceToHost);
* This source code was highlighted with Source Code Highlighter .
Event Managment APIの関数を詳しく見てみましょう。
イベントはcudaEventCreate関数を使用して作成されます。プロトタイプの形式は次のとおりです。
cudaError_t cudaEventCreate（cudaEvent_t * event）、ここで
- * event-イベントハンドルを書き込むポインタ。
戻り値：
- cudaSuccess-成功した場合
- cudaErrorInitializationError-初期化エラー
- cudaErrorPriorLaunchFailure-前回の非同期機能開始中のエラー
- cudaErrorInvalidValue-無効な値
- cudaErrorMemoryAllocation-メモリ割り当てエラー
イベントの記録は、次の形式のプロトタイプのcudaEventRecord関数を使用して実行されます。
cudaError_t cudaEventRecord（cudaEvent_tイベント、CUstreamストリーム）、ここで
- event-イベントのハンドル
- stream-書き込むストリームの番号（この場合、これはメインのゼロストリームです）。
戻り値：
- cudaSuccess-成功した場合
- cudaErrorInvalidValue-無効な値
- cudaErrorInitializationError-初期化エラー
- cudaErrorPriorLaunchFailure-前回の非同期機能開始中のエラー
- cudaErrorInvalidResourceHandle-無効なイベントハンドル
イベントの同期は、cudaEventSynchronize関数によって実行されます。 この関数は、すべてのGPUスレッドの完了と、指定されたイベントの通過を期待し、その後、呼び出しプログラムに制御を渡します。 関数のプロトタイプは次のとおりです。
cudaError_t cudaEventSynchronize（cudaEvent_tイベント）、ここで
- event-通過が予想されるイベントハンドル。
戻り値：
- cudaSuccess-成功した場合
- cudaErrorInitializationError-初期化エラー
- cudaErrorPriorLaunchFailure-前回の非同期機能開始中のエラー
- cudaErrorInvalidValue-無効な値
- cudaErrorInvalidResourceHandle-無効なイベントハンドル
次の図からcudaEventSynchronizeの仕組みを理解できます。
図 4.メインプログラムとGPUプログラムの動作の同期。
図4では、ブロック「イベントの通過を待機しています」はcudaEventSynchronize関数呼び出しです。
結論として、結果を画面に表示し、割り当てられたリソースをクリーンアップします。
//
for ( int i = 0; i < SIZE; i++)
{
printf( "Element #%i: %.1f\n" , i , vec3[i]);
}
//
//
//
cudaEventDestroy(syncEvent);
cudaFree(devVec1);
cudaFree(devVec2);
cudaFree(devVec3);
delete[] vec1; vec1 = 0;
delete[] vec2; vec2 = 0;
delete[] vec3; vec3 = 0;
* This source code was highlighted with Source Code Highlighter .
リソースを解放する機能を記述する必要はないと思います。 作業を確認する必要がある場合は、cudaError_t値も返すことを思い出してください。
おわりに
この資料がGPUの仕組みを理解するのに役立つことを願っています。 CUDAで作業するために知っておく必要がある最も重要なポイントについて説明しました。 2つのマトリックスの追加を自分で記述してみてください。ただし、ビデオカードのハードウェア制限を忘れないでください。
PS：それほど簡単には機能しませんでした。 疲れないことを願っています。 すべてのソースコードが必要な場合は、メールで送信できます。
PSS：質問してください。