10. CUDA+マルチGPU

10.1 CUDA+マルチGPUプログラミング

1台のコンピュータに複数のグラフィックスボードが実装されている環境をマルチGPUと呼びます。
NVIDIAのグラフィックスボードでCUDAを用いて計算する方法にはいくつかありますが、 ここではstreamを使用する方法とOpenMPを使用する方法を説明します。

10.2 CUDA+stream

CUDAの言語仕様にはstreamが含まれています。 これは複数のstreamを同時に起動して計算時間を短縮する技術です。
GPUの数だけのstream配列を作成し、各GPUにcudaSetDevice関数でデバイス番号を与え、 cudaStreamCreate関数でstreamとデバイスを関連付けます。
cudaMalloc関数で配列を作成する前にcudaSetDevice関数でデバイス番号を指定すると、 配列は指定したデバイス番号のビデオメモリーに置かれます。
カーネルを起動するときにexecution configurationの第4引数にstreamを指定します。
各streamは非同期に実行されますので、cudaStreamSynchronize関数で各streamの終了を待つ必要があります。

10.3 CUDA+OpenMP

OpenMPを使用して複数のスレッドを起動し、 cudaSetDevice関数を用いてスレッドとデバイスを関連付けます。
streamと同様、cudaMalloc関数で配列を作成する前にcudaSetDevice関数でデバイス番号を指定すると、 配列は指定したデバイス番号のビデオメモリーに置かれます。
streamに比べるとstreamの生成、同期、廃棄が不要になりプログラムは少し簡単になります。
ただし、マルチGPUプログラミングの作業の中心は計算をスレッドで分解する所であり、 その手間はstreamとOpenMPで変わりません。

10.4 CUDA+マルチGPUプログラム例

リスト10-1にベクトルの和をCUDA+マルチGPUで計算するプログラムを示します。
streamとOpenMPのプログラムには共通点が多いのでここでは一つのプログラムに記述し、 コンパイルオプションで使い分けています。


