OpenCLとMPIを併用すると、複数個のNVIDIAまたはAMDのグラフィックスボードを用いて並列計算することができます。
1台のコンピュータが複数個のグラフィックスボードを持つ場合(マルチGPU)と、
複数台のコンピュータが1個以上のグラフィックスボードを持つ場合に対応することができます。
グラフィックスボードではOpenCLを用いて高速に計算を行い、プロセス間の通信はMPIを用いて行います。
プログラミング方法と作業手順はCUDA+MPIと同じなので11.も参考にしてください。
リスト12-1にベクトルの和をOpenCL+MPIで並列計算するプログラムを示します。
なお、カーネルプログラムvadd.clはリスト9-2と同じです。
リスト12-1 OpenCL+MPIプログラム(ocl_mpi_vadd.c)
1 /*
2 ocl_mpi_vadd.c (OpenCL + MPI)
3
4 Compile + Link:
5 > cl.exe /O2 ocl_mpi_vadd.c OpenCL.lib msmpi.lib
6
7 Usage:
8 > mpiexec -n <proc> ocl_mpi_vadd <n> <loop> <platform> <device> ...
9 */
10
11 #include <stdio.h>
12 #include <stdlib.h>
13 #include <string.h>
14 #include <time.h>
15
16 #ifdef __APPLE__
17 #include <OpenCL/opencl.h>
18 #else
19 #include <CL/cl.h>
20 #endif
21
22 #ifdef MPI
23 #include <mpi.h>
24 #endif
25
26 #define MAX_PLATFORMS (10)
27 #define MAX_DEVICES (10)
28 #define MAX_SOURCE_SIZE (100000)
29
30 static int setup_ocl(cl_uint, cl_uint, char *);
31 static void vadd_calc(void);
32 static void vadd(void);
33
34 cl_command_queue Queue;
35 cl_kernel k_vadd;
36 int N;
37 float *h_A, *h_B, *h_C;
38 cl_mem d_A, d_B, d_C;
39 int OCL;
40
41 int main(int argc, char **argv)
42 {
43 int *platform = NULL;
44 int *device = NULL;
45 int comm_size = 1;
46 int comm_rank = 0;
47 int n = 1000;
48 int nloop = 1000;
49
50 clock_t t0 = 0, t1 = 0;
51 #ifdef MPI
52 MPI_Status status;
53 int g_ret;
54 #endif
55
56 // initialize (MPI)
57 #ifdef MPI
58 MPI_Init(&argc, &argv);
59 MPI_Comm_size(MPI_COMM_WORLD, &comm_size);
60 MPI_Comm_rank(MPI_COMM_WORLD, &comm_rank);
61 #endif
62
63 // arguments
64 platform = (int *)calloc(comm_size, sizeof(int));
65 device = (int *)calloc(comm_size, sizeof(int));
66 if (argc >= 3 + 2 * comm_size) {
67 n = atoi(argv[1]);
68 nloop = atoi(argv[2]);
69 for (int i = 0; i < comm_size; i++) {
70 platform[i] = atoi(argv[3 + 2 * i]);
71 device[i] = atoi(argv[4 + 2 * i]);
72 }
73 }
74 OCL = (platform[0] >= 0);
75
76 // local size : N
77 N = (n + (comm_size - 1)) / comm_size;
78
79 // alloc host memory
80 size_t size = N * sizeof(float);
81 h_A = (float *)malloc(size);
82 h_B = (float *)malloc(size);
83 h_C = (float *)malloc(size);
84
85 // setup problem
86 for (int i = 0; i < N; i++) {
87 int gid = (comm_rank * N) + i;
88 if (gid < n) {
89 h_A[i] = gid + 1.0f;
90 h_B[i] = gid + 1.0f;
91 }
92 }
93
94 // setup OpenCL
95 if (OCL) {
96 char msg[BUFSIZ];
97 int ret = setup_ocl((cl_uint)platform[comm_rank], (cl_uint)device[comm_rank], msg);
98 #ifdef MPI
99 // show device name
100 if (comm_rank == 0) {
101 printf("%d: %s\n", 0, msg);
102 for (int i = 1; i < comm_size; i++) {
103 MPI_Recv(msg, BUFSIZ, MPI_CHAR, i, 0, MPI_COMM_WORLD, &status);
104 printf("%d: %s\n", i, msg);
105 }
106 fflush(stdout);
107 }
108 else {
109 MPI_Send(msg, BUFSIZ, MPI_CHAR, 0, 0, MPI_COMM_WORLD);
110 }
111 // error check
112 MPI_Allreduce(&ret, &g_ret, 1, MPI_INT, MPI_LOR, MPI_COMM_WORLD);
113 if (g_ret) {
114 MPI_Finalize();
115 exit(1);
116 }
117 #else
118 printf("%s\n", msg);
119 if (ret) {
120 exit(1);
121 }
122 #endif
123 }
124
125 // timer
126 if (comm_rank == 0) {
127 t0 = clock();
128 }
129
130 // copy host to device
131 if (OCL) {
132 clEnqueueWriteBuffer(Queue, d_A, CL_TRUE, 0, size, h_A, 0, NULL, NULL);
133 clEnqueueWriteBuffer(Queue, d_B, CL_TRUE, 0, size, h_B, 0, NULL, NULL);
134 }
135
136 // run
137 for (int loop = 0; loop < nloop; loop++) {
138 vadd_calc();
139 }
140
141 // copy device to host
142 if (OCL) {
143 clEnqueueReadBuffer(Queue, d_C, CL_TRUE, 0, size, h_C, 0, NULL, NULL);
144 }
145
146 // timer
147 if (comm_rank == 0) {
148 t1 = clock();
149 }
150
151 // sum
152 double sum = 0;
153 for (int i = 0; i < N; i++) {
154 sum += h_C[i];
155 }
156
157 // reduction (MPI)
158 #ifdef MPI
159 double l_sum = sum;
160 MPI_Reduce(&l_sum, &sum, 1, MPI_DOUBLE, MPI_SUM, 0, MPI_COMM_WORLD);
161 #endif
162
163 // output
164 if (comm_rank == 0) {
165 double exact = n * (n + 1.0);
166 double cpu = (double)(t1 - t0) / CLOCKS_PER_SEC;
167 printf("N=%d L=%d %.6e(%.6e) err=%.1e %dGPUs %.3f[sec]\n",
168 n, nloop, sum, exact, fabs((sum - exact) / exact), comm_size, sec);
169 fflush(stdout);
170 }
171
172 #ifdef MPI
173 MPI_Finalize();
174 #endif
175
176 // release
177 if (OCL) {
178 clReleaseMemObject(d_A);
179 clReleaseMemObject(d_B);
180 clReleaseMemObject(d_C);
181 clReleaseKernel(k_vadd);
182 clReleaseCommandQueue(Queue);
183 }
184
185 // free
186 free(h_A);
187 free(h_B);
188 free(h_C);
189
190 return 0;
191 }
192
193 // setup OpenCL
194 static int setup_ocl(cl_uint platform, cl_uint device, char *msg)
195 {
196 cl_context context = NULL;
197 cl_program program = NULL;
198 cl_platform_id platform_id[MAX_PLATFORMS];
199 cl_device_id device_id[MAX_DEVICES];
200
201 FILE *fp;
202 char *source_str;
203 char str[BUFSIZ];
204 size_t source_size, ret_size, size;
205 cl_uint num_platforms, num_devices;
206 cl_int ret;
207
208 // alloc
209 source_str = (char *)malloc(MAX_SOURCE_SIZE * sizeof(char));
210
211 // platform
212 clGetPlatformIDs(MAX_PLATFORMS, platform_id, &num_platforms);
213 if (platform >= num_platforms) {
214 sprintf(msg, "error : platform = %d (limit = %d)", platform, num_platforms - 1);
215 return 1;
216 }
217
218 // device
219 clGetDeviceIDs(platform_id[platform], CL_DEVICE_TYPE_ALL, MAX_DEVICES, device_id, &num_devices);
220 if (device >= num_devices) {
221 sprintf(msg, "error : device = %d (limit = %d)", device, num_devices - 1);
222 return 1;
223 }
224
225 // device name (option)
226 clGetDeviceInfo(device_id[device], CL_DEVICE_NAME, sizeof(str), str, &ret_size);
227 sprintf(msg, "%s (platform = %d, device = %d)", str, platform, device);
228
229 // context
230 context = clCreateContext(NULL, 1, &device_id[device], NULL, NULL, &ret);
231
232 // command queue
233 Queue = clCreateCommandQueue(context, device_id[device], 0, &ret);
234
235 // source
236 if ((fp = fopen("vadd.cl", "r")) == NULL) {
237 sprintf(msg, "kernel source open error");
238 return 1;
239 }
240 source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
241 fclose(fp);
242
243 // program
244 program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
245 if (ret != CL_SUCCESS) {
246 sprintf(msg, "clCreateProgramWithSource() error");
247 return 1;
248 }
249
250 // build
251 if (clBuildProgram(program, 1, &device_id[device], NULL, NULL, NULL) != CL_SUCCESS) {
252 sprintf(msg, "clBuildProgram() error");
253 exit(1);
254 }
255
256 // kernel
257 k_vadd = clCreateKernel(program, "vadd", &ret);
258 if (ret != CL_SUCCESS) {
259 sprintf(msg, "clCreateKernel() error");
260 return 1;
261 }
262
263 // memory object
264 size = N * sizeof(float);
265 d_A = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &ret);
266 d_B = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &ret);
267 d_C = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &ret);
268
269 // release
270 clReleaseProgram(program);
271 clReleaseContext(context);
272
273 // free
274 free(source_str);
275
276 return 0;
277 }
278
279 // entry point
280 static void vadd_calc(void)
281 {
282 if (OCL) {
283 size_t global_item_size, local_item_size;
284
285 // args
286 clSetKernelArg(k_vadd, 0, sizeof(cl_mem), (void *)&d_A);
287 clSetKernelArg(k_vadd, 1, sizeof(cl_mem), (void *)&d_B);
288 clSetKernelArg(k_vadd, 2, sizeof(cl_mem), (void *)&d_C);
289 clSetKernelArg(k_vadd, 3, sizeof(int), (void *)&N);
290
291 // work item
292 local_item_size = 256;
293 global_item_size = ((N + local_item_size - 1) / local_item_size)
294 * local_item_size;
295
296 // run
297 clEnqueueNDRangeKernel(Queue, k_vadd, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL);
298 }
299 else {
300 // non-OpenCL code
301 vadd();
302 }
303 }
304
305 // non-OpenCL code
306 static void vadd(void)
307 {
308 for (int i = 0; i < N; i++) {
309 h_C[i] = h_A[i] + h_B[i];
310 }
311 }
ソースコードの説明
77行目:全体の問題サイズnから各プロセスの問題サイズNを求めます。
80-83行目:各プロセスは部分配列を持ちます。
これによってプロセス数を増やせばそれに比例して扱える問題サイズが大きくなります。
87行目:各プロセスのランクと問題サイズから全体のインデックスを求めます。
コンパイル・リンク方法
コンパイル・リンク方法は以下の通りです。
> cl /Ox /DMPI vadd_ocl_mpi.c OpenCL.lib msmpi.lib
プログラムの実行方法
プログラムの実行方法は以下の通りです。
1台のコンピュータに複数のグラフィックスボードがあるとき:
> mpiexec -n GPU数 vadd_ocl_mpi 配列の大きさ 繰り返し回数 platform番号 device番号 ...
"platform番号 device番号"はGPUの数だけ必要です。
例えば以下のようになります。
> mpiexec -n 2 vadd_ocl_mpi 100000000 1000 0 0 0 1 (2GPUで計算するとき)
複数台のコンピュータを使用するとき:
> mpiexec -hosts ホスト数 ホスト名 プロセス数 ... vadd_ocl_mpi 配列の大きさ 繰り返し回数 platform番号 device番号 ...
"ホスト名 プロセス数"はホストの数だけ必要であり、
"platform番号 device番号"は全プロセスの数だけ必要です。
例えば以下のようになります。
> mpiexec -hosts 2 localhost 1 PC2 1 vadd_ocl_mpi 100000000 1000 0 0 0 0 (2台で計算するとき)
繰り返し回数は計算時間の測定誤差を小さくするためです。
platform番号とdevice番号は環境によって異なるので、
メッセージを見て適当に設定してください。
mpiexecの使い方と実行環境の設定については7.5を参考にしてください。
表12-1にベクトル和の計算時間を示します。
配列の大きさ(=N)と繰り返し回数(=L)の積は一定(=1011)です。
従って全体の演算量は同じです。
本環境は1GPUで1～2個のプロセスを起動しているので、
1プロセスと2プロセスの計算時間が変わらなければ、
2GPUが実装された環境では1GPUの2倍速くなることが予想できます。
No.1-3では2プロセスの計算時間は1プロセスより少し大きいので、
2GPUでは2倍弱速くなると予想できます。
No.4の2プロセスで計算時間が大幅に増えている理由は不明です。
表12-1の2プロセスの計算時間は表11-1の2プロセスとほぼ同じです。
すなわちCUDA+MPIとOpenCL+MPIの性能はほぼ同じと言えます。
| No. | 配列の大きさN | 繰り返し回数L | 1プロセス | 2プロセス |
|---|---|---|---|---|
| 1 | 100,000,000 | 1,000 | 5.22秒 | 6.99秒 |
| 2 | 10,000,000 | 10,000 | 5.02秒 | 6.85秒 |
| 3 | 1,000,000 | 100,000 | 5.21秒 | 6.71秒 |
| 4 | 100,000 | 1,000,000 | 8.97秒 | 20.37秒 |