Geekly Articles each Day

CUDA：メモリーの操作。 パートI

CUDAを使用する過程で、ビデオカードメモリの使用に関する質問に実際には触れませんでした。 このギャップを埋める時が来ました。

トピックは非常に膨大であるため、いくつかの部分に分けることにしました。 このパートでは、ビデオカードで利用できるメモリの主な種類について説明し、メモリの種類の選択がGPUでの計算のパフォーマンスにどのように影響するかの例を示します。

ビデオカードとメモリの種類


GPUを使用する場合、開発者は、レジスタ、ローカル、グローバル、共有、定数、およびテクスチャメモリのいくつかのタイプのメモリを使用できます。 これらのタイプのメモリにはそれぞれ特定の目的があり、その目的は技術的なパラメーター（速度、読み取りおよび書き込みアクセスのレベル）によって決まります。 メモリタイプの階層を図に示します。 1。


図 1.ビデオカードメモリの種類
  1. レジスタメモリは、あらゆる種類の中で最速です。 GPUで使用可能なレジスタの数は、すでによく知られているcudaGetDeviceProperties関数を使用して決定できます。 また、1つのGPUスレッドで使用可能なレジスタの数を計算することは難しくありません。これは、レジスタの合計数をブロック内のスレッド数とグリッド内のブロック数の積で割る必要があるためです。 すべてのGPUレジスタは32ビットです。 CUDAには、レジスタメモリを使用する明示的な方法はなく、レジスタにデータを配置するすべての作業はコンパイラによって行われます。
  2. ローカルメモリー （ローカルメモリー）は、コンパイラーが任意の関数に多数のローカル変数を使用して使用できます。 速度特性の観点では、ローカルメモリはレジスタメモリよりもかなり低速です。 nVidiaのドキュメントでは、最も必要な場合にのみローカルメモリを使用することを推奨しています。 ローカルメモリの使用をブロックする明示的なツールは提供されていません。したがって、パフォーマンスが低下した場合、コードを慎重に分析し、不要なローカル変数を除外する価値があります。
  3. グローバルメモリは、GPUから利用可能な最も遅いタイプのメモリです。 グローバル変数は、__ global__指定子を使用して選択することも、cudMallocXXXファミリの関数を使用して動的に選択することもできます。 グローバルメモリは主に、デバイスからホストから受信した大量のデータを保存するために使用されます。この移動は、関数cudaMemcpyXXXを使用して実行されます。 高性能を必要とするアルゴリズムでは、グローバルメモリを使用する操作の数を最小限に抑える必要があります。
  4. 共有メモリは、高速メモリタイプを指します。 共有メモリは、グローバルメモリへのアクセスを最小限に抑え、ローカル変数関数を格納するために使用することをお勧めします。 ストリームのスレッド間の共有メモリのアドレス指定は、単一ブロック内で同じであり、単一ブロック内のスレッド間でデータを交換するために使用できます。 __shared__指定子は、共有メモリにデータを保存するために使用されます。
  5. コンスタントメモリ （コンスタントメモリ）は、利用可能なGPUの中で非常に高速です。 コンスタントメモリの特徴は、ホストからデータを書き込む機能ですが、同時に、GPU内ではこのメモリからの読み取りのみが可能であり、GPUがその名前を決定します。 __constant__指定子は、データを定数メモリに保存するために提供されています。 定数メモリで配列を使用する必要がある場合、グローバルメモリとは異なり、動的メモリは定数メモリではサポートされないため、そのサイズを事前に指定する必要があります。 cudaMemcpyToSymbol関数は、ホストから定数メモリへの書き込み、およびデバイスからcudaMemcpyFromSymbolホストへのコピーに使用されます。ご覧のように、このアプローチはグローバルメモリを操作するアプローチとは多少異なります。
  6. 名前が示すように、 テクスチャメモリは 、主にテクスチャを操作するためのものです。 テクスチャメモリには、データのアドレス指定、読み取り、書き込みに固有の機能があります。 GPU画像処理を検討する場合、テクスチャメモリについて詳しく説明します。

共有メモリの使用例


もう少し上に、GPUのプログラミング時に使用できるさまざまなタイプのメモリについて簡単に説明しました。 次に、マトリックスの転置で共有メモリを使用する例を示します。

メインコードの記述を開始する前に、これをデバッグする簡単な方法を示します。 ご存知のように、CUDAランタイムAPIの関数はさまざまなエラーコードを返すことがありますが、前回はこれを考慮していませんでした。 生活を簡素化するために、次のマクロを使用してエラーをキャッチできます。
#define CUDA_DEBUG

#ifdef CUDA_DEBUG

#define CUDA_CHECK_ERROR(err) \
if (err != cudaSuccess) { \
printf( "Cuda error: %s\n" , cudaGetErrorString(err)); \
printf( "Error in file: %s, line: %i\n" , __FILE__, __LINE__); \
} \

#else

#define CUDA_CHECK_ERROR(err)

#endif

* This source code was highlighted with Source Code Highlighter .