リスト10-1 CUDA+マルチGPUプログラム (vadd_cuda_multi.cu)
     1	/*
     2	add two vectors (CUDA + multiGPU : stream or OpenMP)
     3	
     4	Compile,Link:
     5	> nvcc -O3 -o vadd_cuda_stream vadd_cuda_multi.cu cuda_memory.cu
     6	> nvcc -O3 -Xcompiler /openmp -o vadd_cuda_omp vadd_cuda_multi.cu cuda_memory.cu
     7	
     8	Usage:
     9	> vadd_cuda_stream [-gpu|-cpu] [-hdm|-um] <ngpu> <n> <loop>
    10	> vadd_cuda_omp [-gpu|-cpu] [-hdm|-um] <ngpu> <n> <loop>
    11	*/
    12	
    13	// GPU/CPU
    14	__host__ __device__
    15	static void vadd_calc(float a, float b, float *c)
    16	{
    17		*c = a + b;
    18	}
    19	
    20	// GPU
    21	__global__
    22	static void vadd_gpu(int n, const float *a, const float *b, float *c)
    23	{
    24		int tid = threadIdx.x + (blockIdx.x * blockDim.x);
    25		if (tid < n) {
    26			vadd_calc(a[tid], b[tid], &c[tid]);
    27		}
    28	}
    29	
    30	// CPU
    31	static void vadd_cpu(int n, const float *a, const float *b, float *c)
    32	{
    33		for (int i = 0; i < n; i++) {
    34			vadd_calc(a[i], b[i], &c[i]);
    35		}
    36	}
    37	
    38	#include <stdlib.h>
    39	#include <stdio.h>
    40	#include <string.h>
    41	#include <time.h>
    42	#ifdef _OPENMP
    43	#include <omp.h>
    44	#endif
    45	
    46	static void divide_array(const int, const int, int [], int []);
    47	extern void cuda_malloc(int, int, void **, size_t);
    48	extern void cuda_free(int, void *);
    49	extern void cuda_memcpy(int, void *, const void *, size_t, cudaMemcpyKind);
    50	
    51	int main(int argc, char **argv)
    52	{
    53		int    gpu = 1;
    54		int    um = 0;
    55		int    ngpu = 1;
    56		int    nvector = 1000;
    57		int    nloop = 1000;
    58		int    *offset, *length;
    59		size_t size;
    60	#ifndef _OPENMP
    61		cudaStream_t *stream;
    62	#endif
    63	
    64		// arguments
    65		while (--argc) {
    66			argv++;
    67			if      (!strcmp(*argv, "-gpu")) {
    68				gpu = 1;
    69			}
    70			else if (!strcmp(*argv, "-cpu")) {
    71				gpu = 0;
    72			}
    73			else if (!strcmp(*argv, "-hdm")) {
    74				um = 0;
    75			}
    76			else if (!strcmp(*argv, "-um")) {
    77				um = 1;
    78			}
    79			else if (argc == 3) {
    80				ngpu = atoi(*argv);
    81				if (ngpu < 1) ngpu = 1;
    82			}
    83			else if (argc == 2) {
    84				nvector = atoi(*argv);
    85				if (nvector < 1) nvector = 1;
    86			}
    87			else if (argc == 1) {
    88				nloop = atoi(*argv);
    89				if (nloop < 1) nloop = 1;
    90			}
    91		}
    92	
    93		// setup GPU
    94		if (gpu) {
    95			// check device
    96			int ndevice;
    97			cudaGetDeviceCount(&ndevice);
    98			if (ndevice < 1) {
    99				printf("No CUDA device.\n");
   100				exit(1);
   101			}
   102	
   103			// GPU info
   104			for (int igpu = 0; igpu < ngpu; igpu++) {
   105				cudaDeviceProp prop;
   106				cudaGetDeviceProperties(&prop, igpu);
   107				printf("GPU-%d : %s, C.C.%d.%d, U.M.%s\n", igpu, prop.name, prop.major, prop.minor, (um ? "ON" : "OFF"));
   108			}
   109	
   110	#ifndef _OPENMP
   111			// streams
   112			size = ngpu * sizeof(cudaStream_t);
   113			stream = (cudaStream_t *)malloc(size);
   114			for (int igpu = 0; igpu < ngpu; igpu++) {
   115				cudaSetDevice(igpu);  // 重要
   116				cudaStreamCreate(&stream[igpu]);
   117			}
   118	#endif
   119		}
   120	
   121		// host memory
   122		size = nvector * sizeof(float);
   123		float *h_a = (float *)malloc(size);
   124		float *h_b = (float *)malloc(size);
   125		float *h_c = (float *)malloc(size);
   126	
   127		// offset and length
   128		size = ngpu * sizeof(int);
   129		offset = (int *)malloc(size);
   130		length = (int *)malloc(size);
   131		divide_array(nvector, ngpu, offset, length);
   132	
   133		// device memory
   134		size = ngpu * sizeof(float *);
   135		float **d_a = (float **)malloc(size);
   136		float **d_b = (float **)malloc(size);
   137		float **d_c = (float **)malloc(size);
   138		for (int igpu = 0; igpu < ngpu; igpu++) {
   139			if (gpu) cudaSetDevice(igpu);  // 重要
   140			size = length[igpu] * sizeof(float);
   141			cuda_malloc(gpu, um, (void **)&d_a[igpu], size);
   142			cuda_malloc(gpu, um, (void **)&d_b[igpu], size);
   143			cuda_malloc(gpu, um, (void **)&d_c[igpu], size);
   144		}
   145	
   146		// setup problem
   147		for (int i = 0; i < nvector; i++) {
   148			h_a[i] = i;
   149			h_b[i] = i + 1;
   150			h_c[i] = 0;
   151		}
   152	
   153		// copy host to device
   154		for (int igpu = 0; igpu < ngpu; igpu++) {
   155			size = length[igpu] * sizeof(float);
   156			cuda_memcpy(gpu, d_a[igpu], h_a + offset[igpu], size, cudaMemcpyHostToDevice);
   157			cuda_memcpy(gpu, d_b[igpu], h_b + offset[igpu], size, cudaMemcpyHostToDevice);
   158		}
   159	
   160		// timer
   161		for (int igpu = 0; igpu < ngpu; igpu++) {
   162			cudaSetDevice(igpu);
   163			cudaDeviceSynchronize();
   164		}
   165		clock_t t0 = clock();
   166	
   167		// calculation
   168		for (int loop = 0; loop < nloop; loop++) {
   169			if (gpu) {
   170				// GPU
   171	#ifdef _OPENMP
   172				// OpenMP
   173				omp_set_num_threads(ngpu);
   174	#pragma omp parallel
   175				{
   176					int igpu = omp_get_thread_num();
   177					cudaSetDevice(igpu);
   178					int block = 256;
   179					int grid = (length[igpu] + (block - 1)) / block;
   180					vadd_gpu<<<grid, block>>>(length[igpu], d_a[igpu], d_b[igpu], d_c[igpu]);
   181				}
   182	#else
   183				// streams
   184				for (int igpu = 0; igpu < ngpu; igpu++) {
   185					cudaSetDevice(igpu);
   186					int block = 256;
   187					int grid = (length[igpu] + (block - 1)) / block;
   188					vadd_gpu<<<grid, block, 0, stream[igpu]>>>(length[igpu], d_a[igpu], d_b[igpu], d_c[igpu]);
   189				}
   190	
   191				// synchronize streams
   192				for (int igpu = 0; igpu < ngpu; igpu++) {
   193					cudaSetDevice(igpu);
   194					cudaStreamSynchronize(stream[igpu]);
   195				}
   196	#endif
   197			}
   198			else {
   199				// CPU
   200				for (int igpu = 0; igpu < ngpu; igpu++) {
   201					vadd_cpu(length[igpu], h_a + offset[igpu], h_b + offset[igpu], h_c + offset[igpu]);
   202				}
   203			}
   204		}
   205	
   206		// timer
   207		for (int igpu = 0; igpu < ngpu; igpu++) {
   208			cudaSetDevice(igpu);
   209			cudaDeviceSynchronize();
   210		}
   211		clock_t t1 = clock();
   212	
   213		// copy device to host
   214		for (int igpu = 0; igpu < ngpu; igpu++) {
   215			cuda_memcpy(gpu, h_c + offset[igpu], d_c[igpu], size, cudaMemcpyDeviceToHost);
   216		}
   217	
   218		// sum
   219		double sum = 0;
   220		for (int i = 0; i < nvector; i++) {
   221			sum += h_c[i];
   222		}
   223	
   224		// output
   225		double exact = (double)nvector * nvector;
   226		double sec = (double)(t1 - t0) / CLOCKS_PER_SEC;
   227		printf("nvector=%d nloop=%d %e(%e) %s[sec]=%.3f\n",
   228			nvector, nloop, sum, exact, (gpu ? "GPU" : "CPU"), sec);
   229	
   230		// free
   231	#ifndef _OPENMP
   232		if (gpu) {
   233			for (int igpu = 0; igpu < ngpu; igpu++) {
   234				cudaStreamDestroy(stream[igpu]);
   235			}
   236			free(stream);
   237		}
   238	#endif
   239		for (int igpu = 0; igpu < ngpu; igpu++) {
   240			cuda_free(gpu, d_a[igpu]);
   241			cuda_free(gpu, d_b[igpu]);
   242			cuda_free(gpu, d_c[igpu]);
   243		}
   244		free(h_a);
   245		free(h_b);
   246		free(h_c);
   247		free(d_a);
   248		free(d_b);
   249		free(d_c);
   250		cudaDeviceReset();
   251	
   252		return 0;
   253	}
   254	
   255	static void divide_array(const int n, const int ndiv, int offset[], int length[])
   256	{
   257		int l_n = (n + (ndiv - 1)) / ndiv;
   258		for (int i = 0; i < ndiv; i++) {
   259			offset[i] = i * l_n;
   260			int istop = (i + 1) * l_n;
   261			if (istop > n) istop = n;
   262			length[i] = istop - offset[i];
   263		}
   264	}

