12. OpenCL+MPI

12.1 OpenCLとMPIの併用

OpenCLとMPIを併用すると、複数個のNVIDIAまたはAMDのグラフィックスボードを用いて並列計算することができます。
1台のコンピュータが複数個のグラフィックスボードを持つ場合(マルチGPU)と、 複数台のコンピュータが1個以上のグラフィックスボードを持つ場合に対応することができます。
グラフィックスボードではOpenCLを用いて高速に計算を行い、プロセス間の通信はMPIを用いて行います。
プログラミング方法と作業手順はCUDA+MPIと同じですので11.を参考にしてください。

12.2 OpenCL+MPIプログラミング例

リスト12-1にベクトルの和をOpenCL+MPIで並列計算するプログラムを示します。
なお、カーネルプログラムvadd.clはリスト9-2と同じです。


リスト12-1 OpenCL+MPIプログラム(vadd_ocl_mpi.c)
     1	/*
     2	add two vectors (OpenCL + MPI)
     3	
     4	Compile + Link:
     5	> cl /Ox vadd_ocl_mpi.c OpenCL.lib msmpi.lib
     6	
     7	Usage:
     8	> mpiexec -n <proc> vadd_ocl_mpi <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, nloop=%d, nproc=%d, %e(%.6e), cpu[sec]=%.3f\n",
   168				n, nloop, comm_size, sum, exact, cpu);
   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.3 OpenCL+MPIの計算時間

表12-1に計算時間を示します。
1台のコンピュータに2個のGPUが実装されている場合です。
2GPUで計算すると1GPUのほぼ2倍速くなります。

表12-1 OpenCL+MPIの計算時間(Windows, GTX 1070 X 2, ()内は1GPUとの速度比)
No.配列の大きさN繰り返し回数L1GPU2GPU
1100,000,0001,0006.10秒 (1.0)3.17秒 (1.92)
210,000,00010,0005.88秒 (1.0)2.98秒 (1.97)