データを2、3次元に分割し、分割した部分に対して @@ -298,14 +295,7 @@

- HTaskPtr twice = manager->create_task(Twice); - twice->set_cpu(GPU); - twice->set_inData(0,(memaddr)input, sizeof(float)*length); - twice->set_outData(0,(memaddr)output, sizeof(float)*length); - // paramに0~length-1(index)をsetしたtaskをlength個spawnする - twice->iterate(length); -

iterate(length)とすることで、TaskManagerがデータ並列用にTaskを生成

iterate(Kernel)

-
-            long i = (long)scheduler->x; // (long)scheduler->get_param(0);
-            output[i]=input[i]*2;
-

-            long i = get_global_id(0);
-            output[i]=input[i]*2;
-

MultiCoreではschedulerの持つメンバ変数x、OpenCLはget_global_idというAPIを用いて、自分に割り振られたid(index)を取得する。その後、そのindexに対して処理を行う。 @@ -347,18 +329,8 @@

iterateの実装

-            void
-            HTask::iterate(long x) {
-                tl->dim=1;
-                tl->x=x;
-                // 1次元なのでy軸、z軸の要素数は1
-                tl->y=1;
-                tl->z=1;
-                mimpl->spawn_task(this);
-            }
-

渡されたlengthはwork item のx座標の要素数となる。

渡されたlengthはwork item の要素数となる。

taskにデータ分割のためのdata(dimension, や workitem size)をsetする。

Schedulerはtaskが持ってるdataを元にtaskを複数生成し、idを割り当てる。 @@ -371,16 +343,7 @@

iterateの実装(多次元)

-            void
-            HTask::iterate(long x, long y) {
-                tl->dim=2;
-                tl->x=x;
-                tl->y=y;
-                tl->z=1;
-                mimpl->spawn_task(this);
-            }
-

引数を複数渡せば多次元のデータ分割ができる。3次元までサポートする。

@@ -395,14 +358,10 @@

GpuSchedulerがOpenCLのAPIを呼び出し、GPUの制御を行う

- TaskManagerからTaskを受け取り、Command Queueにenqueueする + TaskManagerから受け取ったTaskは必要なパラメータを各種持っている

-
-clEnqueueNDRangeKernel(command_queue, kernel[cur], task->dim, NULL,
-                       &task->x, &task->y, &task->z, NULL, NULL);
-

Taskの持つメンバ変数

@@ -536,7 +495,7 @@

- CPUとGPUの同時実行は、実行するTaskに対してset_cpu(SPE_ANY)とすれば良い + CPUとGPUの同時実行は、実行するTaskに対してset_cpu(ANY_ANY)とすれば良い

TaskはCPUとGPU、交互に割り振られる @@ -549,47 +508,33 @@

ベンチマーク

- - - - - - - -

- -	- 結果 - - - 現段階ではSchedulingを行ってないため、GPU単体で動かした時よりも遅くなる。 - - - CPUとGPUは実行速度に差が出る場合がある。 - Taskの計算内容によってはどちらかのアーキテクチャに任せた方が良い。 - それを確認するベンチマークを動かす。 - - - - - -

- - - - - -

ベンチマーク

- -

結果

+ 現段階ではSchedulingを行ってないため、GPUやCPUを単体で動かした時よりも遅くなる。 +

+ CPUとGPUは実行速度に差が出る場合がある。 + Taskの計算内容によってはどちらかのアーキテクチャに任せた方が良い。 + それを確認するベンチマークを動かす。 +

+ +

+ + + + + +

ベンチマーク

+ +

FFTはSpinFactやButterfly演算等、様々なTaskで構成されている。それぞれのTaskについて、実行時間を計測した。

- SpinFactのTaskに関しては、CPUの方が実行速度が早い。 + 大体のTaskはGPUの方が速いが、SpinFactのTaskに関してはCPUの方が実行速度が速い。

@@ -605,7 +550,7 @@

Kernelの記述

iterate(Kernel)

iterateの実装

iterateの実装(多次元)

Taskの持つメンバ変数

ベンチマーク

結果

ベンチマーク

結果

ベンチマーク

今後の課題