キーポイント
- TornadoVM is a programming and execution framework for offloading and running JVM applications on heterogeneous hardware (multi-core CPU, GPUs and FPGAs)
- TornadoVM extends the Graal JIT compiler with a new backend for OpenCL
- Applications written for TornadoVM are single-source - the same code is used to express the host code and the accelerated code
- TornadoVM can perform live-task migration across computing devices
3月に私は、Qcon-LondonでTornadoVMについて講演して、TornadoVMの紹介とともに、その動作について説明しました。この記事では、Qcon Londonでの講演の内容を拡張して、異種(heterogeneous)ハードウェア上でJavaを自動的に実行することによって開発者が期待できるメリットについて、さらに詳しく説明したいと思います。
最初に、TornadoVMプロジェクトとアーキテクチャの概要について説明します。次に、TornadoVMのさまざまな部分について、実際的な例をあげて解説します。
TornadoVMはなぜ必要か？
あらゆるタイプのワークロードを最高の効率で実行できるような、単一のコンピュータアーキテクチャというものは存在しません。近年の異種ハードウェアの急増はここから来ています。従って、私たちがプログラミングするすべてのシステムに、コンピュータ要素が混在している可能性があるのです。
これらの要素はそれぞれ、異なったハードウェア特性を持っています。ハードウェアの異種性はプログラマにとって、アプリケーションのパフォーマンス改善とエネルギ消費の低減をもたらしてくれます。
このようなコンピューティングにおける新しい異種デバイスの中には、マルチコアCPU、GPU(Graphics Processing Unit)、FPGA(Field Programmable Gate Array)などがあります。多様性はすばらしいことですが、これら新しいデバイスを効率的にプログラムする方法が必要です。
その最たるものとして挙げられるのが、ヘテロジニアス(異種)プログラミング言語の代表格である、CUDAとOpenCLの２つです。しかしこれらの言語では、APIに低レベルな機能が現れているため、専門家でないユーザが使うには非常に難しいものになっています。その例として、OpenCL 3.0標準の中から次の一節を引用してご紹介しましょう。
OpenCLが対象とするのは、ポータブルかつ効率的なコードの記述を望む専門的プログラマです。[...] 従ってOpenCLでは、低レベルのハードウェア抽象化とプログラミングをサポートするフレームワークを提供し、基盤となるハードウェアの詳細部分の多くを公開しています。
この文章の内容は、CUDAや同種の並列(parallel)プログラミングモデルにも当てはまります。一方、産業界や学術界の開発者たちには、低レベルのプログラミング言語を使う代わりに、高レベルのオブジェクト指向プログラミング言語を使用する傾向があります。マネージドランタイム環境で動作するJavaやR、Python、JavaScriptなどはその典型です。このようなプログラミング言語ならば、異種ハードウェア上での透過的実行に対応できているのでは、と期待するプログラマは多いかも知れませんが、実際には極めて限定的か、あるいはまったくサポートされていないのが現実です。
今回の記事では、ヘテロジニアスコンピューティング用の低レベルな並列プログラミング言語に代わるものとして、TornadoVMについて探っていきます。ご存じのように、並列コンピューティングアーキテクチャや並列プログラミングモデルについての知識がなくても、開発者はマルチコアCPUやGPUを活用することができます。
簡単に言うとTornadoVMは、透過的かつ動的にJavaバイトコードをOpenCLにオフロードし、生成したコードを異種ハードウェア上で実行することの可能なJVM言語用の並列プログラミングフレームワークです。さらにTornadoVMは、最適化ランタイムの統合、デバイスバッファの再利用とデバイス間の転送データ保存が可能、コンピューティングデバイス間でのライブタスクマイグレーションを実施する最新の動的アプリケーション再構成コンポーネント、といった特徴を備えています。
では始めましょう！
次の図は、TornadoVMプロジェクトの概要を高レベルで示したものです。ご覧のようにTornadoVMは、多層マイクロカーネルソフトウェアアーキテクチャで構成されています。その中でコアに当たるのがTornadoVM実行エンジン(execute engine)です。最上位レベルでは、開発者にAPIを公開しています。これは現在のTornadoVMが並列性を検出する(自動並列化/auto-parallelization)のではなく、並列性を利用する方式であるためです。そのためTornadoVMでは、どのメソッドあるいは関数がGPUおよびFPGAでの実行対象なのか、指定する方法が必要なのです。
さらにTornadoVMはコアランタイムを持っています。コアランタイムは、a) 新たなバイトコードジェネレータを備えたデータフローオプティマイザ b) 新しいバイトコードを実行するコンパクトなバイトコードインタプリタ c) JITコンパイラとメモリ管理、というコンポーネントに分割されています。今回の記事では、API、ランタイム、およびJITコンパイラの概要を紹介したいと思います。
最後に、先程の図に示されているように、TornadoVMは、現時点では最新のJDK(u242)とJVMCIを使用したJava 8、OpenJDK 11、GraalVM 19.3.0をサポートしています。Tornado VMはOpenCL 1.2とも互換性があるため、GPU(AMD、NVIDIA)、FPGA(Xilinx、Intel)、統合GPU(Mail ARM、Intel HD Graphics)や、マルチコアCPUといった、さまざまなデバイスセット上で実行可能です。
実際のTornadoVM
実例をもとに、詳細を見ていきましょう。ここで説明するのは、TornadoVMを使ってマルチコアCPU、GPU、統合GPU上で行列計算をプログラムし、実行する方法についてです。行列計算は、簡単なコードでTornadoVMのさまざまな概念の説明を始められる上に、多くのマシンラーニングやディープラーニングアプリケーションで使用されているものです。
注記: TornadoVMはJavaでプログラムされていますが、その計算カーネルは、GraalVMの提供するPolyglotフレームワーク(Truffle)を経由することで、他のJVM言語でも使用することが可能です。
以下のコードスニペットは、Javaでプログラムされた行列の乗算を示すものです。
class Compute {
public static void matrixMultiplication(final float[] A, final float[] B, final float[] C, final int size) {
for (int i = 0; i < size; i++) {
for (int j = 0; j < size; j++) {
float sum = 0.0f;
for (int k = 0; k < size; k++)
sum += A[(i * size) + k] * B[(k * size) + j];
C[(i * size) + j] = sum;
}
}
}
}
このコードスニペットは、GPUコンピューティングの古典的かつ標準的な行列乗算の例です。このコードスニペットをTornadoVMで高速化するため、まず最初に、並列化可能なループをアノテートします。このケースでは、最外周の２ループのイテレーションには依存性がなく、完全な並列化が可能なので、TornadoVMのアノテーション@Parallelを使って、次のようにコードをアノテートします。
class Compute {
public static void matrixMultiplication(final float[] A, final float[] B, final float[] C, final int size) {
for (@Parallel int i = 0; i < size; i++) {
for (@Parallel int j = 0; j < size; j++) {
float sum = 0.0f;
for (int k = 0; k < size; k++)
sum += A[(i * size) + k] * B[(k * size) + j];
C[(i * size) + j] = sum;
}
}
}
}
@Parallelアノテーションは、TornadoVM JITコンパイラ(JavaバイトコードをOpenCLに変換する)がヒントとして使用します。
TornadoVM JITコンパイラは強制的に並列化を実行するのではなく、アノテートされたループが並列化可能であるかをチェックした上で、forループから等価なOpenCLの並列インデックス(parallel indexing)への置き換え(get_global_id(dimension))を行います。forループが並列化不能な場合は、何もせずにシーケンシャルなコードを実行します。
高速化可能なJavaメソッドも指示する必要があります。そのためにTornadoVMは、軽量なタスクベースのAPIを公開しています。このAPIは高速化するメソッドのリストをセットするためのもので、それぞれのメソッドがタスクに対応します。タスクスケジューラを通じて、タスクのグループを生成することができます。以下のコードスニペットでは、行列乗算の例でタスクスケジュールを生成する方法を示しています。
TaskSchedule t = new TaskSchedule("s0")
.task("t0", Compute::matrixMultiplication, matrixA, matrixB, result, size)
.streamOut(result);
ここでは、タスクスケジュールオブジェクト(t)を生成しています。コンストラクタには、タスクの名称として任意の文字列を渡します。この名称は、タスクを実行するデバイスを変更する場合に使用することができます。次に、タスクセットを定義します。この例ではひとつだけですが、タスクの数は任意です。
タスクに渡すパラメータは次のとおりです — まず名前(この場合は"t0")と、高速化したいメソッドへの参照(この場合は、JavaクラスComputeのmatrixMultiplicationメソッドを示しています)もパラメータとして渡します。その他のパラメータは、そのメソッドの実際のパラメータセットに対応します。
そして最後に、ホスト(CPU)と同期したい変数ないし配列を指示します。これが必要なのは、GPUやFPGAは一般的に、CPUと同じメモリを共有していないためです。そのためTornadoVMランタイムは、ターゲットデバイス上の全変数のためのスペースを確保して、ホスト(CPU)からデバイス(GPUなど)へのデータ転送を行います。その上で、最終的な結果を取得するために、TornadoVM APIコールのstreamOutを通じて変数リストを同期するのです。
ここまでで、タスクを宣言して、並列処理が可能なコードにそれを置くところまできました。アプリケーションをTornadoVMで実行するには、TaskScheduleオブジェクトの"execute()"メソッドをコールする必要があります。
これはブロッキングコールで、OpenCLバッファを生成し、実行グラフを作成し、すべてのタスクをJavaバイトコードからOpenCLにコンパイルした上で、生成したOpenCLプログラムをターゲットデバイス上で実行します。さらにTornadoVMでは、多数のメソッドを同時にコンパイルしてひとつのコンパイルユニットに統合し、同じデバイス(同じGPUなど)上で実行することも可能です。これによって、ホストと、通常は(AMD APU、ARM Mail、Intel HD Graphic GPUといった統合GPUでない限り)プライマリホストとのメモリ共有を行わない異種デバイスとの間のデータ転送を最適化する機会が生まれるのです。
ソースコード内には、デバイスを特定する情報が一切含まれていない点に注目してください。同じコードをマルチコアCPU、GPU、FPGA上で共有することができるのです。TornadoVMランタイムとJITコンパイラは、アーキテクチャに応じてコードを自動的に最適化します。
それではコード例を実行してみましょう。まずは、TornadoVM環境のセットアップ方法を紹介します。これらの例はすべてGitHub上のリポジトリにあります。
行列乗算の実行: TornadoVMの設定
ここではTornadoVMを、Graal 19.3.0をSDKに使用して実行します。私たちがGraalのバージョンを頻繁にアップデートしていることに注意してください。Graal 20.xのTornadoVMへの統合は、今年末を予定しています。コードの実行には、OpenCLがインストールされていることが前提になります。すべての要件をこちらで確認しておいてください。
$ mkdir -p TornadoVM
$ cd TornadoVM
$ wget https://github.com/graalvm/graalvm-ce-builds/releases/download/vm-19.3.0/graalvm-ce-java11-linux-amd64-19.3.0.tar.gz
$ tar -xf graalvm-ce-java11-linux-amd64-19.3.0.tar.gz
$ export JAVA_HOME=$PWD/graalvm-ce-java11-19.3.0
$ git clone --depth 1 https://github.com/beehive-lab/TornadoVM
$ cd TornadoVM
$ export PATH=$PWD/bin/bin:$PATH
$ export TORNADO_SDK=$PWD/bin/sdk
$ export CMAKE_ROOT=<SET YOUR PATH TO CMAKE ROOT>
$ make graal-jdk-11
$ export TORNADO_ROOT=$PWD
例の含まれているリポジトリをダウンロードしましょう。
$ git clone https://github.com/jjfumero/qconlondon2020-tornadovm
$ cd qconlondon2020-tornadovm/
$ export JAVA_HOME=/path/to/graalvm-ce-java11-19.3.0
$ export PATH="${PATH}:${TORNADO_ROOT}/bin/bin/" ## Defined previously
$ export TORNADO_SDK=${TORNADO_ROOT}/bin/sdk
$ export CLASSPATH=target/tornado-1.0-SNAPSHOT.jar
$ mvn clean install
これで、例を実行する条件がすべて整いました。まず最初に、TornadoVMで確認できている、利用可能なデバイスを探索してみます。
$ tornado --devices
Number of Tornado drivers: 1
Total number of devices : 3
Tornado device=0:0
NVIDIA CUDA -- GeForce GTX 1050
Global Memory Size: 3.9 GB
Local Memory Size: 48.0 KB
Workgroup Dimensions: 3
Max WorkGroup Configuration: [1024, 1024, 64]
Device OpenCL C version: OpenCL C 1.2
Tornado device=0:1
Intel(R) OpenCL -- Intel(R) Core(TM) i7-7700HQ CPU @ 2.80GHz
Global Memory Size: 31.0 GB
Local Memory Size: 32.0 KB
Workgroup Dimensions: 3
Max WorkGroup Configuration: [8192, 8192, 8192]
Device OpenCL C version: OpenCL C 1.2
Tornado device=0:2
Intel(R) OpenCL HD Graphics -- Intel(R) Gen9 HD Graphics NEO
Global Memory Size: 24.8 GB
Local Memory Size: 64.0 KB
Workgroup Dimensions: 3
Max WorkGroup Configuration: [256, 256, 256]
Device OpenCL C version: OpenCL C 2.0
私の場合は、使用しているラップトップ上に3つのデバイス — NVIDIA GPU、IntelマルチコアCPU、Intel HD Graphics(統合GPU)があります。TornadoVMは、既定値としてデバイス０を選択しますが、タスクをデバイスに関連付けることで変更が可能です。ここではデフォルトの設定で始めることにしましょう。
$ tornado qconlondon.MatrixMultiplication 512 tornado
このプログラムは、行列乗算メソッドを100回実行して、合計時間を繰り返し毎に表示します。このメソッドは、何ができるかを示す単純なサンプルに過ぎません — この後、JMHを使って適切なパフォーマンス比較を行ってみます。
$ tornado qconlondon.MatrixMultiplication 512 tornado
Computing MxM of 512x512
Total time: 77568790 (ns), 0.0776 (s)
Total time: 3133182 (ns), 0.0031 (s)
Total time: 3126146 (ns), 0.0031 (s)
…
最初の実行が2回目以降より長くかかっていることに注目してください — これはJITコンパイルのウォームアップによるもので、JMHを使用する効果によって解消します。
最初にタスクスケジュールを実行すると、TornadoVMがOpenCL JITコンパイラを起動して、最適化と、JavaバイトコードからのOpenCL Cコード生成を行います。コードが生成されれば、そのコードをコードキャッシュにインストールして、ランタイム間の任意のポイントで同じタスクが再実行されれば、そのバイナリが再利用されるのです。TornadoVMがGPU(デバイス0)上で動作していることを確認するために、次のようなデバッグ情報を有効にすることができます。
$ tornado --debug qconlondon.MatrixMultiplication 512 tornado
Computing MxM of 512x512
task info: s0.t0
platform : NVIDIA CUDA
device : GeForce GTX 1050 CL_DEVICE_TYPE_GPU (available)
dims : 2
global work offset: [0, 0]
global work size : [512, 512]
local work size : [32, 32, 1]
よろしい、TornadoVMは行列乗算のJavaコードを、NVIDIA GTX 1050上で実行しています。比較のために、シーケンシャルなアプリケーションも実行してみましょう。そのためには、コードを高速化するためのTornadoVM JITコンパイラを実行しなければよいので、パラメータをもうひとつ追加してそれを指示します。
$ tornado qconlondon.MatrixMultiplication 512 sequential
Computing MxM of 512x512
Total time: 259398036 (ns), 0.2594 (s)
Total time: 247857535 (ns), 0.2479 (s)
...
TornadoVM JITコンパイラが動作する初回でも、3.3倍高速であることが分かります。2回目以降は、Javaのシーケンシャルコードよりも80倍以上高速になっています。ただしこの数値は、あくまでも参考値であることに注意してください。次のセクションでは、Java JMHを使用したパフォーマンス比較を導入します。
デバイスの変更方法は？
アプリケーションを実行するデバイスは、コマンドから変更することができます。例えばIntel Integrated Graphicsを使用するには、次のようなオプションで起動することが可能です。
$ tornado -Ds0.t0.device=0:2 --debug qconlondon.MatrixMultiplication 512 tornado
Computing MxM of 512x512
task info: s0.t0
platform : Intel(R) OpenCL HD Graphics
device : Intel(R) Gen9 HD Graphics NEO CL_DEVICE_TYPE_GPU (available)
dims : 2
global work offset: [0, 0]
global work size : [512, 512]
local work size : [16, 16, 1]
構文は次のとおりです -D<taskScheduleName>:<taskName>.device=0:<deviceIndex>
Dell XPS 15ラップトップ上でMxMの行列乗算を行った場合の、TornadoVMのパフォーマンス
これらのオプションを使えば、パフォーマンスの結果はすぐに得られます。次の図は、TornadoVMをさまざまなOpenCLデバイス上で実行した場合の、Javaのシーケンシャルな実装に対する速度の向上を示したものです(高いほどよい値です)。挙げられている速度向上(speedup)の値は、Java JMHフレームワークをベンチマークとして使用した平均値に対応しています。速度向上が大きいため、Y軸は対数スケールで表示されている点に注意してください。JMHを使用したベンチマークはすべて、例と同じリポジトリに置いてあります。この図から分かるように、マルチコアCPUでTornadoVMを使用することにより、Java Hostpotに比較して最大で3.6倍の向上が実現できています。GPU上で実行した場合には、Javaに比較してIntel HDグラフィックスで最大39倍、NVIDIA 1050で最大270倍を達成しています。
実行モデルとコンパイル
ここまででは、TornadoVM APIと、TornadoVMを使ってユーザレベルでアプリケーションを実行する方法について、簡単に説明を行ってきました。次にもう少し詳しく、TornadoVMがターゲットデバイス上でコードを実行する方法について見ていきましょう。
次に示す図は、JVMとTornadoVM間の実行フローを表現したものです。
タスクスケジュールの定義とTornadoVM APIからの実行メソッドの起動は、ひとつのJavaスレッド(マスタスレッド)上で実行されます。実行メソッドはブロッキングコールなので、メソッド実行から戻った時点で、並列デバイスでの実行完了が保証されています。実行メソッドが起動された時、TornadoVMが最初に行うのは、タスクスケジュール内のさまざまなタスク間でデータがコミュニケーションされる方法を表現する、データフローグラフの構築です。このグラフは、データ転送を最適化するために使用されます。
次にTornadoVMは、新しいバイトコード(COPY_IN、LAUNCH、COPY_OUT、BARRIERなど、ターゲットデバイス上での実行をオーケストレーションするための簡単なインストラクション)を生成します。コードが(LAUNCHバイトコードによって)最初にローンチされた時は、OpenCL JITコンパイラが起動されて、各タスク(速度向上の対象となる各Javaメソッド)から入力されたJavaバイトコードを、最適化されたOpen CL Cコードに変換します。
TornadoVMの生成するOpenCL Cコードは、ターゲットデバイスに応じて特化されています。すなわち、GPU用に生成されるコードは、CPUやFPGAのものとは違います。これは、OpenCLのコードはデバイス間で互換性があるが、パフォーマンスについては同じではない、という事実に基づくものです。そのためTornadoVMでは、デバイス毎に特化した、異なる最適化を適用することによって、パフォーマンスの向上を図っているのです。
注記: TortnadoVM JITコンパイラは単一スレッドで動作します。そのため、HotSpotでも経験しているように、重い負荷の下での動作ではコンパイラのリソースが枯渇する可能性がある、という問題があります。
コンパイルの最後のステップでは、最適化とデバイス特化の行われたOpenCL Cコードをターゲットプラットフォーム用にコンパイルする処理が、OpenCLドライバの起動を通じて実行されます。例えば、アプリケーションがNVIDIA GPU上で実行されるならば、このステップは対応するPTXコードを生成します。
OpenCLコードが生成、コンパイルされれば、TornadoVMがアプリケーションをターゲットデバイス上でローンチします。そのためにTortnadoVMは、カーネルを実行するスレッドを多数デプロイしています。デプロイするスレッドの数は、アプリケーションの入力数とハードウェアの特性によって決まります。
先に示した行列乗算の例であれば、512×512のスレッドのブロックを使ってデプロイされます。つまりTornadoVMは、元々プログラムされた単一スレッドのJavaアプリケーションから、512×512スレッドのブロックをデプロイする、ということです。ターゲットデバイスがマルチコアCPUの場合は、使用可能なCPUコアの最大値と同じ数のスレッドをデプロイします。
並列デバイス上の実行が終了すれば、TornadoVMは結果をJavaのヒープに(ホスト側から見えるように、バイトコードCOPY_OUTを使って)コピーした上で、VM内のマスタスレッドにコントロールを戻します。
アプリケーション毎にTornadoVMが生成したバイトコードを問い合わせることができます。例えば次のコードスニペットは、TornadoVMバイトコードのデバッグ情報付きで行列乗算を実行した場合の出力の一部です。
$ tornado --printBytecodes qconlondon.MatrixMultiplication 512 tornado
vm: COPY_IN [F@3e694b3f on NVIDIA -- GeForce GTX 1050
vm: COPY_IN [F@397fbdb on NVIDIA -- GeForce GTX 1050
vm: COPY_IN [F@33d512c1 on NVIDIA -- GeForce GTX 1050
vm: LAUNCH task s0.t0-matrixMultiplication on NVIDIA -- GeForce GTX 1050
vm: STREAM_OUT_BLOCKING [F@33d512c1 on NVIDIA -- GeForce GTX 1050
私たちが前に取り入れた行列乗算メソッドは、3つのパラメータ(行列A,B,C)を受けるようになっていました。それぞれの変数に対して、TornadoVMは、ホストからデバイスへのデータ転送(COPY_IN)を行います。その次に、LAUNCHバイトコードを使ってアプリケーションを実行します。
繰り返しになりますが、最初にLAUNCHが実行された時にはOpenCL JITコンパイラが起動されて、計算デバイスに従ったコードの特殊化と最適化を行います。そして最後に、デバイスからメインホストへのコピー(STREAM_OUT_BLOCKING)を実行して、結果を取得するのです。
生成されたOpenCLコードを分析する
TornadoVMが生成するOpenCLカーネルの中を調べてみましょう。TornadoVMでは、次に示すような--printKernelフラグを使用することで、生成されたカーネルのデバッグとチェックが可能です。
$ tornado --printKernel qconlondon.MatrixMultiplication 512 tornado
TornadoVMは、タスクスケジュール内のタスク毎にひとつのカーネルを生成します。さらに、VMのブートストラップ中に実行されるlookupBufferAddressというカーネルを生成します。このカーネルが必要な理由は、TornadoiVMがひとつの大きなバッファのみを確保して、ターゲットデバイス上のヒープとして使用していることから来ています。そのために、ターゲットデバイスからTornadoVMがデータ転送を実行する時にベースアドレスとして使用可能な、有効なポインタが必要なのです。lookupBufferAddressカーネルはこのベースポインタを返します。
第2のカーネルは、私たちが高速化したJavaメソッドから生成されたOpenCLコードに関わるものです。次のコードスニペットは、生成されたカーネルを簡略化したものを、Javaからの要点に関するコメントとOpenCLコードで示したものです。生成されるカーネルは、ターゲットのアーキテクチャによって異なる可能性がある点に注意してください。また、TornadoVMはOpenCL Cコードを静的単一代入(SSA、Static Single Assignment)表現で生成する、すなわち、各変数への割り当ては正確に1回である、という点にも注意が必要です。これはTornadoVMが、(HotSpotのJITコンパイラの主流であるC2と同じく)SSA表記を取り入れたGraal-IRの拡張である、という事情によるものです。
__kernel void lookupBufferAddress(...parameters) {
__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];
_frame[0] = (ulong) _heap_base;
}
__kernel void matrixMultiplication(...parameters) {
// Variables declaration …
// Access to the stack-frame
__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];
// Access elements within the stack-frame
ul_0 = (ulong) _frame[6]; // base address of input matrix A
ul_1 = (ulong) _frame[7]; // base address of input matrix B
ul_2 = (ulong) _frame[8]; // base address of input matrix C
i_3 = get_global_id(1); // Parallel OpenCL indexing (2nd dimension)
i_4 = i_3;
for(;i_4 < 512;) {
i_5 = get_global_id(0); // Parallel OpenCL indexing (1st dimension)
i_6 = i_5;
for(;i_6 < 512;) {
i_7 = i_4 << 9;
f_8 = 0.0F;
i_9 = 0;
for(;i_9 < 512;) {
i_10 = i_9 + 1;
i_11 = i_7 + i_9;
l_12 = (long) i_11;
l_13 = l_12 << 2;
l_14 = l_13 + 24L; // Skip Java object header
ul_15 = ul_0 + l_14;
f_16 = *((__global float *) ul_15); // Load element from matrix A
i_17 = i_9 << 9;
i_18 = i_17 + i_6;
l_19 = (long) i_18;
l_20 = l_19 << 2;
l_21 = l_20 + 24L;
ul_22 = ul_1 + l_21;
f_23 = *((__global float *) ul_22);// Load element from matrix B
f_24 = fma(f_16, f_23, f_8); // Computation (fuse-multiple-add)
f_8 = f_24;
i_9 = i_10;
}
i_25 = i_6 + i_7;
l_26 = (long) i_25;
l_27 = l_26 << 2;
l_28 = l_27 + 24L;
ul_29 = ul_2 + l_28;
*((__global float *) ul_29) = f_8; // Store the result in Matrix C
i_30 = get_global_size(0);
i_31 = i_30 + i_6;
i_6 = i_31;
}
i_32 = get_global_size(1);
i_33 = i_32 + i_4;
i_4 = i_33;
}
}
TornadoVmはどのように使用するのか？
この記事では、TornadoVMランタイムとJITコンパイラの違う部分を簡単に示すために、単純な例である行列乗算を中心として見てきました。しかしTonradoVMでは、単一タスクと単純なデータ型以上のプログラミングも可能です。Microsoft Kinect Fusionを使用したSLAM(Simultaneous Localization and Mapping)アプリケーションでの使用例では、NVIDIA GPU上において、Javaに比較して最大90フレーム/秒の高速化を実現しています。このアプリケーションでは約7,000行のJavaコードがTornadoVMによって高速化されているのですが、複雑なJava構造をTornadoVMが生成可能である点が注目されます。
一般論としてTornadoVMは、SIMD(Single Instruction Multiple Data)に従うワークロードや、パイプラインアプリケーションの高速化に適しています。意外なことに、このカテゴリには、ディープラーニングやマシンラーニング、数学や物理学のシミュレーション、計算写真学(computational photography)、コンピュータビジョン、金融アプリケーション、信号処理、化学など、さまざまなアプリケーションが含まれます。
TornadoVMはさらに、PythonやR、Ruby、JavaScript、その他のGraalVM上の言語から起動することができます(QCon-LondonではNode.jsアプリケーションの高速化を紹介しました)。
TornadoVMは学術界で生まれました(現在はマンチェスター大学で開発されています)が、すでに何社かの企業が、ディープラーニングアプリケーションを高速化するために使用しています。
そのひとつである、ロンドンを拠点とする企業のExus Ltd.では、患者の再入院を予測するUK NHS(医療)システムの改善に使用されていて、200万人の患者のデータセットによるトレーニングフェーズのパフォーマンスを、TornadoVMを使うことで14倍に改善することに成功しています。
産業界におけるアーリーアダプションのもうひとつの例は、ルクセンブルクのNEUROCOMです。同社では、TornadoVMをGPU上で使用して、自然言語処理で使用される重要な演算処理(具体的には、コサイン類似性メトリクスアルゴリズムを用いたレーベンシュタイン距離および階層分類子)を、それぞれ10倍および28倍高速化しています。
要約
TornadoVMは、マルチコアCPUやGPU、FPGAといった異種ハードウェアにJVMアプリケーションをオフラインする、OpenJDKとGraalVMのプラグインです。さらに、デバイス間のライブタスクマイグレーションを行うことで、アプリケーション全体のパフォーマンスを最大化します。今回の記事では、例を通じてTornadoVMの機能を探りました。TornadoVMがどのように動作するかを調査し、生成されるコードがどのようなものかを見てきました。
ですが、この記事ではTornadoVMとは何か、何ができるのかについて、表面的に論じたに過ぎません。今回の紹介記事では取り上げられなかった重要なトピックが、まだたくさんあります。例えば、アーキテクチャ毎のコンパイラの特殊化に関する説明、コンパイルの実行を効率的に削減する方法、FPGAコンパイルパイプライン、ライブタスクマイグレーションなどです。これらトピックの一部については、以下のリンクに詳細な説明があります。
参考文献
謝辞
TornadoVM開発は、欧州連合(European Union)のHorizon 2020 E2Data 780245のサポートを受けています。
著者について
Juan Fumero氏はマンチェスター大学の博士研究員として、異種高レベル言語仮想マシン、GPGPU、分散コンピューティングを研究しています。現在はTornadoVMおよびE2Data Europeanプロジェクトの一員として、GPUおよびFPGA用自動JITコンパイルとJavaプログラムでの実行処理を開発しています。氏はJava、R、Rubyといったインタプリタプログラミング言語の高速化によって、エジンバラ大学からPh.Dの学位を受けています。Oracle LabsやCERNにもインターンとして従事しており、コンパイラの実装やマルチコアシステム用の並列化テクニックの評価などを行っています。