一方、どうやって2種類(CPUとAccelerator)のコードを混在するか、であるが、OpenMPのフロントエンドを拡張して、ソース中に埋め込む形をとる(Photo22)。案外に驚きがない、というか割と当たり前な方法だが、違うのはここからである。
このFat binaryをロードしたローダは、内部でこれを分離してx86 BinaryとAcceleratorのBinaryをそれぞれに割り振って実行する形になる(Photo23)。このローダ部をOSが管理するのか、それともApplicationもしくは外部のRuntimeが実施するのかは現状はっきりしないが、一番可能性がありそうなのはApplication内部に埋め込まれたRuntimeであろう。
ところでこのローダ、単に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のグラフで言えば一番左端に近い状態になる筈で、なるほどこうした工夫が無ければオーバーヘッドが大きくなりすぎる訳である。
開発に当たっては、既存のプログラムの範囲を崩さないために、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を吐くような特殊バージョンのようである。
CrearSpeedのAcceleratorを使った例がPhoto27に示されたが、こんな具合にx86 BinaryとCrearSpeed Binaryが渾然とするケースまでExoskeletonはカバーすることを考えているようである。