ソースコードの説明
112-117行目: streamではstream配列を生成します。
139行目: デバイスメモリーに配列を作成する前にその配列が存在するデバイス番号を指定することが必要です。
173行目: OpenMPではスレッド数を設定します。(既定値は論理スレッド数)
174-181行目: OpenMPでは"#pragma omp parallel"指示文で並列領域を作成し、 その中でomp_get_thread_num関数でスレッド番号を取得し、 cudaSetDevice関数でスレッド番号とデバイス番号を関連付けます。
184-189行目: streamではGPUの数のfor文の中で、 cudaSetDevice関数でstream番号とデバイス番号を関連付けます。
execution configurationの第4引数にstreamを指定します。
192-195行目: streamは非同期で実行されるために、 cudaStreamSynchronize関数でカーネルの終了を待機します。
255-264行目: 自作関数divide_arrayで配列をブロック分割したときのオフセットと長さを計算します。

コンパイル・リンク方法
コンパイル・リンク方法は以下の通りです。
streamのとき:
> nvcc -O3 -o vadd_cuda_stream -arch=sm_30 vadd_cuda_multi.cu cuda_memory.cu
OpenMPのとき:
> nvcc -O3 -Xcompiler "/openmp" -o vadd_cuda_omp -arch=sm_30 vadd_cuda_multi.cu cuda_memory.cu
VC++では多数の"warning C4819"が出ることがあります。 そのときは以下のようにコンパイルオプションを追加してください。
> nvcc -O3 -o vadd_cuda_stream -Xcompiler "/wd4819" -arch=sm_30 vadd_cuda_multi.cu cuda_memory.cu
> nvcc -O3 -Xcompiler "/openmp /wd4819" -o vadd_cuda_omp -arch=sm_30 vadd_cuda_multi.cu cuda_memory.cu

