メモリ・アクセスを速くする
お次はGPU内で動くデバイス側コード（kernel）の高速化。cudaMemcpyでコピー元／先に指定されるのはデバイス内メモリ、すなわちビデオ・カード上のビデオ・メモリです。kernel呼び出し時に引数に渡しているのはビデオ・メモリ上の位置を指すポインタで、kernelはGPUの"外部"にあるビデオ・メモリをアクセスしています。
GPUはチップ"内部"にもメモリを持っていて、GPUの外にあるビデオ・メモリよりずっと高速。kernelがこのGPUチップ内メモリを使うよう仕向けてやればメモリ・アクセスの頻度が多いほど高速になるはずです。
簡単な例としてSMA（Simple Moving Average：単純移動平均）を考えます。SMAは株価や気温といった均等間隔に並ぶデータの平滑化、デコボコを埋めて滑らかにする計算です。
例えば毎日正午の気温を測定した値がsize日分、float input[size]に納められているとしましょう。input[0]が本日、input[i]がi日前の正午の気温データです。そうするとi日前における過去N日間分の平均気温 result[i] は、
result[i] = (input[i] + input[i+1] + …… + input[i+N-1]) / N
ですね。これを i = 0..size-N について計算しようと。
まず、素直に書いたkernelはこんなところでしょうか。
__global__ void kernel_SMA(float* result, const float* input, size_t size, size_t N) {
size_t i = blockDim.x * blockIdx.x + threadIdx.x;
if ( i <= size - N ) {
float sum = 0.0f;
for ( size_t t = 0; t < N; ++t ) {
sum += input[i+t];
}
result[i] = sum / N;
}
}
入力データinputの各要素はN回読み込まれます。いったん読みだしてGPU内部のメモリに置いとけば、それ以降より高速な読み出しができるハズ。
・高速化の方法1つめ、read-onlyキャッシュを紹介しましょう、こんなコードです。
__global__ void kernel_SMA_ROCache(float* result, const float* __restrict__ input, size_t size, size_t N) {
size_t i = blockDim.x * blockIdx.x + threadIdx.x;
if (i <= size - N) {
float sum = 0.0f;
for (size_t t = 0; t < N; ++t) {
sum += input[i + t];
}
result[i] = sum / N;
}
}
素直に書いたコードとの違いは引数：const float* inputが__restrict__で修飾されているところ、それだけ。__restrict__は"このポインタが指す先の値は変化しない"ことをコンパイラに教えてあげるキーワードです。これによってinput[～]は変化しない、すなわちread-only であるとみなされ、GPU内のread-onlyキャッシュに置いて最適化するヒントとなります。
ただし、__restrict__が効果を現すのはGPUのCompute Capabilityが3.5以降のもの。有効にするにはコンパイル時のオプションを調整しておかなくてはなりません。
・高速化の方法2つめはshared-memory（共有メモリ）の利用です。
kernelコードでshared-memoryを利用するには、変数を__shared__修飾します。例えば、
__shared__ int data; __shared__ float buffer[100];
のように。__shared__修飾された変数はkernel呼び出し時に指定したブロックに属する全スレッドで共有されます（ブロックをまたいだ共有はできません）。1つのスレッドで書いたbuffer[]を同一ブロック内の他のスレッドが読めるわけ。
shared-memoryを利用したSMAはこうなりました。
__global__ void kernel_cached_SMA(float* result, const float* input, size_t size, size_t N) {
size_t gid = blockDim.x * blockIdx.x + threadIdx.x;
extern __shared__ float cache[]; // block内の全threadで共有される
if (gid < size) // 値が存在すれば書いてよい
cache[threadIdx.x] = input[gid];
// blockDim.xで足りない分、オフセットを与えて読み込む。
if ((threadIdx.x < N - 1) && (gid + blockDim.x < size))
cache[threadIdx.x + blockDim.x] = input[gid + blockDim.x];
syncthreads(); // block内の全threadがここで待ち合わせ
if (gid < size - N + 1) {
// cacheを入力データとしてSMAを求める
float sum = 0.0f;
for ( size_t t = 0; t < N; ++t )
sum += cache[threadIdx.x + t];
result[gid] = sum / N;
}
}
extern __shared__ float cache[];に要素数が書かれていません。extern __shared__としておくと、その大きさ（バイト数）を動的に外から与えることができます。呼び出し時、kernel_cached_SMA<<<>>>の3つめのパラメータがextern __shared__したブロック内共有領域のバイト数となります。
extern __shared__ float cache[]の中身を各スレッドで手分けして埋めています。syncthreads()は、ブロック内の全スレッドがこの地点で待ち合わせるため。キャッシュが全部埋まるのを待たないと、その後に続く計算結果がおかしくなっちゃいますからね。
500回ほど繰り返し、キャッシュの効果を調べてみたところ……。
enum KernelType {
SMA = 0,
SMA_ROCache,
cached_SMA,
};
chrono::milliseconds device_SMA(float* result, const float* input, size_t size, size_t N, KernelType kernelType, int iteration) {
chrono::high_resolution_clock::time_point start, stop;
float* dinput = nullptr;
float* dresult = nullptr;
CUERRCHK(cudaMalloc((void**)&dinput, size*sizeof(float)));
CUERRCHK(cudaMalloc((void**)&dresult, size*sizeof(float)));
CUERRCHK(cudaMemcpy(dinput, input, size*sizeof(float), cudaMemcpyHostToDevice));
CUERRCHK(cudaMemset(dresult, 0, size*sizeof(float))); // テストのため。0に初期化しておく。
start = chrono::high_resolution_clock::now();
const unsigned int block = 256; // 1-blockあたりのthread数
assert( block > N );
const unsigned int grid = (size + block -1) / block; // 1-grid あたりの block数
for ( int i = 0; i < iteration; ++i ) {
switch (kernelType) {
case SMA:
kernel_SMA <<<grid, block >>>(dresult, dinput, size, N);
break;
case SMA_ROCache:
kernel_SMA_ROCache<<<grid, block >>>(dresult, dinput, size, N);
break;
case cached_SMA:
kernel_cached_SMA<<<grid, block, (block + N - 1) * sizeof(float) >>>(dresult, dinput, size, N);
break;
}
}
CUERRCHK(cudaGetLastError());
cudaError_t err = cudaDeviceSynchronize();
assert(err == cudaSuccess);
stop = chrono::high_resolution_clock::now();
CUERRCHK(cudaMemcpy(result, dresult, size*sizeof(float), cudaMemcpyDeviceToHost));
CUERRCHK(cudaFree(dinput));
CUERRCHK(cudaFree(dresult));
return chrono::duration_cast<chrono::milliseconds>(stop - start);
}
なるほど、確かにGPU内メモリが効いてますね。処理の過程で特定の領域を何度もアクセスしなきゃいけないときは、かなりの効果が期待できそうです。
