-title: PS3 上でのゲームプログラミング

--abstract:

本大学では、学生の実習にPS3の実機を用いたゲーム作成を導入している。
Cell.B.E は、安価に手に入る Many Core 機であるが、従来とは異なる
プログラミングを要求される。特に、PS3のGPUであるRSXの仕様は公開されてない
ので、教育目的を含めて、Cerium Rendering Engine を作成した。
Cerium Task Manager はメインメモリとのやりとりをDMAで隠した
関数であるTaskのQueueを管理し、恒常的な並列性を維持する。
Cerium Engine での若干の並列プログラミングの例について報告する。

--abstract-e:

Game Programming on PS3 is introduced in our University as programming
training. PS3 is a cheap Many Core Architecture, but it requires
different programming from conventional one. Especially specification of PS3's
GPU RSX is not opened, so we decided to create Cerium Rendering
Engine as an educational tool. 
In order to keep
constant parallelism,
Cerium Task Manager provides Task
queue management of Task Functions, which has hiding mechanism of
memory copy from main memory  using DMA. 
We report examples of parallel programming on our Cerium Engine.




--歴史的経緯

当研究室では、主に PlayStation を用いた実機上のゲームプログラミング
を取り扱って来た。2007年から、PS3 上のゲーム作成を行っており、
Cell B.E. \cite{cell} でのプログラミングを学生に行わせている。

PS3 上では(最近のPS3ではサポートされてないが)、Fedora Linux が動作し、
その上で、6個のCellにアクセスすることが可能である。Graphic Processor
RSX の仕様は公開されないので、初期は、Open GL Mesa を用いて、PPE
で描画を行っていた。

Open GL Mesa の描画部分をCell のSPEを行なったが、Open GL Mesa の
実装は、描画ルーチンを「すべて C のマクロで記述」するような手法であり、
一度、マクロを展開したソースを変更して SPE に対応させるような
手法となる。全体の対応は現実的なものでないと判断した。Galium \cite{galium}
では、同様のアプローチが行われている。

そこで、Open GL とは
独立なEngine を、SPE 上のTaskを管理するマネージャ(Cerium Engine)
と一緒に作成する
ことにした。その際に、SPURS Engine \cite{spurs}を参考にしている。
ただし、SPURS Engine の仕様も公開されていないために、独自実装となっている。

その後、2009年になって Open CLの最初の仕様が公開されたが、Cerium
Engine は、比較的似た構成となっているが、どちらかと言えば、Cell
に特化した構成となっている。

狙いとしては、SPEの繁雑なSIMD処理は、なるべく使わずに、学生に
Many Core CPU のプログラムに入っていけるライブラリを目指している。
特定のアプリケーションで最高速を目指すのではなく、ゲームで
出て来る汎用的な処理を自然に並列に実行させることが出来る。

現在は簡単なゲームをプログラムを作成できるようになっており、Cell
だけでなく、Mac OS X上、Linux 上でも動作する。Core i7 等のマルチコア
にも対応可能であるが、現在は実機がないので対応していない。

本論文では、Many Core Programming の実際の問題(特に、Rendering Engine
と Word Count を例題に) を取り上げ、Cell 固有の問題と、Cerium Engine
での解決を示す。

また、Scala / Erlang などのアプローチや、Open CL との比較を行う。

--Many Core Programming の特徴

PS3 で用いられた Cell B.E. では、2 thread のPower PCである PPE と、128bit
vector register を持つ SPE が6個付いている。SPE は 256Kbyte Local Store を PPE の
メインメモリとは別に持っていて、明示的なDMAで通信を行う。
GeForce GTX 280 等の GPGPU では、16kBを共有する8個の演算ユニットSPEをまとめたSM
を30個用意しており、メインメモリにはキャッシュを通してアクセスする。
Intelの Nehalem では、256KBのL2キャッシュがあり、1 Chip 内の 4 Coreが8MB L3
キャッシュ を共有し、さらに、 Quick Path と呼ばれるInterconnectで
Chip 間を相互接続する。
個々のCore での処理は、256Kbyte程度にまとめるとキャッシュに収まることになる(図\ref{manycore})。

<center><img src="fig/ManyCore.eps" alt="manycore"></center>

並列処理では、Amdahl 則があり、部分的に並列化するのでは、全体的な速度の
向上の比率が低くなってしまう。プロセッサの数S、並列化の比率Pに対して、
処理時間は、$ (1-P)+(P/S) $となる。並列化が50\% では、どんなに多くの
プロセッサを投入しても効果は2倍程度になってしまう。従って、並列化率
を上げることが重要になる。

