9. OpenCL

9.1 OpenCLとは

グラフィックスボード(GPU)が搭載されたコンピュータでは、 その高い演算能力を汎用的な科学技術計算に用いることができます。 そのためのプログラミング言語をOpenCL(Open Computing Language)と呼びます。
CPUがフロントエンドとして動作し、計算の主要部はGPUが行います。
OpenCLプログラムはCPUでも動作しますが性能上の利点は少なく主にGPUで使用します。

9.2 OpenCLプログラミング

OpenCLプログラム
OpenCLプログラムは通常のC/C++で記述されるホストプログラムと、 OpenCL C言語で記述されるカーネルプログラム(拡張子.cl)から成ります。
ホストプログラムはCPUで実行され、カーネルプログラムはGPUで実行されます。 両者の関係はCUDAと同じです。

OpenCL C言語
カーネルを記述するOpenCL C言語のC/C++との違いは以下の通りです。

  1. 関数修飾子に"__kernel"が必要です。
  2. カーネル関数の引数にはアドレス空間修飾子 ("__global","__local","__constant","_private"のいずれか)が必要です。 "__"は省略できます。指定しないときは"__private"になります。
  3. "__local"は共有メモリーになります。(CUDAの"__shared__"に対応)
  4. "__constant"は読み取り専用になります。(CUDAの"__constant__"に対応)
  5. 多数の数学関数が用意されています。
  6. ベクタ型(float4,int4など)が使えます。
  7. 環境によってサイズの変わる型(size_tなど)は使えません。
  8. #define文が使えます。
  9. #include文は使えません。従って必要部分をコピーする必要があります。 (コピー元が変更されたときは忘れずに修正することが必要です)

オンラインコンパイルとオフラインコンパイル
OpenCLプログラムにはオンラインコンパイルとオフラインコンパイルの2種類のモードがあります。 それぞれの意味は以下の通りです。

  1. オンラインコンパイル
    実行プログラムと一緒にカーネルのソースコードを配布する必要があります。 実行時にカーネルがコンパイルされます。プログラムの移植性は上がります。
  2. オフラインコンパイル
    開発時にカーネルをコンパイルしてリンクします。 カーネルのソースコードを配布する必要はありませんが、 動作環境ごとに実行プロラムを作成しテストする必要があります。
ここではオンラインコンパイルを使用します。
この場合カーネルに文法エラーがあると実行時にエラーが発生しますが、 エラー発生場所がわからないので開発効率が悪くなります。 そこで、IntelのSDK[16]に付属する"Intel Kernel Builder"を使用します。 図9-1に使用例を示します。カーネルファイルを開いた後、 [Build]→[Compile File]メニューをクリックします。 文法エラーがないときは左のように緑色になり、 文法エラーがあるときは右のように赤色になりエラー発生場所が表示されます。


図9-1 Intel Kernel Builder (左:コンパイル正常、右:コンパイルエラー)

OpenCLプログラミングの指針
OpenCLプログラミングでは以下の点が最も重要です。これはCUDAと同じです。
(1)並列計算できるアルゴリズムを採用する。
(2)CPU/GPU間のデータ転送を最小限にする。
(3)カーネルコードではメモリアクセスをスレッド順とする。
これらを満たさないときは速度は数分の一以下になりGPUを使う意味がなくなります。

変数の命名ルール
CPUコードとGPUコードの両方から呼ばれる同じ意味をもつ変数については、 前者の頭には何もつけないか"h_"(host memoryの意味)を付け、 後者の頭に"d_"(device memoryの意味)を付ける方法がよく用いられます。
このようにすればその変数がCPUにあるかGPUにあるか一目でわかります。
またカーネル変数の名前はカーネル関数の名前と1対1に対応させるとわかりやすくなります。

データ並列とタスク並列
並列処理にはデータ並列とタスク並列の2種類がありますが、数値計算では通常前者を使用します。 OpenCLではそのための関数がclEnqueueNDRangeKernel関数です。

