CUDA 4.0でマルチGPU化が限りなく簡単になっているという話

以前、複数のGPUでマンデルブロ集合を並列計算というブログ記事を書いた。このときはTeslaが1枚で7.45秒、Teslaが2枚で4.26秒で計算できたと報告したが、マルチGPU化は多少面倒だったし、ある程度のオーバーヘッドが掛かっていた。

しかし、CUDA 4.0がリリースされてから状況が一変した。自由にcudaSetDeviceが呼び出せるようになり、またUVA (Unified Virtual Addressing)を利用することにより、それはそれは簡単にマルチGPU化や、ホスト・複数GPU間での一元的なメモリの参照を実現できることになった。

というわけで、CUDA 4.0を使ったマンデルブロ集合のプログラムを以下に示す。シングルGPUからマルチGPUへ変更した部分については赤字で示してある。ただし、簡便のために、GPUデバイスは2枚で決め打ちとし、マンデルブロ集合の幅は2の倍数のピクセルとなるようにしてある。また、UVAは使わなくても良かったのだが、使い方を示すために利用している。

CUDA 4.0によるマルチGPU化によりマンデルブロ集合計算の実行時間(計算条件は前回と同じ)は7.45秒から3.73秒となった。ほぼ2倍の速度である。シングルGPUからの修正箇所はほんの僅か。カーネル関数に至っては一文字たりとも変更していない。良い時代になったものだ。

mandelbrot_multigpu.cu

// madelbrot using multi GPU devices by nox, 2011.06.06 // nvcc -lcutil_x86_64 -arch sm_13 -use_fast_math -prec-sqrt=false -keep -L ~/NVIDIA_GPU_Computing_SDK/C/lib -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc -g mandelbrot_multigpu.cu -o mandelbrot_multigpu #include <iostream> #include <fstream> #include "cutil_inline.h" using namespace std; const int BLOCK_SIZE_X = 16; const int BLOCK_SIZE_Y = 16; __device__ uchar4 coloring(int n, int b) { int d = n % b; d *= 256 / b; int m = static_cast<int>(d / 42.667f); switch (m) { case 0: return make_uchar4(0, 6 * d, 255, 0); case 1: return make_uchar4(0, 255, 255 - 6 * (d - 43), 0); case 2: return make_uchar4(6 * (d - 86), 255, 0, 0); case 3: return make_uchar4(255, 255 - 6 * (d - 129), 0, 0); case 4: return make_uchar4(255, 0, 6 * (d - 171), 0); case 5: return make_uchar4(255 - 6 * (d - 214), 0, 255, 0); default: return make_uchar4(0, 0, 0, 0); } } __global__ void mandelbrot(float t, float l, float w, float h, int sw, int sh, int max_iter, float th, uchar4* d_color) { int ix = blockIdx.x * blockDim.x + threadIdx.x; int iy = blockIdx.y * blockDim.y + threadIdx.y; if (ix >= sw || iy >= sh) return; float ci = t + (static_cast<float>(iy) / sh) * h; float cr = l + (static_cast<float>(ix) / sw) * w; float zi = 0.0f; float zr = 0.0f; float zrzi, zr2, zi2; for (int i = 0; i < max_iter; i++) { zrzi = zr * zi; zr2 = zr * zr; zi2 = zi * zi; zr = zr2 - zi2 + cr; zi = zrzi + zrzi + ci; if (zi2 + zr2 >= th) { d_color[ix*sh+iy] = coloring(i, 256); return; } } d_color[ix*sh+iy] = make_uchar4(0, 0, 0, 0); } void write_mandelbrot(const string outfile, float t, float l, float w, float h, int sw, int sh, int max_iter=256, float th=2.0f) { dim3 num_blocks((sw - 1) / BLOCK_SIZE_X + 1, (sh - 1) / BLOCK_SIZE_Y + 1, 1); dim3 num_threads(BLOCK_SIZE_X, BLOCK_SIZE_Y, 1); uchar4* h_color; uchar4* d_color[2]; cudaSetDevice(0); cudaMallocHost((void**)&h_color, sizeof(uchar4) * sw * sh); for (int i = 0; i < 2; i++) { cudaSetDevice(i); cudaMalloc((void**)&d_color[i], sizeof(uchar4) * sw * sh / 2); } unsigned int timer; cutCreateTimer(&timer); cutStartTimer(timer); for (int i = 0; i < 2; i++) { cudaSetDevice(i); mandelbrot<<<num_blocks, num_threads>>>(t, l + (w / 2) * i, w / 2, h, sw / 2, sh, max_iter, th, d_color[i]); } cudaThreadSynchronize(); cutStopTimer(timer); cout << "Time: " << cutGetTimerValue(timer) / 1000 << endl; for (int i = 0; i < 2; i++) cudaMemcpy(h_color + sw * sh / 2 * i, d_color[i], sizeof(uchar4) * sw * sh / 2, cudaMemcpyDefault); fstream fs(outfile.c_str(), ios_base::out); fs << sw << endl; fs << sh << endl; for (int i = 0; i < sw * sh; i++) fs << h_color[i].x << h_color[i].y << h_color[i].z; fs.close(); for (int i = 0; i < 2; i++) { cudaSetDevice(i); cudaFree(d_color[i]); } cudaSetDevice(0); cudaFreeHost(h_color); } int main(int argc, char* argv[]) { string outfile("mandelbrot.dat"); if (argc >= 2) outfile = argv[1]; // for UVA cudaSetDevice(0); cudaDeviceEnablePeerAccess(1, 0); cudaSetDevice(1); cudaDeviceEnablePeerAccess(0, 0); // write_mandelbrot(output filename, // top position, // left position, // width, // height, // screen width, // screen height, // max number of iterations [256] // threshold [2.0]); //write_mandelbrot(outfile, -1.0f, -1.5f, 2.0f, 2.0f, 480, 480, 100); //write_mandelbrot(outfile, 0.57f, 0.34f, 0.036f, 0.027f, 6400, 4800, 1000); write_mandelbrot(outfile, -0.005f, -0.005f, 0.01f, 0.01f, 4800, 4800, 10000); return 0; }

出力ファイルの画像への変換は複数のGPUでマンデルブロ集合を並列計算の記事に記載されているread_mandelbrot.pyが利用できる。