つまり、一般的なアプリケーションで、
for文等で部分的に並列実行するのでは十分な効果を上げることは出来ない。
しかし、科学技術計算や画像処理などでは、巨大なメインループがあるので、
並列に実行する処理自体は十分にある。

処理は、データと処理の組で規定する Task で構成する。Open/CLでは Kernel
と呼ばれている。

実際に問題になるのは、メモリの待ち時間である。Many Core では、
Local Store やキャッシュは限られていて、メインメモリから、
そこにデータを持って来る必要がある。これらのコピー時間の間、
プロセッサが待っているようだと、並列化率が落ちてしまう。

これを防ぐには、キャッシュやLocal Storeへのデータ転送の間に、
別な処理を行えば良い。つまり、ソフトウェアパイプライン的に
データ転送と計算処理を重ねて処理する。つまり、

   Many Core Programming = 
      Data Parallel +
      Pipe Line

と言う構成になる。東芝/SCEIでは、Cell 用に SPURS Engine \cite{spurs}という
フレームワークを提供している。
Apple の Open/CL\cite{opencl} では、Kernel と呼ばれるタスクをQueue
に入れることにより、パイプライン実行を可能にしている。

ここでは、キャッシュとLocal Storeを合わせて LS と呼ぶ。

---メモリのアクセスパターン

汎用的なプログラムでは、Local Store に閉じた処理を行うことは、ほとんどない。
従来のように、巨大なメモリをランダムアクセスするような処理では、
Many Core Programming の効率が上がることはない。

コードは、Task を実行している間、必ずLSに常駐する必要がある。これは、
前もって予測することは用意である。しかし、Task は細かく分割する必要
がある。64kbyte程度が望ましいが、これは経験的には5,000行程度であり、単一の
Task では、 足りなくなる
ことは少ない。しかし、全てのTaskのコードがLSに収まることはありえない。

順序良くアクセスされるデータは、処理をしながら並列にDMAや、キャッシュ
のアクセスなどでロードすれば良い。(図\ref{task})

<center><img src="fig/Task.eps" alt="task"></center>

問題は、処理によってアクセスする場所が変わる場合である。
描画処理する時のTexture画像はプログラムの実行時に初めて、
どこが必要かが判明する。これらは、キャッシュと同様の扱いをする
必要がある。Cell では、SPEがキャッシュを持たないので、これを
ソフトウェアで実現する必要がある。キャッシュを暗黙に扱う場合でも、
何らかの工夫が必要な場合が多い。

データの入力時だけでなく、データの書き出し場所が実行時に決まる
場合がある。データを決まった形でメモリ上に分散させなければ、
次のTaskに引き渡す処理が繁雑になる。この処理は、メモリ上を広く
アクセスする必要があり、SPEやGPGPUでは自分自身では処理できない
か、非常に遅くなる。この場合には、書込キャッシュを実装する必要
がある。

--Copy するかしないか

従来の逐次型プログラムでは、データのコピーはコストであり、
コピーをいかに減らすかが重要だった。並列処理が多用される
Many Core では、データは基本的に LS にコピーする必要が
ある。
計算主導なアプリケーションでは、CPUが動いている限り、
コピーのコストは見えない。実際、Cacheへの移動や、
DMA により隠されることになる。

時分割処理では、同期コストは、待ち時間中に他のTaskが走る
ので、それほど高くはない。しかし、Many Core の場合は、
Cache または、LS 中のTask の入れ換えを伴う。例えば、
Core 側から read / write のシステムコールを呼ぶことは、
明示的に待ってしまう。

待ち合わせや競合を防ぐには、コピーを取ることが簡単な
解決になる。実際、LSにコピーされて処理され、メインメモリ
に書き戻すという処理なので、自動的にコピーされることになる。

パイプライン処理は、命令レベル、Core レベル、Taskレベル
で行われるが、パイプラインバッファは読み書きペアが必要に
なる。同じ場所に書き戻す方法ではパイプライン処理を実現する
ことは出来ない。(図\ref{pipeline})

<center><img src="fig/pipeline.eps" alt="pipeline"></center>

つまり、オブジェクト指向プログラムをするとしても、ポインタを
IDとして持ち、インスタンス変数をその場で書き換えるような
方法は Many Core との相性は良くない。

--Scene Graph

ゲームプログラミングは、3Dグラフィックス処理と Real-time
プログラミングの組合せとなる。

ゲームプログラムで、繁雑な並列処理に集中するのは望ましくない。