ワークアイテム
OpenCLではスレッドのことをワークアイテムと呼び、スレッドの集合をグループと呼びます。 ワークアイテムの大きさはclEnqueueNDRangeKernel関数の引数で指定します。
表9-1にOpenCLとCUDAの対応関係を示します。 カーネル関数でこれらを通して自分のスレッド番号を取得します。 OpenCLではNo.5,No.6も関数が用意されており便利です。

表9-1 OpenCLとCUDAのスレッド関係の比較
No.OpenCL
ワークアイテム
CUDA
execution configuration
1get_local_id(0)
get_local_id(1)
get_local_id(2)
threadIdx.x
threadIdx.y
threadIdx.z
2get_local_size(0)
get_local_size(1)
get_local_size(2)
blockDim.x
blockDim.y
blockDim.z
3get_group_id(0)
get_group_id(1)
get_group_id(2)
blockIdx.x
blockIdx.y
blockIdx.z
4get_num_groups(0)
get_num_groups(1)
get_num_groups(2)
gridDim.x
gridDim.y
gridDim.z
5get_global_id(0)
get_global_id(1)
get_global_id(2)
threadIdx.x+blockIdx.x*blockDim.x
threadIdx.y+blockIdx.y*blockDim.y
threadIdx.z+blockIdx.z*blockDim.z
6get_global_size(0)
get_global_size(1)
get_global_size(2)
blockDim.x*gridDim.x
blockDim.y*gridDim.y
blockDim.z*gridDim.z

9.3 OpenCLプログラミング例(1)

ホストプログラム
リスト9-1にベクトルの和をOpenCLで並列計算するプログラムを示します。


