Cerium による並列処理向け I/O の設計と実装

Masataka Kohagura 12th, February

担当教官 : 河野 真治

研究背景と目的

近年のCPUのほとんどはマルチコアであり、それらの性能を引き出すためには並列プログラミングが必須となっている。そこで当研究室では Cerium Library の開発を行い、提供することによって並列プログラミングを容易にしている。

先行研究では Task の並列化によって、プログラム全体の処理速度は向上しているが、ファイル読み込み等の I/O に対して並列に Task が動作するようにはされていない。

本研究では I/O と Task の並列化の設計、実装によってプログラム全体の処理速度、処理効率を上げていく。

Cerium とは

  1. Taskを生成
  2. 依存関係のチェック
  3. Schedulerに転送
  4. 並列実行

CpuThreads、Schedulerに対応させる形でGpuThreadsとGpuSchedulerを作成した

GPU Task実行の流れ


kernel fileの記述

__kernel void // kernel.cl(kernel file)
twice(__global int *input_data,__global int *output_data) {
    long count = (long)data_count[0];
    for (int i = 0; i< count; i++) {
        output_data[i] = input_data[i] * 2;;
    }
}

GPU Task 実行の流れ


kernelをTaskとしてCeriumに登録

void
task_init(void) { // task_init.cc
    GpuSchedRegister(Twice, "./twice.cl", "twice");
}
第1引数
Twice		 Taskのid。enumで定義されている
TaskManagerはこの値でTaskを識別する
第2引数
"./twice.cl"		 OpenCLが処理するkernelが記述されているファイルのパス
第3引数
"twice"		 関数の指定。kernel file内にある、実行する関数名を指定
Taskにあたる部分

GPU Task 実行の流れ


GPU Task生成

// main.cc
HTaskPtr twice = manager->create_task(Twice);
twice->set_inData(0, indata, sizeof (int)*length);
// twice->set_inData(1, indata2, sizeof (int)*length);
twice->set_outData(0, outdata, sizeof (int)*length);
twice->set_cpu(GPU);
twice->spawn();

Cerium OpenCL API比較

ベンチマーク

Time
1 CPU  796 ms
2 CPU  439 ms
6 CPU  153 ms
12 CPU  96 ms
24 CPU  89 ms
GPU(改良前)  330752 ms
GPU(改良後)  5306 ms

10万入力によるBitonic Sort

実験環境

OS : MacOS 10.8.2
CPU : 2*2.66GHz 6-CoreIntel Xeon
Memory : 16GB
Compiler : Apple clang version 4.1
　　　　　(based on LLVM 3.1svn)
GPU : AMD ATI Radeon HD 5870 1024MB

結果

1coreのCPUよりも10倍遅い

充分な性能が出なかったため、一度に送信する data のサイズを増やす改善を行ったところ、 約60倍程実行速度が向上した

考察

性能向上は見られたが、CPUと比べると未だ差が開いている GPU向けに適切なチューニングが今後の課題となる


改善案

データ並列

データを2、3次元に分割し、分割した部分に対して並列処理する並列化手法。

OpenCL ではin/outするデータ郡をWork Itemと呼ぶ。

各Work Item のサイズを指定するとOpenCLがデータ並列で実行する。

同期機構

GpuSchedulerはCommand Queueの内部でパイプライン的に実行を行っている。 パイプラインを構成するには処理にwaitをかける必要がある。現在はclWaitForEvent APIを使用

API 機能
clFlush() Command Queueに投入したTask全てをDeviceで実行する
clWaitForEvent() 特定の処理の終了を待つ

clFlushは実行は保証するが、終了は保証しない仕様になっている

新しい同期

FrameWork Dependency
Cerium Task Dependency
OpenCL Data Dependency

Task Dependency:Schedulerで依存関係が決定
Data Dependency:GPUに読み込まれた時に決定

GPGPUはなるべくGPU内部で処理を行う方が高速なため、性能向上が見込める

まとめ

今後の課題

ベンチマーク

Time
1 CPU  67 ms
2 CPU  34 ms
6 CPU  12 ms
12 CPU  9 ms
24 CPU  6 ms
GPU  10201 ms

word count

10MBのテキストファイルを分割
各Taskがcountしていく

スペースと改行区切りでword countしていく

結果

CPUの方が150倍早い