一つの方法は、3Dグラフィックス部分をオブジェクトのノード
とする Scene Graph にすることである。ゲームプログラムは、
Scene Graph のノードの動作(Move)と、相互作用(collision)
からなる。これらと、動的な Scene Graph の構成により、
ほとんどのゲームプログラムを行うことが出来る。

Scene Graph には、Open Scene Graph や Java 3D などがある。
Scene Graph のノードには、視点(カメラ)、光源(複数)、
親子関係がある。Scene Graph が持つ情報には、3D のPolygon
と、そのnormal vector(Polygonの向き、曲面である場合は複数)、
そして、Polygon に貼られた2次元の画像(Texture)からなる。

Collision と Move を記述するだけで、十分な並列性を
ゲームプログラムから自動的に抽出するのが、Cerium Engine 
の一つの目標である。(図\ref{move})

<center><img src="fig/move.eps" alt="move"></center>




--Cell のアーキテクチャ

PS3 のCell では、メインメモリは256MBある。使用可能な6個のSPE 
にはLocal Storeが256Kbyteずつある。SPEには、128bit register
が128個あり、Local Store はキャッシュされていないがレジスタ
並に高速にアクセス出来る。LS はレジスタと考えて良い。(図\ref{cell})

SPEとPPEはリングバスで結ばれている。MFC(Memory Flow Controller)
に、32bit Mail によりコマンドを送ることにより、複数のDMAを
起動し、SPEとメインメモリ間での転送を行うことが出来る。Mail
は FIFO であり、SPEとPPEの通信に使うことも出来る。

DMAには、転送メモリの位置と数を指定するDMA と、複数の位置と
数をLS上に置き、それをまとめて転送するList DMAが用意されている。

<center><img src="fig/Cell.eps" alt="cell"></center>

--Cell B.E. 上のプログラミング


libSPE2 とは、PPE が SPE を扱うためのライブラリ群である \cite{libspe2} 。
libSPE2 は SPE Context Creation、SPE Program Image Handling、
SPE Run Control、SPE Event Handling、SPE MFC Problem State Facilities、
Direct SPE Access for Applications という基本構成でできている。
Cell の基本プログラムは次の様になる。

\begin{enumerate}
\item create N SPE context
\item Load the appropriate SPE executable object into each SPE context's local store
re
\item Create N threads
\item Wait for all N threads to terminate
\end{enumerate}

各SPEは、TLBを持っていて、起動したPPE のThreadが持つ仮想メモリ空間(64bit addressing) にDMAでアクセスすることが出来る。


\subsection{SPE C/C++ 言語拡張}
SPE では基本的な C 言語の機能の他に、Cell 特有の拡張が行われている
\cite{cell-cpp} 。
\tabref{tab:cell-cpp} に主な API を記す。

\begin{table}[htb]
\begin{center}
\caption{SPE C/C++ 言語拡張 API}
\label{tab:cell-cpp}
\begin{tabular}{|l|l|}
\hline
spu\_mfcdma32 & DMA 転送を開始する \\
\hline
spu\_read\_in\_mbox & PPE からの mail を取得する \\
\hline
spu\_write\_out\_mbox & PPE へ mail を送信する \\
\hline
spu\_add、spu\_sub、spu\_mul & SIMD 演算 (加算、減算、乗算) \\
\hline
\end{tabular}
\end{center}
\end{table}

gcc は、SPEのVector演算をサポートしており、\verb+__attribute__ ((vector_size (16)))+ を使用することにより適切な命令を出してくれる場合がある。
\verb+spu_add()+等で手動で処理することも可能である。


--Cerium Engine

Cerium Engine は、PS3 上のScene Graphの描画と変更によって
記述されたゲームを実行する Engine である。Task の生成を
管理する TaskManager, Task をScheduleする Scheduler,
ソフト的なキャッシュを管理する MemeorySegment そして、
SceneGraph ライブラリ、 複数の Task で実行される
Rendering Engine からなる。

C++ で記述されており、Mac OS X, Linux, PS3 上で同一の
ソースで動作する。

SceneGraph は、Blender (Open Source な3D modeling Tool)
から、Python Script を作って生成された XML である。
内部でも生成変更可能である。

SceneGraphのパラメータを変更することによりゲームが
進行し、それと同時に Pipeline 的に SceneGraphの描画を
行う。

Cerium Engine は、
Open CL とは独立に設計しているが、結果的には似たものになっている。

---Task Manager

Cerium の Task は、型のない入力と出力を持つシンプルな関数である。
入力と出力は、Cell のDMAによって用意される。List DMAがあるので、
断片化されたデータを読み書きすることも可能になっている。