リスト9-1 OpenCLホストプログラム(vadd_ocl_v1.c)
     1	/*
     2	vadd_ocl_v1.c (test program of OpenCL, version 1)
     3	
     4	Compile + Link:
     5	> cl /Ox vadd_ocl_v1.c OpenCL.lib
     6	
     7	Usage:
     8	> vadd_ocl_v1 <n> <loop> <platform> <device>
     9	*/
    10	
    11	#include <stdio.h>
    12	#include <stdlib.h>
    13	#include <time.h>
    14	#include <CL/cl.h>
    15	
    16	#define MAX_PLATFORMS (10)
    17	#define MAX_DEVICES (10)
    18	#define MAX_SOURCE_SIZE (100000)
    19	
    20	int main(int argc, char **argv)
    21	{
    22		// OpenCL
    23		cl_context       context       = NULL;
    24		cl_command_queue command_queue = NULL;
    25		cl_program       program       = NULL;
    26		cl_kernel        kernel        = NULL;
    27		cl_platform_id   platform_id[MAX_PLATFORMS];
    28		cl_device_id     device_id[MAX_DEVICES];
    29	
    30		// memory object
    31		cl_mem d_a = NULL;
    32		cl_mem d_b = NULL;
    33		cl_mem d_c = NULL;
    34	
    35		FILE *fp;
    36		char *source_str;
    37		size_t source_size;
    38		size_t global_item_size, local_item_size;
    39		size_t ret_size;
    40		cl_uint num_platforms;
    41		cl_uint num_devices;
    42		cl_int ret;
    43		char str[BUFSIZ];
    44	
    45		cl_uint platform = 0;
    46		cl_uint device = 0;
    47		int nloop = 1000;
    48		int n = 1000;
    49	
    50		// arguments
    51		if (argc >= 5) {
    52			n        = atoi(argv[1]);
    53			nloop    = atoi(argv[2]);
    54			platform = atoi(argv[3]);
    55			device   = atoi(argv[4]);
    56		}
    57	
    58		// alloc
    59		source_str = (char *)malloc(MAX_SOURCE_SIZE * sizeof(char));
    60	
    61		// setup host arrays
    62		float *a = (float *)malloc(n * sizeof(float));
    63		float *b = (float *)malloc(n * sizeof(float));
    64		float *c = (float *)malloc(n * sizeof(float));
    65		for (int i = 0; i < n; i++) {
    66			a[i] = (float)(1 + i);
    67			b[i] = (float)(1 + i);
    68		}
    69	
    70		// platform
    71		clGetPlatformIDs(MAX_PLATFORMS, platform_id, &num_platforms);
    72		if (platform >= num_platforms) {
    73			printf("error : platform = %d (limit = %d)\n", platform, num_platforms - 1);
    74			exit(1);
    75		}
    76	
    77		// device
    78		clGetDeviceIDs(platform_id[platform], CL_DEVICE_TYPE_ALL, MAX_DEVICES, device_id, &num_devices);
    79		if (device >= num_devices) {
    80			printf("error : device = %d (limit = %d)\n", device, num_devices - 1);
    81			exit(1);
    82		}
    83	
    84		// device name (option)
    85		clGetDeviceInfo(device_id[device], CL_DEVICE_NAME, sizeof(str), str, &ret_size);
    86		printf("%s\n", str);
    87	
    88		// context
    89		context = clCreateContext(NULL, 1, &device_id[device], NULL, NULL, &ret);
    90	
    91		// command queue
    92		command_queue = clCreateCommandQueue(context, device_id[device], 0, &ret);
    93	
    94		// source
    95		if ((fp = fopen("vadd.cl", "r")) == NULL) {
    96			fprintf(stderr, "kernel source open error\n");
    97			exit(1);
    98		}
    99		source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
   100		fclose(fp);
   101	
   102		// program
   103		program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
   104		if (ret != CL_SUCCESS) {
   105			fprintf(stderr, "clCreateProgramWithSource() error\n");
   106			exit(1);
   107		}
   108	
   109		// build
   110		if (clBuildProgram(program, 1, &device_id[device], NULL, NULL, NULL) != CL_SUCCESS) {
   111			fprintf(stderr, "clBuildProgram() error\n");
   112			exit(1);
   113		}
   114	
   115		// kernel
   116		kernel = clCreateKernel(program, "vadd", &ret);
   117		if (ret != CL_SUCCESS) {
   118			fprintf(stderr, "clCreateKernel() error\n");
   119			exit(1);
   120		}
   121	
   122		// memory object
   123		d_a = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret);
   124		d_b = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret);
   125		d_c = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret);
   126	
   127		// host to device
   128		clEnqueueWriteBuffer(command_queue, d_a, CL_TRUE, 0, n * sizeof(float), a, 0, NULL, NULL);
   129		clEnqueueWriteBuffer(command_queue, d_b, CL_TRUE, 0, n * sizeof(float), b, 0, NULL, NULL);
   130	
   131		// args
   132		clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_a);
   133		clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&d_b);
   134		clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&d_c);
   135		clSetKernelArg(kernel, 3, sizeof(int),    (void *)&n);
   136	
   137		// timer
   138		clock_t t0 = clock();
   139	
   140		// work item
   141		local_item_size = 256;
   142		global_item_size = ((n + local_item_size - 1) / local_item_size) * local_item_size;
   143	
   144		// run
   145		for (int loop = 0; loop < nloop; loop++) {
   146			clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL);
   147		}
   148	
   149		// device to host
   150		clEnqueueReadBuffer(command_queue, d_c, CL_TRUE, 0, n * sizeof(float), c, 0, NULL, NULL);
   151	
   152		// timer
   153		clock_t t1 = clock();
   154		double cpu = (double)(t1 - t0) / CLOCKS_PER_SEC;
   155	
   156		// output
   157		double sum = 0;
   158		for (int i = 0; i < n; i++) {
   159			sum += c[i];
   160		}
   161		double exact = n * (n + 1.0);
   162		printf("n=%d nloop=%d %e(%.6e) cpu[sec]=%.3f\n",
   163			n, nloop, sum, exact, cpu);
   164	
   165		// release
   166		clFlush(command_queue);
   167		clFinish(command_queue);
   168		clReleaseMemObject(d_a);
   169		clReleaseMemObject(d_b);
   170		clReleaseMemObject(d_c);
   171		clReleaseKernel(kernel);
   172		clReleaseProgram(program);
   173		clReleaseCommandQueue(command_queue);
   174		clReleaseContext(context);
   175	
   176		// free
   177		free(source_str);
   178		free(a);
   179		free(b);
   180		free(c);
   181	
   182		return 0;
   183	}

