グラフィックスボード(GPU)が搭載されたコンピュータでは、
その高い演算能力を汎用的な科学技術計算に用いることができます。
そのためのプログラミング言語をOpenCL(Open Computing Language)と呼びます。
CPUがフロントエンドとして動作し、計算の主要部はGPUが行います。
OpenCLプログラムはCPUでも動作しますが性能上の利点は少なく主にGPUで使用します。
OpenCLプログラム
OpenCLプログラムは通常のC/C++で記述されるホストプログラムと、
OpenCL C言語で記述されるカーネルプログラム(拡張子.cl)から成ります。
ホストプログラムはCPUで実行され、カーネルプログラムはGPUで実行されます。
両者の関係はCUDAと同じです。
OpenCL C言語
カーネルを記述するOpenCL C言語のC/C++との違いは以下の通りです。
オンラインコンパイルとオフラインコンパイル
OpenCLプログラムにはオンラインコンパイルとオフラインコンパイルの2種類のモードがあります。
それぞれの意味は以下の通りです。
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も関数が用意されており便利です。
| No. | OpenCL ワークアイテム | CUDA execution configuration |
|---|---|---|
| 1 | get_local_id(0) get_local_id(1) get_local_id(2) | threadIdx.x threadIdx.y threadIdx.z |
| 2 | get_local_size(0) get_local_size(1) get_local_size(2) | blockDim.x blockDim.y blockDim.z |
| 3 | get_group_id(0) get_group_id(1) get_group_id(2) | blockIdx.x blockIdx.y blockIdx.z |
| 4 | get_num_groups(0) get_num_groups(1) get_num_groups(2) | gridDim.x gridDim.y gridDim.z |
| 5 | get_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 |
| 6 | get_global_size(0) get_global_size(1) get_global_size(2) | blockDim.x*gridDim.x blockDim.y*gridDim.y blockDim.z*gridDim.z |
ホストプログラム
リスト9-1にベクトルの和をOpenCLで並列計算するプログラムを示します。
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にカーネルプログラムを示します。
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行目:配列の大きさはスレッド数の倍数とは限らないのでこの条件判定が必要です。
プログラミングの改良
上で示したプログラムはOpenCLを用いた最小のプログラムですが、
大規模なアプリケーションを開発するための出発点としては不適切です。
すわなちOpenCL固有の処理とその他の処理が分離されておらず、
また計算アルゴリズムを検証するための逐次計算が実装されていないために開発効率が悪くなります。
その点を考慮したプログラムをリスト9-3に示します。
なお、カーネルプログラムはリスト9-2と共通です。
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行目:逐次コードを先に作成して十分テストし、計算式が正しく実装されていることを確認します。
逐次版と並列版の開発
プログラムの最終目的が並列計算であってもいきなり並列計算プログラムを実装してもうまく動くことはほとんどありません。
最初に並列化を想定したプログラム構造の逐次版を作成し、
計算式が正しく実装されていることを十分確認してから並列版を実装することが開発の近道です。
並列版固有の処理は以下の通りです。これらはすべて定型的な作業ですので注意深く実装すれば間違いは少なくなります。
コンパイル・リンク方法
コンパイル・リンク方法は以下の通りです(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-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と評価することができます。
(通常のアプリケーションではカーネルをこのように多数回呼ぶことは少ないです)
| No. | 配列の大きさN | 繰り返し回数L | CPU1コア | GPU |
|---|---|---|---|---|
| 1 | 10,000,000 | 1,000 | 11.33秒 (1.0) | 0.63秒 (18.0) |
| 2 | 1,000,000 | 10,000 | 7.39秒 (1.0) | 0.63秒 (11.7) |
| 3 | 100,000 | 100,000 | 6.07秒 (1.0) | 0.65秒 (9.3) |
| 4 | 10,000 | 1,000,000 | 5.30秒 (1.0) | 6.06秒 (0.9) |
前節では計算が最も速い外付け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に比べて大幅に遅く実用的な意味はありません。
なお、プラットフォーム番号とデバイス番号は環境によって変わりますので、
いろいろ変えて実行し表示されるデバイス名から番号を決定して下さい。
| No. | ハードウェア | プラットフォーム番号 | デバイス番号 | 計算時間 |
|---|---|---|---|---|
| 0 | CPU逐次版(参考) | - | - | 10.99秒 |
| 1 | 外付GPU(dGPU)GTX1070,1台目 | 0 | 0 | 0.63秒 |
| 2 | 外付GPU(dGPU)GTX1070,2台目 | 0 | 1 | 0.63秒 |
| 3 | CPU内蔵GPU(iGPU)HD4600 | 1 | 0 | 7.71秒 |
| 4 | CPU i7-4770K | 1 | 1 | 11.33秒 |