create\_task() で、HTask という構造を作成し、spawn()することで
active queue に登録する。

Task は、番号で登録される。SPE と PPE では独立な番号を用いている。
手動による Overlay で、SPE 上のコードの入れ換えを行っているので、
ポインタで指定することは出来ない。Open/CL では、文字列で指定した
プログラムを llvm に引き渡す形式だが、その点は異なる。

Cerium は、PPEのTaskと、SPEのTaskの二種類を持っており、それぞれ
別なキューで管理されている。さらに、Task の終了時に呼ばれる
Task がある。

PPE Task は相互のDependencyを持ち、Dependency が満たされた時点で
PPEまたはSPE上で実行される。

普通のOSと異なり、高度なスケジューリング
は行われない。Task は十分小さく、基本的に preempt されずに
実行されるので、Round Robin などの工夫は不要である。通常の
スレッドやプロセスとは異なる。

SPE上には最小限のSchedulerが存在している。SPE上では基本的にTask
は生成しない。

SPE Task からはメインメモリは参照できないが、MemHash というLRU
キャッシュが実装されている。しかし、特に構文的な制限があるわけではな
く、PPE Task からはメインメモリ全域をアクセスすることが出来る。

一つのPPE Taskから生成されたSPE/PPE Taskは、まとめて投入される。
投入した Task は Cell の Mail を使って、SPE/PPE スケジューラに
伝えられる。終了した Task は、Dependency を解消し、Active に
なった Task を Active Queue に移動する。

DMAは、Taskの中から起動することも可能になっている。PPE や、Mac OS X
上でも、DMAのEmulationを行っており、同じコードで、PS3、Mac OS X
で動作させることが出来る。

HTask は、実行時は SchedTask と言う型のオブジェクトになる。
Cerium Engine に対する処理は、SchedTask へのメソッドで行われる。

---Task 間の同期

Task はメインメモリ上にQueueを持っており、

   task->wait_for(another_task)

という形で待ち合わせを行う。これは、task queue にwaiting queue
を持つこととで実現されている。実際の実行は、Core 上で行われる。
Core から見たメインメモリはかなり遠いので注意が必要である。

   task->set_post(post_task,in, out)

という形で、Taskの継続を指定することが出来る。継続のTaskでは、
DMA は行われない。主に、終ったタスクが書き出したデータの整理、
次のタスクの起動などに使われる。スケジューラの前に起動される
ので、継続で接続された Task のチェーンは、見掛け上、シングル
スレッドで実行される。

何もしないDummy taskを作り、それをwait\_for することにより、
同期を取るような手法が多用されることになる。これは、一種の
バリアである。

taskの終了は Mail による通知であり、Singe Thread なTask Manager
がMail 待ちループを持っていて処理する。

    do {
        ppeManager->schedule(ppeTaskList);
        ppeTaskList = mail_check(waitTaskQueue);
    } while (ppeTaskList || spe_running >0);


つまり、wait\_for/set\_post はメインメモリ上の
シングルスレッド Task として実行される。Task Manager を
複数のスレッド、あるいは分散したプロセスとして実装することも
可能だと思われるが、それが、wait\_for/set\_post の意味に
どう係わるかは、かなり難しいと思われる。現状では、シングル
スレッドが使いやすい。

---Task

一つの Task は、read/exec/write の三つに分解されて、
それぞれが、さらにパイプライン的に実行される。TaskManager
は、それが可能なように、Core に複数のTaskの集合(TaskList)
を投入する。

    void
    Scheduler::run()
    {
        task1 = new SchedNop();
        task2 = new SchedNop();
        task3 = new SchedNop();
        // main loop
        do {
            task3->write();
            task2->exec();
            task1->read();
            delete task3; 

            task3 = task2;
            task2 = task1;
            task1 = task1->next(this, 0);
        } while (task1);
        delete task3;
        delete task2;
    }

ここで、read/exec/write は Task の状態によって変わる
ステートパターン用の仮想関数である。何もしないもの、
DMAの開始/待ち合わせ、次のタスクの取得、次の TaskList
の取得などを行う。つまり、read/write は非同期で、
ほとんど待ち時間はないと想定されている。exec が
実際のユーザが定義したTaskの関数を呼び出す。

TaskList は、Core 側がEmptyになったのをMailでMain側に
知らせて、Main側から Mail によりTaskListへのポインタ
を取得して、DMAにより転送を行う。これは一つのTask
として実装されていて、上のパイプラインループの中で
実行される。


---Memory Segment Manager 