ソースコードの説明
14行目:OpenCLプログラムにはこのinclude文が必須です。
70-120行目:OpenCLに必須の前処理です。OpenCLはいろいろな環境で動くためにこのような処理が必要になります。 コードが煩雑になりますが定型的な処理です。
platform→device→context→command queue→カーネルソースコード→program→build→kernelの順に処理します。
78行目:"CL_DEVICE_TYPE_ALL"を"CL_DEVICE_TYPE_GPU"に変えるとGPUだけが対象デバイスになります。
123-125行目:GPUで使用する配列を確保します。(CUDAのcudaMallocに対応)
128-129行目:ホストからデバイスにメモリーを転送します。(CUDAのcudaMemcpy(...,cudaMemcpyHostToDevice)に対応)
132-135行目:カーネルの引数を代入します。(CUDAにない機能)
141-142行目:ワークアイテム(スレッドの構成)を指定します。global_item_sizeはlocal_item_sizeの整数倍であることが必要です。(CUDAのExecution configurationに対応)
150行目:デバイスからホストにメモリーを転送します。(CUDAのcudaMemcpy(...,cudaMemcpyDeviceToHost)に対応)
168-170行目:GPUで使用した配列を解放します。(CUDAのcudaFreeに対応)

カーネルプログラム
リスト9-2にカーネルプログラムを示します。


リスト9-2 OpenCLカーネルプログラム(vadd.cl)
     1	__kernel void vadd(global const float *a, global const float *b, global float *c, int n)
     2	{
     3		int i = get_global_id(0);
     4	
     5		if (i < n) {
     6			c[i] = a[i] + b[i];
     7		}
     8	}

ソースコードの説明
3行目:globalなスレッド番号を取得します。それがこの場合配列のインデックスになります。
5行目:配列の大きさはスレッド数の倍数とは限らないのでこの条件判定が必要です。

9.4 OpenCLプログラミング例(2)

プログラミングの改良
上で示したプログラムはOpenCLを用いた最小のプログラムですが、 大規模なアプリケーションを開発するための出発点としては不適切です。
すわなちOpenCL固有の処理とその他の処理が分離されておらず、 また計算アルゴリズムを検証するための逐次計算が実装されていないために開発効率が悪くなります。
その点を考慮したプログラムをリスト9-3に示します。 なお、カーネルプログラムはリスト9-2と共通です。


