1台のコンピュータに複数のグラフィックスボードが実装されている環境をマルチGPUと呼びます。
NVIDIAのグラフィックスボードでCUDAを用いて計算する方法にはいくつかありますが、
ここではstreamを使用する方法とOpenMPを使用する方法を説明します。
CUDAの言語仕様にはstreamが含まれています。
これは複数のstreamを同時に起動して計算時間を短縮する技術です。
GPUの数だけのstream配列を作成し、各GPUにcudaSetDevice関数でデバイス番号を与え、
cudaStreamCreate関数でstreamとデバイスを関連付けます。
cudaMalloc関数で配列を作成する前にcudaSetDevice関数でデバイス番号を指定すると、
配列は指定したデバイス番号のビデオメモリーに置かれます。
カーネルを起動するときにexecution configurationの第4引数にstreamを指定します。
各streamは非同期に実行されるので、cudaStreamSynchronize関数で各streamの終了を待つ必要があります。
OpenMPを使用して複数のスレッドを起動し、
cudaSetDevice関数を用いてスレッドとデバイスを関連付けます。
streamと同様、cudaMalloc関数で配列を作成する前にcudaSetDevice関数でデバイス番号を指定すると、
配列は指定したデバイス番号のビデオメモリーに置かれます。
streamに比べるとstreamの生成、同期、廃棄が不要になりプログラムは少し簡単になります。
ただし、マルチGPUプログラミングの作業の中心は計算をスレッドで分解する所であり、
その手間はstreamとOpenMPで変わりません。
リスト10-1にベクトル和をCUDA+マルチGPUで計算するプログラムを示します。
streamとOpenMPのプログラムには共通点が多いのでここでは一つのプログラムに記述し、
コンパイルオプションで使い分けています。
リスト10-1 CUDA+マルチGPUプログラム (cuda_multi_vadd.cu)
1 /*
2 cuda_multi_vadd.cu (CUDA + multiGPU : stream or OpenMP)
3
4 Compile,Link:
5 > nvcc -O2 -o cuda_stream_vadd cuda_multi_vadd.cu cuda_memory.cu
6 > nvcc -O2 -Xcompiler "/openmp" -o cuda_omp_vadd cuda_multi_vadd.cu cuda_memory.cu
7
8 Usage:
9 > cuda_stream_vadd [-gpu|-cpu] [-hdm|-um] <ngpu> <n> <loop>
10 > cuda_omp_vadd [-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 // create 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 -O2 -o cuda_stream_vadd cuda_multi_vadd.cu cuda_memory.cu
OpenMPのとき:
> nvcc -O2 -Xcompiler "/openmp" -o cuda_omp_vadd cuda_multi_vadd.cu cuda_memory.cu
VC++では多数の"warning C4819"が出ることがあります。
そのときは以下のようにコンパイルオプションを追加してください。
> nvcc -O2 -o cuda_stream_vadd -Xcompiler "/wd4819" cuda_multi_vadd.cu cuda_memory.cu
> nvcc -O2 -Xcompiler "/openmp /wd4819" -o cuda_omp_vadd cuda_multi_vadd.cu cuda_memory.cu
プログラムの実行方法
プログラムの実行方法は以下の通りです。
streamのとき:
> cuda_stream_vadd [-gpu|-cpu] [-hdm|-um] GPU数 配列の大きさ 繰り返し回数
OpenMPのとき:
> cuda_omp_vadd [-gpu|-cpu] [-hdm|-um] GPU数 配列の大きさ 繰り返し回数
例えば以下のようになります。
> cuda_stream_vadd 2 100000000 1000 (2GPU+streamで計算するとき)
> cuda_omp_vadd 2 100000000 1000 (2GPU+OpenMPで計算するとき)
> cuda_stream_vadd -um 2 100000000 1000 (GPUのunified memoryで計算するとき)
> cuda_stream_vadd -cpu 1 100000000 1000 (CPUで計算するとき)
繰り返し回数は計算時間の測定誤差を小さくするためです。
表10-1に計算時間を示します。
配列の大きさ(=N)と繰り返し回数(=L)の積は一定(=1011)です。
従って全体の演算量は同じです。
本環境は1GPUで2個のstreamやOpenMPスレッドを起動しているので、
1GPUと2GPUの計算時間が変わらなければ、
2GPUが実装された環境では1GPUの2倍速くなることが予想できます。
No.1-2では1GPUと2GPUの計算時間は同じです。
No.3-4ではstreamまたはOpenMPスレッドを起動するオーバーヘッドのために2GPUは1GPUより遅くなっています。
No.1-4のすべてのケースの2GPUでstreamとOpenMPの計算時間が同じです。
プログラミングの手間もほぼ同じです。
すなわち、マルチGPU環境では、streamとOpenMPのどちらを使ってもかまいません。
マルチGPU環境では、CUDAカーネルの起動時間に加えて、
スレッドの起動時間がネックにならないようにプログラムを設計することが大切です。
| No. | 配列の大きさN | 繰り返し回数L | 1GPU | 2GPU(stream) | 2GPU(OpenMP) |
|---|---|---|---|---|---|
| 1 | 100,000,000 | 1,000 | 4.98秒 | 4.99秒 | 4.98秒 |
| 2 | 10,000,000 | 10,000 | 5.00秒 | 5.02秒 | 5.02秒 |
| 3 | 1,000,000 | 100,000 | 5.20秒 | 5.48秒 | 5.46秒 |
| 4 | 100,000 | 1,000,000 | 6.77秒 | 53.62秒 | 53.24秒 |