一方、どうやって2種類(CPUとAccelerator)のコードを混在するか、であるが、OpenMPのフロントエンドを拡張して、ソース中に埋め込む形をとる(Photo22)。案外に驚きがない、というか割と当たり前な方法だが、違うのはここからである。

Photo22:"#pragma omp"(OpenMPのPragma宣言)内にAcceleratorのNative Codeを記述し、埋め込む形。一言で言ってしまえば、単なるプリプロセッサである。

このFat binaryをロードしたローダは、内部でこれを分離してx86 BinaryとAcceleratorのBinaryをそれぞれに割り振って実行する形になる(Photo23)。このローダ部をOSが管理するのか、それともApplicationもしくは外部のRuntimeが実施するのかは現状はっきりしないが、一番可能性がありそうなのはApplication内部に埋め込まれたRuntimeであろう。

Photo23:PragmaでAccelerator Native Commandを埋め込み、コンパイル・リンクしたものは、最終的にx86のBinaryとAccelerator NativeのBinaryを含む"fat binary"となる。fat binaryといえば、昔のMacOSを思い出すが、昔は同じ処理を行うPowerPC用と68K用のBinaryを2重持ちし、実行時にどちらを使うか選ぶというものだった。対して今度は別々に動く2種類のBinaryを持ち、どちらも使われる訳でちょっと意味合いが違う気がする。

ところでこのローダ、単にApplication BinaryをAcceleratorに投入するだけではない。これはPhoto24を見ると判る。これはGXA3000(IntelのG965に内蔵されたグラフィックコア)を叩くサンプルである。#pragma omp parallel targer(x3000)以下に続くのが、Accelerator側で使われる変数列、__asmの下が実際のAccelerator用のコードとなるが、例えば14行目であれば、

(1) GXA3000に shl.1.w vr1=i, 3という命令を投入する
(2) 投入した命令の完了を待つ

の2つが必要になる。つまりAccelerator Binaryの実行の同期も一緒にとる必要がある。実際にこれを内部でどう行うかはまだはっきり言えないが、このケースでは__asm{}のセクションの完了のタイミングで、恐らくはSemaphoreなどで同期をとる形になると思われる。ここで役に立つのが、先にPCI Express Protocol Extensionの処で出てきたAtomic Read-Modify-Writeである。Root Complexに置いたSemaphoreを使って同期を取ることで、より少ないレイテンシで実現できるようになる。この例だと、所要Cycle数は非常に少ないと思われ、先のPhoto10のグラフで言えば一番左端に近い状態になる筈で、なるほどこうした工夫が無ければオーバーヘッドが大きくなりすぎる訳である。

Photo24:まとめて同期を取るのか、それとも毎ステートメントに必ず同期を取るのか、は当然作業によって異なってくる。明示的な説明はまだだが、Photo27などと見比べると、__asm{}でまとめたブロックはそれ全体をAtomicとして扱って同期を取り、そうでない場合はステートメント毎に同期をとるというあたりではないかと思う。もっともこれはあくまでもPseudo Codeなので、実際はどこで同期を取るかまでプログラム中で記述する方向になっても不思議ではないし、現実的にはその方が好ましい気もする。

開発に当たっては、既存のプログラムの範囲を崩さないために、Visual Studioなどの開発ツールをそのまま利用できる、としている(Photo25)。実際、今回の場合で言えばPreProcessorでExoskeleton拡張部をまとめて処理するから、こうしたものと親和性は悪くない。ただ、fat binaryをどうやってVisual Studioなどで扱うか? というのは疑問が残るところ。実際にはfat binaryのうちAccelerator binaryとAccelerator dataについては、x86 binaryのdata sectionに押し込めてしまい、Debuggerからは単なるデータとしか見えないような工夫がなされているのかもしれない。セッションでは、コンパイル/リンクのサンプルなども示されており、開発もシームレスであることが示された(Photo26)。ちなみにこの環境だと、Visual StudioにIntel C++を組み合わせているが、このIntel C++がGXA3000 Binaryを吐くような特殊バージョンのようである。

Photo25:x86 BinaryはVisual Studioで、GMA3000 BinaryはEclipse環境で動作させているサンプル。いずれにしても、Accelerator向けに何かしらの開発環境を用意し、それを既存のVisual Studioなどと組み合わせられる様にカスタマイズする手間は避けられないのは仕方ないところか。

Photo26:謎なのは、linearfilter.exeを動かす前に起動する、謎の"run"コマンド。これがひょっとしてローダとして常駐して、fat binaryをx86 binaryとGXA3000 binaryにDispatchする作業を担っているのかもしれない。

CrearSpeedのAcceleratorを使った例がPhoto27に示されたが、こんな具合にx86 BinaryとCrearSpeed Binaryが渾然とするケースまでExoskeletonはカバーすることを考えているようである。

Photo27:このケースで言えば、Array1/2/3はClearSpeed側の領域となる。まずmemcpym2p()(これはMemory→Accceleratorの転送だろう。恐らく処理はClearSpeedで行われる)を使い領域をまとめて転送後、ClearSpeedでArray3=Array1+Array2という配列の加算をまとめて行い、結果を再びmemcpyp2m()を使ってメモリに戻すという仕組みである。ただint i=0;は(宣言だけしつつ、内部では使って無いが)x86 binaryであり、こうした混在も出来るという話。このケースでは、ステートメント毎にClearSpeedと同期を取っているように思われる。