リスト9-3 OpenCLホストプログラム(vadd_ocl_v2.c)
     1	/*
     2	vadd_ocl_v2.c (test program of OpenCL, version 2)
     3	
     4	Compile + Link:
     5	> cl /Ox vadd_ocl_v2.c OpenCL.lib
     6	
     7	Usage:
     8	> vadd_ocl_v2 <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	#define MAX_PLATFORMS (10)
    23	#define MAX_DEVICES (10)
    24	#define MAX_SOURCE_SIZE (100000)
    25	
    26	static int setup_ocl(cl_uint, cl_uint, char *);
    27	static void vadd_calc();
    28	static void vadd();
    29	
    30	cl_command_queue Queue;
    31	cl_kernel k_vadd;
    32	int N;
    33	float *A, *B, *C;
    34	cl_mem d_A, d_B, d_C;
    35	int OCL;
    36	
    37	int main(int argc, char **argv)
    38	{
    39		int platform = 0;
    40		int device = 0;
    41		int nloop = 1000;
    42	
    43		// arguments
    44		N = 1000;
    45		if (argc >= 5) {
    46			N        = atoi(argv[1]);
    47			nloop    = atoi(argv[2]);
    48			platform = atoi(argv[3]);
    49			device   = atoi(argv[4]);
    50		}
    51		OCL = (platform >= 0);
    52	
    53		// alloc host arrays
    54		size_t size = N * sizeof(float);
    55		A = (float *)malloc(size);
    56		B = (float *)malloc(size);
    57		C = (float *)malloc(size);
    58	
    59		// setup problem
    60		for (int i = 0; i < N; i++) {
    61			A[i] = (float)(1 + i);
    62			B[i] = (float)(1 + i);
    63		}
    64	
    65		// setup OpenCL
    66		if (OCL) {
    67			char msg[BUFSIZ];
    68			int ret = setup_ocl((cl_uint)platform, (cl_uint)device, msg);
    69			printf("%s\n", msg);
    70			if (ret) {
    71				exit(1);
    72			}
    73		}
    74	
    75		// timer
    76		clock_t t0 = clock();
    77	
    78		// copy host to device
    79		if (OCL) {
    80			clEnqueueWriteBuffer(Queue, d_A, CL_TRUE, 0, size, A, 0, NULL, NULL);
    81			clEnqueueWriteBuffer(Queue, d_B, CL_TRUE, 0, size, B, 0, NULL, NULL);
    82		}
    83	
    84		// run
    85		for (int loop = 0; loop < nloop; loop++) {
    86			vadd_calc();
    87		}
    88	
    89		// copy device to host
    90		if (OCL) {
    91			clEnqueueReadBuffer(Queue, d_C, CL_TRUE, 0, size, C, 0, NULL, NULL);
    92		}
    93	
    94		// timer
    95		clock_t t1 = clock();
    96		double cpu = (double)(t1 - t0) / CLOCKS_PER_SEC;
    97	
    98		// sum
    99		double sum = 0;
   100		for (int i = 0; i < N; i++) {
   101			sum += C[i];
   102		}
   103	
   104		// output
   105		double exact = N * (N + 1.0);
   106		printf("n=%d nloop=%d %e(%.6e) cpu[sec]=%.3f\n",
   107			N, nloop, sum, exact, cpu);
   108	
   109		// release
   110		if (OCL) {
   111			clReleaseMemObject(d_A);
   112			clReleaseMemObject(d_B);
   113			clReleaseMemObject(d_C);
   114			clReleaseKernel(k_vadd);
   115			clReleaseCommandQueue(Queue);
   116		}
   117	
   118		// free
   119		free(A);
   120		free(B);
   121		free(C);
   122	
   123		return 0;
   124	}
   125	
   126	// setup OpenCL
   127	static int setup_ocl(cl_uint platform, cl_uint device, char *msg)
   128	{
   129		cl_context     context = NULL;
   130		cl_program     program = NULL;
   131		cl_platform_id platform_id[MAX_PLATFORMS];
   132		cl_device_id   device_id[MAX_DEVICES];
   133	
   134		FILE *fp;
   135		char *source_str;
   136		char str[BUFSIZ];
   137		size_t source_size, ret_size, size;
   138		cl_uint num_platforms, num_devices;
   139		cl_int ret;
   140	
   141		// alloc
   142		source_str = (char *)malloc(MAX_SOURCE_SIZE * sizeof(char));
   143	
   144		// platform
   145		clGetPlatformIDs(MAX_PLATFORMS, platform_id, &num_platforms);
   146		if (platform >= num_platforms) {
   147			sprintf(msg, "error : platform = %d (limit = %d)", platform, num_platforms - 1);
   148			return 1;
   149		}
   150	
   151		// device
   152		clGetDeviceIDs(platform_id[platform], CL_DEVICE_TYPE_ALL, MAX_DEVICES, device_id, &num_devices);
   153		if (device >= num_devices) {
   154			sprintf(msg, "error : device = %d (limit = %d)", device, num_devices - 1);
   155			return 1;
   156		}
   157	
   158		// device name (option)
   159		clGetDeviceInfo(device_id[device], CL_DEVICE_NAME, sizeof(str), str, &ret_size);
   160		sprintf(msg, "%s (platform = %d, device = %d)", str, platform, device);
   161	
   162		// context
   163		context = clCreateContext(NULL, 1, &device_id[device], NULL, NULL, &ret);
   164	
   165		// command queue
   166		Queue = clCreateCommandQueue(context, device_id[device], 0, &ret);
   167	
   168		// source
   169		if ((fp = fopen("vadd.cl", "r")) == NULL) {
   170			sprintf(msg, "kernel source open error");
   171			return 1;
   172		}
   173		source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
   174		fclose(fp);
   175	
   176		// program
   177		program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
   178		if (ret != CL_SUCCESS) {
   179			sprintf(msg, "clCreateProgramWithSource() error");
   180			return 1;
   181		}
   182	
   183		// build
   184		if (clBuildProgram(program, 1, &device_id[device], NULL, NULL, NULL) != CL_SUCCESS) {
   185			sprintf(msg, "clBuildProgram() error");
   186			return 1;
   187		}
   188	
   189		// kernel
   190		k_vadd = clCreateKernel(program, "vadd", &ret);
   191		if (ret != CL_SUCCESS) {
   192			sprintf(msg, "clCreateKernel() error");
   193			return 1;
   194		}
   195	
   196		// memory object
   197		size = N * sizeof(float);
   198		d_A = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &ret);
   199		d_B = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &ret);
   200		d_C = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &ret);
   201	
   202		// release
   203		clReleaseProgram(program);
   204		clReleaseContext(context);
   205	
   206		// free
   207		free(source_str);
   208	
   209		return 0;
   210	}
   211	
   212	// entry point
   213	static void vadd_calc()
   214	{
   215		if (OCL) {
   216			size_t global_item_size, local_item_size;
   217	
   218			// args
   219			clSetKernelArg(k_vadd, 0, sizeof(cl_mem), (void *)&d_A);
   220			clSetKernelArg(k_vadd, 1, sizeof(cl_mem), (void *)&d_B);
   221			clSetKernelArg(k_vadd, 2, sizeof(cl_mem), (void *)&d_C);
   222			clSetKernelArg(k_vadd, 3, sizeof(int),    (void *)&N);
   223	
   224			// work item
   225			local_item_size = 256;
   226			global_item_size = ((N + local_item_size - 1) / local_item_size)
   227			                 * local_item_size;
   228	
   229			// run
   230			clEnqueueNDRangeKernel(Queue, k_vadd, 1, NULL, &global_item_size, &local_item_size, 0, NULL, NULL);
   231		}
   232		else {
   233			// serial code
   234			vadd();
   235		}
   236	}
   237	
   238	// serial code
   239	static void vadd()
   240	{
   241		for (int i = 0; i < N; i++) {
   242			C[i] = A[i] + B[i];
   243		}
   244	}

ソースコードの説明
30-35行目:複数の関数から呼ばれる変数(プログラムの核をなす変数)をグローバル変数にします。
48行目:3番目の引数に負の値を代入するとOpenCLを使用しない逐次計算プログラムになります。
66-73行目:OpenCL固有の前処理は一か所に集めます。
213行目:この関数が計算の入り口になります。この中で並列処理と逐次処理を場合分けしています。
219-222行目:引数の内容は計算の途中で変わることがありますのでカーネルを呼び出す直前に設定します。
225-227行目:ここではワークアイテムをカーネルを呼び出す直前に設定していますが、 これは計算の途中で変わることはありませんので外部で設定しても構いません。
239-244行目:逐次コードを先に作成して十分テストし、計算式が正しく実装されていることを確認します。

逐次版と並列版の開発
プログラムの最終目的が並列計算であってもいきなり並列計算プログラムを実装してもうまく動くことはほとんどありません。
最初に並列化を想定したプログラム構造の逐次版を作成し、 計算式が正しく実装されていることを十分確認してから並列版を実装することが開発の近道です。
並列版固有の処理は以下の通りです。これらはすべて定型的な作業ですので注意深く実装すれば間違いは少なくなります。

  1. OpenCLの前処理、CommandQueueとカーネル変数の作成まで
  2. メモリーオブジェクトの作成
  3. ホスト→デバイス間とデバイス→ホスト間のメモリーコピー
  4. カーネル引数の代入
  5. ワークアイテムの作成(これは計算時間に影響しますのでいろいろ変えてテストする必要があります)
  6. 逐次版関数を参考にカーネル関数を作成(オフラインコンパイラで文法チェック)

9.5 コンパイル・リンク・実行

コンパイル・リンク方法
コンパイル・リンク方法は以下の通りです(VC++の場合)。
> cl /Ox vadd_ocl_v2.c OpenCL.lib

プログラムの実行方法
プログラムの実行方法は以下の通りです。
> vadd_ocl_v2 配列の大きさ 繰り返し回数 プラットフォーム番号 デバイス番号
例えば以下のようになります。
> vadd_ocl_v2 1000000 1000 1 0 (プラットフォーム番号=1, デバイス番号=0)
> vadd_ocl_v2 1000000 1000 -1 0 (OpenCLを用いないで逐次計算するとき)
繰り返し回数は計算時間の測定誤差を小さくするためです。

9.6 OpenCLの計算時間

表9-2にOpenCLの計算時間を示します。
GPUは2.の2種類で計算しています。
配列の大きさ(=N)と繰り返し回数(=L)の積は一定(=10^10)です。従って全体の演算量は同じです。
GPUではNo.1-3より問題のサイズによらず計算時間はほぼ一定であることがわかります。
GPUのNo.4ではカーネル起動回数が増えるために計算時間が増えます。 リスト9-3よりカーネル起動回数は繰り返し回数と同じです。 繰り返し回数が1,000,000回のときGPUで5秒余分に時間がかかっています。 これからGPUのカーネル起動のオーバーヘッドはおよそ5μsecと評価することができます。 (通常のアプリケーションではカーネルをこのように多数回呼ぶことは少ないです)

表9-2 OpenCLの計算時間(Windows,()内はCPU1コアとの速度比)
No.配列の大きさN繰り返し回数LCPU1コアGPU
110,000,0001,00011.33秒 (1.0)0.63秒 (18.0)
21,000,00010,0007.39秒 (1.0)0.63秒 (11.7)
3100,000100,0006.07秒 (1.0)0.65秒 (9.3)
410,0001,000,0005.30秒 (1.0)6.06秒 (0.9)

9.7 OpenCLのデバイスについて

前節では計算が最も速い外付けGPUを用いましたが、 OpenCLプログラムはCPUまたはCPU(IvyBridge以降)内蔵のGPUで計算することもできます。[21] (以下では簡単のため、CPU内蔵GPUをiGPU、外付GPUをdGPUと呼びます)
計算するハードウェアはOpenCLのプラットフォーム番号とデバイス番号によって選択することができます。
表9-3にテスト結果を示します。計算条件は表9-2のNo.1と同じです。
No.0,No.4よりCPUではOpenCLプログラムの速度は逐次版と同じになっています。 これからCPUで計算するならOpenCLで並列化するよりその他の方法で並列化するほうが勝っていることがわかります。 (CPUは並列化の余地があるため)
No.1,No.2より複数のGPUが実装されているときはそれぞれ異なるデバイス番号が与えられます。 プログラムがマルチGPUに対応していないときは一つのGPUが使用されます。
No.3,No.4よりiGPUはCPUより少し速いですが、No.2,No.3のdGPUには遠く及びません。
以上からOpenCLプログラムはCPUまたはiGPUでも動きますが、 計算速度はdGPUに比べて大幅に遅く実用的な意味はありません。
なお、プラットフォーム番号とデバイス番号は環境によって変わりますので、 いろいろ変えて実行し表示されるデバイス名から番号を決定して下さい。

表9-3 OpenCLのデバイスと計算時間(計算条件No.1、Windows)
No.ハードウェアプラットフォーム番号デバイス番号計算時間
0CPU逐次版(参考)--10.99秒
1外付GPU(dGPU)GTX1070,1台目000.63秒
2外付GPU(dGPU)GTX1070,2台目010.63秒
3CPU内蔵GPU(iGPU)HD4600107.71秒
4CPU i7-4770K1111.33秒