ご覧のとおり、環境変数CUDA_DEBUGが定義されている場合、エラーコードがチェックされ、ファイルとそれが発生した行に関する情報が表示されます。 この変数は、デバッグ用にコンパイル中に有効にし、リリース用にコンパイルするときに無効にすることができます。

メインタスクに到達する。

共有メモリの使用が計算の速度にどのように影響するかを確認するには、グローバルメモリのみを使用する関数も記述する必要があります。
この関数を書きます：

//
//
// inputMatrix -
// outputMatrix -
// width - ( -)
// height - ( -)
//
__global__ void transposeMatrixSlow( float * inputMatrix, float * outputMatrix, int width, int height)
{
int xIndex = blockDim.x * blockIdx.x + threadIdx.x;
int yIndex = blockDim.y * blockIdx.y + threadIdx.y;

if ((xIndex < width) && (yIndex < height))
{
//
int inputIdx = xIndex + width * yIndex;

// -
int outputIdx = yIndex + height * xIndex;

outputMatrix[outputIdx] = inputMatrix[inputIdx];
}
}

* This source code was highlighted with Source Code Highlighter .


この関数は、元のマトリックスの行を結果マトリックスの列に単純にコピーします。 唯一の困難な瞬間は、マトリックス要素のインデックスの決定です。ここでは、カーネルを呼び出すときに、さまざまな次元のブロックとグリッドを使用できることを覚えておく必要があります。これには、組み込み変数blockDim、blockIdxが使用されます。

共有メモリを使用する転置関数を作成します。

#define BLOCK_DIM 16

// c
//
// inputMatrix -
// outputMatrix -
// width - ( -)
// height - ( -)
//
__global__ void transposeMatrixFast( float * inputMatrix, float * outputMatrix, int width, int height)
{
__shared__ float temp[BLOCK_DIM][BLOCK_DIM];

int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
int yIndex = blockIdx.y * blockDim.y + threadIdx.y;

if ((xIndex < width) && (yIndex < height))
{
//
int idx = yIndex * width + xIndex;

//
temp[threadIdx.y][threadIdx.x] = inputMatrix[idx];
}

//
__syncthreads();

xIndex = blockIdx.y * blockDim.y + threadIdx.x;
yIndex = blockIdx.x * blockDim.x + threadIdx.y;

if ((xIndex < height) && (yIndex < width))
{
//
int idx = yIndex * height + xIndex;

//
outputMatrix[idx] = temp[threadIdx.x][threadIdx.y];
}
}

* This source code was highlighted with Source Code Highlighter .


この関数では、共有メモリを2次元配列として使用します。
すでに述べたように、1つのブロック内の共有メモリのアドレス指定はすべてのスレッドで同じです。したがって、アクセスおよび記録中の衝突を避けるため、配列の各要素はブロックの1つのスレッドに対応します。
ソースマトリックスの要素を一時バッファーにコピーした後、__ syncthreads関数が呼び出されます。 この関数は、ブロック内のスレッドを同期します。 他の同期方法との違いは、GPUでのみ実行されることです。
最後に、元の行列の保存された要素が、転置規則に従って結果行列にコピーされます。
この関数は、共有メモリなしのバージョンよりも実行速度が遅いように見えるかもしれませんが、そこでは仲介者はいません。 しかし、実際には、グローバルメモリからグローバルメモリへのコピーは、大量のグローバルメモリ（共有メモリ-グローバルメモリ）よりもはるかに遅くなります。
行列配列の境界を手動で確認する必要があることに注意してください; GPUには配列の境界を追跡するハードウェアがありません。

最後に、CPUでのみ実行される転置関数を作成します。

// , CPU
__host__ void transposeMatrixCPU( float * inputMatrix, float * outputMatrix, int width, int height)
{
for ( int y = 0; y < height; y++)
{
for ( int x = 0; x < width; x++)
{
outputMatrix[x * height + y] = inputMatrix[y * width + x];
}
}
}

* This source code was highlighted with Source Code Highlighter .


ここで、計算用のデータを生成し、それらをホストからデバイスにコピーし、GPUを使用する場合、パフォーマンス測定を行い、リソースをクリーンアップする必要があります。
これらの手順は前回説明した手順とほぼ同じであるため、このフラグメントをすぐに引用します。

#define GPU_SLOW 1
#define GPU_FAST 2
#define CPU 3

#define ITERATIONS 20 //