プログラムの実行方法
プログラムの実行方法は以下の通りです。
streamのとき:
> vadd_cuda_stream [-gpu|-cpu] [-hdm|-um] GPU数 配列の大きさ 繰り返し回数
OpenMPのとき:
> vadd_cuda_omp [-gpu|-cpu] [-hdm|-um] GPU数 配列の大きさ 繰り返し回数
例えば以下のようになります。
> vadd_cuda_stream 2 100000000 1000 (2GPU+streamで計算するとき)
> vadd_cuda_omp 2 100000000 1000 (2GPU+OpenMPで計算するとき)
> vadd_cuda_stream -um 2 100000000 1000 (GPUのunified memoryで計算するとき)
> vadd_cuda_stream -cpu 2 100000000 1000 (CPUで計算するとき)
繰り返し回数は計算時間の測定誤差を小さくするためです。

10.5 CUDA+マルチGPUの計算時間

表10-1に計算時間を示します。
No.1のようにカーネル起動時間が計算時間に比べて十分小さいときは2GPUは1GPUの2倍速くなります。
No.3,No.4ではstreamまたはOpenMPのスレッド起動時間も加わりさらに遅くなります。 スレッド起動時間について見るとOpenMPはstreamのおよそ半分になっています。
カーネル起動時間とスレッド起動時間がネックにならないようにプログラムを設計することが大切です。

表10-1 CUDA+マルチGPUの計算時間(Windows, ()内は1GPUとの速度比)
No.配列の大きさN繰り返し回数L1GPU2GPU(stream)2GPU(OpenMP)
1100,000,0001,0005.78秒 (1.0)2.95秒 (1.96)2.90秒 (1.99)
210,000,00010,0005.80秒 (1.0)3.38秒 (1.72)2.98秒 (1.95)
31,000,000100,0006.00秒 (1.0)6.71秒 (0.86)3.75秒 (1.55)
4100,0001,000,0003.46秒 (1.0)39.19秒 (0.09)22.47秒 (0.15)