PPE上/SPE上のメモリは、MemorySegment によって管理されている。
これは、Hash access 可能な double linked されたメモリ領域
である。

これは、set\_global, get\_global によって、Core(SPE)上に
常駐する領域である。これは明示的に作成削除する必要がある。

Cerium Engine の性質上、ほとんどの処理は Pipeline Buffer 上で
行われる。つまり、 入力 MemorySegment から出力MemorySegment
に書き出される Task の集合である。

MemorySegment は読むだけ、あるいは書くだけとなる。そうでないと、
Pipeline 実行時に途中でデータを変更されてしまうことになる。

Taskの実行時には、必ずメモリのコピーを伴うし、DMAのオーバヘッド
のほとんどは隠されてしまうので、積極的にコピーする。メインメモリ
は、SPEのLSに比べれば広大なので、倍量のメモリを取っても問題ない。
Texture や Polygon のデータ等、変更されない場合はコピーの必要はない。

必ず Copy を伴うので、細かく解放する必要がない。したがって、
常に Copying GC を行っているようなことになる。これは、Pool を
多用する Apache Web Server と同じような実装となる。

--SPE 上のCode 

SPE上のコードは、256Kbyteしかないので、なんらかの方法で入れ換える
必要がある。SPE 上のライブラリは固定されているので、その部分以外
を relocatable にするようなコードをアセンブラから生成する。

コード自体は、Memory Segment Manager により、LRU/Hash で管理される。

現在は、gcc のOverlay code を流用しているので、そのままでgdb でデバッグも
可能であるが、Taskの単体テストが望ましい。

SPE上のプログラムは C++ だが、g++  の仮想関数の実装が relocatable
に出来ないので、SPE上の Task はメソッドではなく、Cの関数である。

Many Core 上でのオブジェクト指向的な利点はほとんどないので、
特に問題ないようである。

現在は、task\_list と言う配列に、そのまま、Taskの関数へのポインタrun
と、その関数を SPE 上へロードする関数をwaitを用いている。

    scheduler->dma_wait(DMA_READ);
    task_list[task->command].wait(scheduler,task->command);
    task_list[task->command].run(this, readbuf, writebuf);

この方法では、Task が増えると task\_list 自体が大きくなり、
メモリを圧迫してしまう。task\_list 自体は、メインメモリ上に
置く方法が望ましい。 現状では、Overlay code 自体がテーブル
持っているので、それは意味がない。おそらくは、Overlay code
自体を捨てる方が望ましいと思われる。


--Cerium Programming

Cerium  では、すべてを Task で記述する。処理は、

    int TMmain(TaskManager *manager, int argc, char *argv[])

から始まる。これは、SDLmain を真似ている。Cerium は SDL\cite{sdl} を使用している
ので、main が衝突しないようにこのようになっている。

ここでは、word count を例に取り上げる。まず、File を mmap によって
メインメモリにマップする。read で書いても良いが、待ち合わせが生じる
ので繁雑になる。mmap したメインメモリを分割し、Task に割り当てる。

分割された word count が出力する領域を確保する。

    int out_size = division_out_size*out_task_num;
    unsigned long long *o_data = 
      (unsigned long long*)manager->allocate(out_size);

まず、各SPEの結果を合計して出力するタスクを起動する。

    HTask *t_print = manager->create_task(TASK_PRINT);
    t_print->add_inData(o_data, out_size);
    t_print->add_param(out_task_num);
    t_print->add_param(status_num);

これで入出力が 確保されたので、word countする Task を起動する。

    /*渡すデータの最後。(スペース、改行以外)*/
    int word_flag = 0;
    int i;
    for (i = 0; i < task_num; i++) {

        t_exec = manager->create_task(TASK_EXEC);
        t_exec->add_inData(file_mmap + i*division_size, 
                   division_size);
        t_exec->add_outData(o_data + i*status_num, 
                   division_out_size);
        t_exec->set_param(0,division_size);
        t_exec->set_param(1,word_flag);
        t_exec->set_cpu(SPE_ANY);
        t_print->wait_for(t_exec);
        t_exec->spawn();

        word_flag = 
  ((file_mmap[(i+1)*division_size-1] != 0x20) 
  && (file_mmap[(i+1)*division_size-1] != 0x0A)); 
        size -= division_size;
    }

add\_inData/add\_outData/set\_param などで、Task に必要なパラメータ
を引き渡している。

実際には端数処理が必要だが省略する。 最後に、print のTaskを起動する。

    t_print->spawn();

