- つまり、GPGPU ではデータ転送を頻繁に行う Task は並列度が低くなる。 -

データ転送

- - - - - @@ -179,15 +165,34 @@

並列化手法とデータ転送

+ 並列化手法は大まかにタスク並列とデータ並列に分けることができる。 +

+ タスク並列では複数のプロセッサで異なる演算を行い、データ並列では複数のプロセッサで同じ演算を行う。 +

+ タスク並列では、同じ演算を行う場合でもタスクの数だけデータ転送を行う必要がある。 +

+ データ並列では、すべてのデータを一度で転送する。タスク側で担当する範囲を計算し、その部分に対して処理を行う。 +

+ 同じ演算を行いたい場合、データ転送が少なくなるのでデータ並列の方が有効である。特に GPGPU ではデータ転送がオーバーヘッドとなるのでその差が顕著に現れる。 +

+ +

OpenCL, CUDA による GPGPU への対応

- OpenCL では CommandQueue、CUDA では Stream をそれぞれ用いて GPU に命令を転送する。 + OpenCL では CommandQueue、CUDA では Stream という同期機構がある。それぞれに命令を投入することで GPU を制御する。

CommandQueue, Stream に投入された命令は発行された順序で実行される。これらを複数用いることで命令を並列に実行することが可能になる。

- OpenCL および CUDA では Task を kernel という単位で記述する。 + OpenCL および CUDA ではタスクを kernel という単位で記述する。

Cerium において、CPU/Cell 上で実行する場合と GPU 上で実行する場合、それぞれほぼ同じ形式の記述が可能である。 @@ -195,20 +200,78 @@

Multi Stream による並列実行

複数の同期機構を用いた並列実行

- Stream に投入された命令は投入された順序で実行されることが保証される。 + CommandQueue, Stream に投入された命令は投入された順序で実行されることが保証される。

- 異なる Stream の命令の依存関係が解消され実行可能な状態の場合、命令は並列に実行される。これは OpenCL で CommandQueue を用いて実行する場合も同様である。 + 異なる CommandQueue, Stream の命令の依存関係が解消され実行可能な状態の場合、命令は並列に実行される。 +

+ HtoD, kernel, DtoH は cuMemcpyHtoD, cuLaunchKernel, cuMemcpyDtoH という CUDA の API である。ホストからデバイスへのデータ転送、kernel の実行、デバイスからホストへのデータ転送をそれぞれ行う。OpenCL では clEnqueueReadBuffer, clEnqueueTask, clEnqueueWriteBuffer に相当する。 +

+ +

命令の投入方法

+ 注意点として CUDA では以下のように命令を投入しないと命令が並列に実行されない。

+// BAD
+for (int i=0;i<stream_num;i++) {
+    cuMemcpyHtoDAsync();
+    cuLaunchKernel();
+    cuMemcpyDtoHAsync();
+}

+// GOOD
+for (int i=0;i<stream_num;i++)
+    cuMemcpyHtoDAsync();
+             
+for (int i=0;i<stream_num;i++)
+    cuLaunchKernel();
+
+for (int i=0;i<stream_num;i++)
+    cuMemcpyDtoHAsync();

+ +

host の記述

+ 2つの input の積を取り、output に返す例題の host +

+void
+multi_init(TaskManager *manager)
+{
+    A = new float[length];
+    B = new float[length];
+    C = new float[length];
+    for(int i=0; i<length; i++) {
+        A[i]=(float)(i+1000);
+        B[i]=(float)(i+1)/10.f;
+    }
+    
+    HTask* multiply = manager->create_task(MULTIPLY_TASK);
+    multiply->set_cpu(CPU_TYPE);
+
+    multiply->set_inData(0,(memaddr)A, sizeof(float)*length);
+    multiply->set_inData(1,(memaddr)B, sizeof(float)*length);
+
+    multiply->set_outData(0,(memaddr)C, sizeof(float)*length);
+
+    multiply->iterate(length); 
+}

+ CPU, GPU ともに完全に同じ記述で Task を生成することができる。set_cpu で実行するデバイスを選択することができる。

kernel の記述

- 2つの input の積を取り、output に返す例題 Multiply + 2つの input の積を取り、output に返す例題の kernel

 // MultiCore CPU
@@ -253,13 +316,13 @@

@@ -312,21 +375,23 @@

Cerium におけるデータ並列のための API(1/2)

CPU におけるデータ並列のための API

- タスク並列で実行する場合、タスクが担当するデータを実行する数だけデータ転送を行う。 -

- データ並列で実行する場合、すべてのデータを一度で転送する。kernel で担当する範囲を計算し、その部分に対して処理を行う。データ転送の回数が少なくなるので、OpenCL, CUDA ではデータ並列による実行が推奨されている。 + つまり、GPGPU ではデータ転送を頻繁に行うタスクは並列度が低くなる。

- 3D グラフィックのような多次元のデータを処理する場合に高い並列度を保つには、データを分割して並列に実行する機能が必要である。 + 1つの記述から ID の異なる複数の kernel を自動生成する。

- これを OpenCL, CUDA ではデータ並列と呼んでいる。 + OpenCL では get_global_id 関数、CUDA では組み込み変数を参照することで ID を取得することができる。

- 1つの記述から ID の異なる複数の kernel を自動生成する。OpenCL では get_global_id 関数、CUDA では組み込み変数の blockDim, blockIdx, threadIdx を参照することで index を取得することができ、index から担当する範囲を決定し処理を行う。 + 取得した ID から担当する範囲を計算し、その部分に対して処理を行う。

@@ -336,27 +401,6 @@

Cerium におけるデータ並列のための API(2/2)

-void
-HTask::iterate(long x, long y, long z) {
-    this->flag.dim_count = 1;
-    TaskList *tl = (TaskList*)rbuf;
-    for (;tl;tl=tl->next) {
-        tl->set_last(last);
-        tl->dim=1;
-        tl->x=x;
-        tl->y=y;
-        tl->z=z;
-    }
-    mimpl->spawn_task(this);
-}

- iterate を使用するとデータ並列で実行する Task として登録される。この時点で TaskList にパラメータが設定され、TaskManager から Scheduler に転送される。 -

- -

ベンチマーク

測定環境

- Task に ID と input/output data を割り当てる「iterate」という API を実装した。 + データ並列で実行する API として「iterate」を実装した。

iterate
CPU におけるデータ並列

1つの記述から複数の Task を生成する。
生成した複数の Task に ID と input/output data を割り当てる。
Multi Core CPU でデータ並列を行う場合は SchedTask に index が設定されており、それを参照することで担当する範囲を決定し処理を行う。
Task を複製する。
複製した Task に index と input/output data を割り当て、各 CPU に転送する。
各 CPU にも ID を割り当てる。
各 CPU に割り当てられた ID から実行する Task を選択する。
Task 側では SchedTask を参照することで index を取得することができる。

はじめに(2/2)