__host__ int main()
{
int width = 2048; //
int height = 1536; //

int matrixSize = width * height;
int byteSize = matrixSize * sizeof ( float );

//
float * inputMatrix = new float [matrixSize];
float * outputMatrix = new float [matrixSize];

//
for ( int i = 0; i < matrixSize; i++)
{
inputMatrix[i] = i;
}

//
printf( "Select compute mode: 1 - Slow GPU, 2 - Fast GPU, 3 - CPU\n" );
int mode;
scanf( "%i" , &mode);

//
printMatrixToFile( "before.txt" , inputMatrix, width, height);

if (mode == CPU) // CPU
{
int start = GetTickCount();
for ( int i = 0; i < ITERATIONS; i++)
{
transposeMatrixCPU(inputMatrix, outputMatrix, width, height);
}
// CPU ( )
printf ( "CPU compute time: %i\n" , GetTickCount() - start);
}
else // GPU
{
float * devInputMatrix;
float * devOutputMatrix;

//
CUDA_CHECK_ERROR(cudaMalloc(( void **)&devInputMatrix, byteSize));
CUDA_CHECK_ERROR(cudaMalloc(( void **)&devOutputMatrix, byteSize));

//
CUDA_CHECK_ERROR(cudaMemcpy(devInputMatrix, inputMatrix, byteSize, cudaMemcpyHostToDevice));

//
dim3 gridSize = dim3(width / BLOCK_DIM, height / BLOCK_DIM, 1);
dim3 blockSize = dim3(BLOCK_DIM, BLOCK_DIM, 1);

cudaEvent_t start;
cudaEvent_t stop;

// event' GPU
CUDA_CHECK_ERROR(cudaEventCreate(&start));
CUDA_CHECK_ERROR(cudaEventCreate(&stop));

// GPU
cudaEventRecord(start, 0);

if (mode == GPU_SLOW) //
{
for ( int i = 0; i < ITERATIONS; i++)
{

transposeMatrixSlow<<<gridSize, blockSize>>>(devInputMatrix, devOutputMatrix, width, height);
}
}
else if (mode == GPU_FAST) //
{
for ( int i = 0; i < ITERATIONS; i++)
{

transposeMatrixFast<<<gridSize, blockSize>>>(devInputMatrix, devOutputMatrix, width, height);
}
}

//
cudaEventRecord(stop, 0);

float time = 0;
//
cudaEventSynchronize(stop);
// GPU
cudaEventElapsedTime(&time, start, stop);

//
printf( "GPU compute time: %.0f\n" , time);

//
CUDA_CHECK_ERROR(cudaMemcpy(outputMatrix, devOutputMatrix, byteSize, cudaMemcpyDeviceToHost));

//
//
//

CUDA_CHECK_ERROR(cudaFree(devInputMatrix));
CUDA_CHECK_ERROR(cudaFree(devOutputMatrix));

CUDA_CHECK_ERROR(cudaEventDestroy(start));
CUDA_CHECK_ERROR(cudaEventDestroy(stop));
}

// -
printMatrixToFile( "after.txt" , outputMatrix, height, width);

//
delete[] inputMatrix;
delete[] outputMatrix;

return 0;
}

* This source code was highlighted with Source Code Highlighter .


計算がCPUでのみ実行される場合、windows.hから接続されているGetTickCount（）関数を使用して計算時間を測定します。 GPUで計算時間を測定するには、プロトタイプの形式が次のcudaEventElapsedTime関数を使用します。

cudaError_t cudaEventElapsedTime（float * time、cudaEvent_t start、cudaEvent_t end）、ここで
  1. time-開始イベントと終了イベントの間の時間を記録するためのフロートへのポインタ（ミリ秒）、
  2. start-最初のイベントのハンドル
  3. end -2番目のイベントのハンドル。

戻り値：
  1. cudaSuccess-成功した場合
  2. cudaErrorInvalidValue-無効な値
  3. cudaErrorInitializationError-初期化エラー
  4. cudaErrorPriorLaunchFailure-前回の非同期機能開始中のエラー
  5. cudaErrorInvalidResourceHandle-無効なイベントハンドル


また、printMatrixToFile関数を使用して、ソースマトリックスと結果をファイルに書き込みます。 結果が正しいことを確認します。 この関数のコードは次のとおりです。

__host__ void printMatrixToFile( char * fileName, float * matrix, int width, int height)
{
FILE* file = fopen(fileName, "wt" );
for ( int y = 0; y < height; y++)
{
for ( int x = 0; x < width; x++)
{
fprintf(file, "%.0f\t" , matrix[y * width + x]);
}
fprintf(file, "\n" );
}
fclose(file);
}

* This source code was highlighted with Source Code Highlighter .


マトリックスが非常に大きい場合、データをファイルに出力すると、プログラムの実行が大幅に遅くなる可能性があります。

おわりに



テストプロセスでは、2048 * 1536 = 3145728要素の次元と20回のロードサイクルの行列を使用しました。 測定結果の後、次の結果を得ました（図2）。


図 2.決済時間。 （少ないほど良い）。

ご覧のとおり、共有メモリを備えたGPUバージョンは、CPU上のバージョンよりもほぼ20倍高速に実行されます。 また、共有メモリを使用する場合、計算はそれを使用しない場合よりも約4倍高速になることに注意してください。
私の例では、ホストからデバイスへ、またはその逆にデータをコピーする時間を考慮していませんが、実際のアプリケーションでは、それらも考慮する必要があります。 可能な限り、CPUとGPU間のデータ転送の数は最小限に抑える必要があります。

PS GPUで得られるパフォーマンスの向上をお楽しみください。

Source: https://habr.com/ru/post/J55461/

More articles:


All Articles