各Taskの処理は以下のように記述する。

    /* これは必須 */
    SchedDefineTask(Exec);

    static int
    run(SchedTask *s, void *rbuf, void *wbuf)
    {
        char *i_data = (char*)s->get_input(rbuf, 0);
        unsigned long long *o_data = 
            (unsigned long long*)
               s->get_output(wbuf, 0);
        long length = (long)s->get_param(0);
        long word_flag = (long)s->get_param(1);
        int word_num = 0;
        int line_num = 0;

        word_flag = 0;

        for (int i=0; i < length; i++) {
            if((i_data[i] != 0x20) && 
                (i_data[i] != 0x0A)) {
                word_num += word_flag;
                word_flag = 0;
            }
            ...
        }
        word_num += word_flag;

        o_data[0] = (unsigned long long)word_num;
        o_data[1] = (unsigned long long)line_num;


get\_input/get\_output/get\_param などでTaskに必要なパラメータを取り出している。

集計 Task は簡単なので省略する。

--Word Count の問題点

この方法では、6台の SPE で走らせるよりも、PPE側で走らせる方が
高速になる。

    SPE上
    ./word_count -file a.txt -cpu 6 
    0.07s user 1.45s system 162\% cpu 0.938 total

    PPE上
    ./word_count -file a.txt -cpu 0 
    0.64s user 0.23s system 99\% cpu 0.872 total


まず、大量の Task を一気に起動することになるので、Task のデータ
量が多い。PPE でのTaskQueue に 線形リストを使っていると、
その処理に時間がかかる。
Task 終了時には、WaitQueue の削除などの線形リストではO(n)かかる
処理があるためである。

Double Linked List を使う以前は、
実際に、Task を少しずつ起動すると高速になる。
しかし、TaskQueue を Double Linked List にすると、その影響は
なくなり、一辺に大量に起動しても構わない。

次に、mmap したメモリに早めにランダムにアクセスしてしまう。これは、
ファイルに高速にアクセスするには、あまり望ましくない方法である。
しかし、ファイル全体がOSのキャッシュに入っている状態では、それほど
のペナルティは存在しない。

キャッシュに入っていない状態では
問題となる。ファイル全体を高速に複数のCoreに提供するようなAPIが
必要となる。

%//ファイル全域に、まず、word_flag を得る爲の処理により
%//アクセスしてしまっている。
%//mmap しても、それが並列処理に適した形でデータ転送されなければ、
%//正しく処理されない。

分割した際の端数処理と総計の処理は、別Taskで集計するが、
これは単一のTaskとなる。これを、分割集計と並列に走らせる
必要があるが、プログラミングは繁雑となる。

SPE上では、128bitレジスタが生かされるように vector 型の宣言を
する必要がある。

   typedef char *cvector 
       __attribute__ ((vector_size (16)));

などを使うと、

       lqx     $16,$20,$5     16byte 一括 load
       ceqb    $15,$8,$16     16byte 一括比較
       gbb     $3,$15         gather Bits from Bytes 

などのvector命令を gcc が生成することが出来る。この宣言は、
vector 命令をサポートしないアーキテクチャでも無害だが、
演算がvector同士で制限されるので注意が必要である。
ただし、spu-gcc の実装は、gcc 4.1.1 では、まだ正しくなく、
うまく動作しない。また、高速にもならない。
実際、Word Count のようなものだと、あまり、うまくvector 命令
に落ちないようである。

パイプライン処理には、パイプラインの切替えがあり、その段階で
並列度が落ちることがある。他の並行して走っているより高度な
パイプライン処理があれば、それは自動的に隠されるはずである。
しかし、Word Count のような例では、それを期待することは難しい。

これらは、ほとんど待ち時間となって現れるが、プログラム上、あるいは、
デバッガ上で、それを確認することは難しい。

Word Count が適切な例題ない(SPE向きではない)と言うのはあるが、
まだ、隠された Overhead が存在するのではないかと考えている。


--より複雑な例 (キャッシュへのアクセス)

zbuffer を用いて、描画を行う場合、描画に必要な Texture のデータを
メインメモリから取って来る必要がある。

{\small
 g->tileList = 
 (TileListPtr)smanager->global_get(GLOBAL_TILE_LIST);

 if (tex_z < g->zRow[localx + (rangex*localy)]) {
    memaddr tex_addr;
    int tex_localx;
    int tex_localy;

    tex_addr = getTile(tex_xpos, tex_ypos,
       span->tex_width, (memaddr)span->tex_addr);
    tex_localx = tex_xpos % TEXTURE_SPLIT_PIXEL;
    tex_localy = tex_ypos % TEXTURE_SPLIT_PIXEL;
    TilePtr tile = 
       smanager->get_segment(tex_addr,g->tileList);
    smanager->wait_segment(tile);

    updateBuffer(g, tex_z, rangex, localx, localy,
		 tex_localx, tex_localy,
		 normal_x, normal_y, normal_z, tile);
 }
}

ここでは、GLOBAL\_TILE\_LIST がSPE上に Task 間で共有されるメインメモリのキャッシュと
して確保されている。Texture を格納している tex\_addr (64bit)のアドレスが確定
すると、get\_segment() により、キャッシュにアクセスする。ここでは、
すぐに wait\_segment してしまっているが、本来は、少し時間があった方が良い。

memaddr は、メインメモリのポインタ型を表している。メインメモリが64bit/32bit
でも、SPEのアドレス空間は256kbyte(20bit)であり、一致しない。

Cerium では明示的にキャッシュを作成しているが、他のキャッシュをサポート
しているMany Coreの場合でも、待ち時間は生じるので、何らかの工夫は必要となる。

--より複雑な例 (明示的なDMA)

以下は、ポリゴンのデータから、Span (同じy座標を持つ直線)を抜き出す
処理に出て来る、SpanPack(Spanの集合)の書き出し部分である。

{\small
 if (charge_y_top <= y && y <= charge_y_end) {
    int index = (y-1) / split_screen_h;
    /**
     * 違う SpanPack を扱う場合、
     * 現在の SpanPack をメインメモリに送り、
     * 新しい SpanPack を取ってくる
     */
    if (index != g->prev_index) {
	tmp_spack = g->spack;
	g->spack = g->send_spack;
	g->send_spack = tmp_spack;

	smanager->dma_wait(SPAN_PACK_STORE);
	smanager->dma_store(g->send_spack, 
              (memaddr)spackList[g->prev_index],
	    sizeof(SpanPack), SPAN_PACK_STORE);
	
	smanager->dma_load(g->spack, 
            (memaddr)spackList[index],
	       sizeof(SpanPack), SPAN_PACK_LOAD);
	g->prev_index = index;
	smanager->dma_wait(SPAN_PACK_LOAD);
    }
}

ここでは、明示的に書き出し領域をDMAで、読み書きしている。この部分を
get\_segment で書き換えることも可能であるが、例として敢えて持って来ている。
dma\_wait が先行しているのは、前のdma\_store との処理がパイプライン的に
行われているためである。

この場合は書込処理なので、SPE間で同期を取ることが必要になる場合があるが、
そのような同期は、ここでは用意していない。

--Scene Graph

%% 書くの?

Scene Graph の処理(Move, Collision) と、Scene Graph のRenderingの
各ステージは、それぞれ、パイプライン化される。ポインタを使った
グラフ構造をそのまま使うと、Core にコピーした時に困ることになる。
したがって、Scene Graph は、コピーしながら生成する手法を取る。

Scene GraphのMoveは、Core 上でノードのプロパティを変更するだけ
で、必要なのはユーザ入力だけである。しかし、Collision の場合は、
$O(N^2)$で処理する必要がある。

Move/Collision は、ステートパターン、つまり、Move/Collision に
状態を表すオブジェクトあるいは関数、Cerium Engine では、
Task 番号を指定することになる。これらの Task は SPE 上で、
Memory Segment Manager によって管理される。



--比較

ここでは、SPURS Engine、Open CL、並列処理言語であるErlang/Scalaと比較してみる。

Cerium は、SPURS Engineの実装の一つであるが、SPURS Engine 自体の情報が
公開されていないので比較することは難しい。Cerium は、sourceforge.jp
上で公開されている。

Cerium では、SPE Task の終了を Mail により PPE に投げて、PPE側で Task
Queue の管理を行っている。Word countのような場合は、それが負荷になる
場合もある。SPURS Engine 等で、どのような工夫が行われているのかは
未知数である。

Open CLは、Taskを登録する代わりに、文字列として渡し、Open CL側で、
GPGPUやThreadに展開される。Mac OS X では、llvm \cite{llvm} を用いて
展開されている。Cerium では、前もってコンパイルされた関数をマクロ
により表に登録する仕組みである。

Kernel は、

    __kernel void UniformAddKernel(
        __global float *output_data, 
        __global float *input_data, 
        __local float *shared_data,
        const uint group_offset, 
        const uint base_index,
        const uint n)
    {
        const uint local_id = get_local_id(0);
        const uint group_id = get_global_id(0) 
            / get_local_size(0);
        const uint group_size = get_local_size(0);

と言う形を持っている。globalがメインメモリ上のデータで、localが
Local Store のデータである。

データのload/storeは、Kernel として明示的に記述し、

{\small
    unsigned int k = PRESCAN_NON_POWER_OF_TWO;
    clSetKernelArg(ComputeKernels[k],  
         a++, sizeof(cl_mem), &input_data);
    clEnqueueNDRangeKernel(ComputeCommands, 
               ComputeKernels[k], 1, 
         NULL, global, local, 0, NULL, NULL);
}

等として、自分で Queue を管理する。実行 Queue に自分で格納する 
Open CL の方がやや繁雑な記述となる。

Kernel 間の依存関係は、Queue で解決されているので、起動する
メインルーチン側で処理することになる。

Kernelの中では、\_\_globalと言う形で、いつでもメインメモリにアクセスする
ことが可能である。なので、明示的なDMAは必要ない。しかし、見えないだけで
コストや待ち時間は生じてしまう。そこで、Pre Scan のような形で、
データを \_\_local に前もって持って来ておく必要がある。

---Erlang, Scala

Erlang と Scala は、Actor に似た感じで並列処理を行う。Erlang は、
Prolog に似た構文を持っており、Scala は Java 上に実装されている。

双方の言語とも通信はチャネルで行われて、Task 上のデータは関数型
言語的な意味で変更されない。再帰的な関数呼び出しにより、Taskの
状態を作るので、両方とも似たような(構文は隨分違うが)並列プログラミング
のスタイルを提供している。

Erlang

    ping(N, Pong_PID) ->
        Pong_PID ! {ping, self()},
        receive
            pong ->
                io:format("Ping received pong~n", [])
        end,
        ping(N - 1, Pong_PID).

Scala
    class Counter extends Actor
    {
      override def act(): Unit = loop(0)
 
      def loop(value: Int): Unit = {
        receive {
          case Incr()   => loop(value + 1)
          case Value(a) => a ! value; loop(value)
          case Lock(a)  => a ! value
	      receive { case UnLock(v) => loop(v) }
          case _        => loop(value)
        }
      }
    }



Cerium では、データの引き渡しは、void *にcastするので、型の
安全性を言語上で保証することは出来ない。Open CL でも状況は
同じで、

    __kernel void UniformAddKernel(
        __global float *output_data, 
        __global float *input_data, 
        __local float *shared_data,
        const uint group_offset, 
        const uint base_index,
        const uint n)

が、
{\small
    err |= clSetKernelArg(ComputeKernels[k],  
         a++, sizeof(cl_mem), &output_data);  
    err |= clSetKernelArg(ComputeKernels[k],  
         a++, sizeof(cl_mem), &partial_sums);
    err |= clSetKernelArg(ComputeKernels[k],  
         a++, sizeof(float),  0);
    err |= clSetKernelArg(ComputeKernels[k],  
         a++, sizeof(cl_int), &group_offset);
    err |= clSetKernelArg(ComputeKernels[k],  
         a++, sizeof(cl_int), &base_index);
    err |= clSetKernelArg(ComputeKernels[k],  
         a++, sizeof(cl_int), &n);
    err |= clEnqueueNDRangeKernel(ComputeCommands, 
      ComputeKernels[k], 1, NULL, 
           global, local, 0, NULL, NULL);
}

などに相当することになる。これらの型を(IDE等で)チェックすることは、難しくはないが、
Scala のよう言語自体がチェックする方が望ましい。

Erlang, Scala は、GCを持っているが、Open CL/Cerium では、明示的なメモリ管理
を行う必要がある。

パイプライン処理をうまく動作させるには、Taskの列(Queue)を作る必要があるが、Erlang、
Scala では、その列を明示的に作ることは強制されていない。分散処理、あるいは、
Thread による並列処理には、Erlang, Scala が適しているが、Many Core 
では、さらになんらかの構文的なサポートが必要だと考えられる。


--まとめ

Many Core でのプログラムは、科学技術計算等に重要であるが、
Desktop PC や、Note PC での性能向上に使用するには、
「普通の」プログラムでも並列処理を恒常的行う必要がある。

Cerium は、Open CL と同様に、プログラムTask に分割し、
Core に投入することによって実行する。

Cerium は、Open CL よりは若干ましな記述が可能となっている。

Many Core ではコピーが頻繁に行われ、そのコピーを隠す
パイプライン実行が重要である。コピー自体は必須となる。

プログラムのチューニングとデバッグは自明ではない。

今後は、Task の単位として、Continuation based C の
code segment を用いた実装を行う予定である。