この記事はGPGPU Advent Calendarの8日目の記事です。
LLVM meets GPU again!
CUDA4.1以降のnvccは、Compute Capability 2.0以上のコードを生成する際にLLVM IRのサブセットであるNVVMを経由して最適化等を行ってから、NVIDIAの規定する中間表現であるPTXへの変換を行なっています。今年に入ってこの成果がLLVM本家にマージされ、バージョン3.2で正式にお目見えすることになりました。LLVM 3.2の正式リリースは2012/12/16と一週間ほど先に予定されていますが、リポジトリにはすでに3.2用のブランチが切られています。今回は一足先にLLVM 3.2を使用して、LLVM IRからPTXを生成してみましょう!
歴史的な経緯
今回LLVMにマージされたNVIDIAの実装したPTXバックエンドはNVPTXと呼ばれています。実は、バージョン3.2以前のLLVMにもPTXバックエンドは存在していました。しかし機能が十分ではなかったり、最新のPTXに追随していなかったという問題があり、NVPTXにリプレースされる形でリポジトリからは姿を消しました。以降、単にPTXバックエンドと呼ぶ場合にはNVPTXを指すこととします。
また今年からNVIDIAはCUDA Compiler SDKというライブラリを一部の開発者向けに公開を開始していました。CUDA Compiler SDKも実質的にはNVVMからPTXを生成するためのライブラリで、LLVM 3.2の機能をより使いやすくラップしたライブラリという位置づけのようです。
準備編
まずはLLVM 3.2ブランチをgitリポジトリから取得しましょう。
% git clone http://llvm.org/git/llvm.git // ビールでも飲んで待つ
% cd llvm
% git checkout -b release_32 remotes/origin/release_32
ビルドはcmakeを使うことにします。無味乾燥なconfigure/makeなんかと違って、コンソール出力が赤や緑の極彩色に色どられて目にやさしい。それにもうすぐクリスマスですからね!
% mkdir build; cd build // ビルド用ディレクトリを掘る
% cmake -D CMAKE_INSTALL_PREFIX=/usr/local/llvm-3.2 ..
% make
% sudo make install
外部ライブラリなどが足りない場合はその旨が表示されるので、適当にパッケージシステムなどを利用して入れてください。インストールが終わったらインストール先のbinディレクトリ(上の例では/usr/local/llvm-3.2/bin)にパスを通して準備完了です。
使ってみる
さあ、あとはあれを良い感じにこうしてPTXが出ればもうゴールはすぐそこ!CUDAのPTXインタフェースを使えばJITコンパイルだって簡単だぞ。えーと、で何を入れればいいんだっけ?
テストを動かす
ソースツリーを眺めてみると、NVPTXのCodeGenテストが幾つかありますね。とりあえずこいつでも動かしてお茶を濁しましょう。
; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
;; These tests should run for all targets
;;===-- Basic instruction selection tests ---------------------------------===;;
;;; f64
define double @fadd_f64(double %a, double %b) {
; CHECK: add.f64 %fl{{[0-9]+}}, %fl{{[0-9]+}}, %fl{{[0-9]+}}
; CHECK: ret
%ret = fadd double %a, %b
ret double %ret
}
ふむ。LLVM IRに慣れていない人でも、2つのdouble型の変数を足し合わせる関数fadd_f64を定義していることがなんとなくわかると思います。冒頭のコメントにある通り、llcでコンパイルしてみましょう。
% llc < arithmetic-fp-sm20.ll -march=nvptx -mcpu=sm_20
//
// Generated by LLVM NVPTX Back-End
//
.version 3.1
.target sm_20, texmode_independent
.address_size 32
// .globl fadd_f64
.func (.param .b64 func_retval0) fadd_f64(
.param .b64 fadd_f64_param_0,
.param .b64 fadd_f64_param_1
) // @fadd_f64
{
.reg .pred %p<396>;
.reg .s16 %rc<396>;
.reg .s16 %rs<396>;
.reg .s32 %r<396>;
.reg .s64 %rl<396>;
.reg .f32 %f<396>;
.reg .f64 %fl<396>;
// BB#0:
ld.param.f64 %fl0, [fadd_f64_param_1];
ld.param.f64 %fl1, [fadd_f64_param_0];
add.f64 %fl0, %fl1, %fl0;
st.param.f64 [func_retval0+0], %fl0;
ret;
}
...
なんかそれっぽいPTXが出ましたね。
さて、上の例はデバイス関数だけしか含まれていないようですが、カーネル関数を含むような場合はどう書いたらいいんでしょうか?この場合のテストがtest/CodeGen/NVPTX/calling-conv.llにあります。
; RUN: llc < %s -march=nvptx -mcpu=sm_10 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_10 | FileCheck %s
; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
;; Kernel function using ptx_kernel calling conv
; CHECK: .entry kernel_func
define ptx_kernel void @kernel_func(float* %a) {
; CHECK: ret
ret void
}
;; Device function
; CHECK: .func device_func
define void @device_func(float* %a) {
; CHECK: ret
ret void
}
;; Kernel function using NVVM metadata
; CHECK: .entry metadata_kernel
define void @metadata_kernel(float* %a) {
; CHECK: ret
ret void
}
!nvvm.annotations = !{!1}
!1 = metadata !{void (float*)* @metadata_kernel, metadata !"kernel", i32 1}
はて、metadataという新しい表現が出てきました。これはいったい何でしょうか?
NVVMの仕様に関するドキュメントがない、ように見える
LLVM IRは立場上ポータブルな中間言語という位置づけですが、バックエンドに対して完全独立というわけではありません。例えば、CUDAでお馴染みのShared MemoryやGlobal Memoryといったメモリ空間はaddrspaceというLLVM IRの枠組みの中で表されるのですが、それぞれのメモリ空間がどういったフォーマットで表現されるべきなのかはバックエンドに依存します。そこで、NVPTXが解釈可能な取り決めがCUDA Compiler SDKの中ではNVVMとして規定されている・・・のですが、今見れる範囲のLLVMのドキュメント中にはどうも見当たらない。そのうち公開されると思いますが、ここではNVPTXのソースコードとテストから大事そうな部分を抜き出してみます。いっしゅんnvcc(正確にはnvvm/cicc)が内部的に保持するであろうLLVM IRをぶっこ抜いたろうかと思いましたが、いい大人はそういうことはしてはいけません。Binary Hacks is watching you.
関数の呼び出し規約
metadataとして関数に情報を付加することで、ホストから呼び出し可能なカーネル関数であることを明示します。metadataには関数の型と関数名を与え(void (float*)* @metadata_kernel)、カーネル関数であることを明示し(metadata !"kernel", i32 1)、最終的にnvvm.annotationsという名前付きmetadataの子要素にします。
!nvvm.annotations = !{!1}
!1 = metadata !{void (float*)* @metadata_kernel, metadata !"kernel", i32 1}
metadataをはじめとするLLVM IRのシンタックスについては、LLVMのリファレンス・マニュアルが参考になります。LLVM Language Reference Manual
メモリ空間
CUDAには、
- Global
- Shared
- Local
- Constant
このようなメモリ空間があります。これらを区別するために、NVVMではaddrspaceという仕組みをつかってポインタに情報を付加します。
define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) {
; PTX32: st.global.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}}
; PTX32: ret
; PTX64: st.global.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}}
; PTX64: ret
store i8 %a, i8 addrspace(1)* %ptr
ret void
}
変数%aをポインタ%ptrに格納する関数st_global_i8を定義しています。%ptrの型にaddrspace(1)という修飾がありますね。これが、%ptrの指す先がGlobal Memoryであることを示しています。
llvm/lib/Target/NVPTX/NVPTX.hを見ると、addrspaceの一覧がenumで定義されています。
enum AddressSpace{
GENERIC = 0,
GLOBAL = 1,
CONSTANT = 2,
SHARED = 3,
PARAM = 4,
LOCAL = 5
};
sm_20以降ではメモリ空間が統一されて、どのメモリに対してもGENERICポインタ、つまりaddrspace(0)なポインタを使用してメモリを読み書きできます。しかし、特定のメモリ空間に属する領域を宣言する時には、やはりaddrspaceを指定する必要があります。たとえば、Global Memoryに載るグローバル変数を定義する場合、以下のようになります。
@a = addrspace(1) global i8 2
組込み関数
CUDAには数学関数や同期命令などの組み込み関数がありますね。これらの一覧は以下のようにして見ることができます。
% cd llvm/lib/Target/NVPTX
% llvm-tblgen -I `llvm-config --includedir` -gen-intrinsic NVPTX.td | grep nvvm | less
さて、これで何をする?
ここでうっかりペリエを絨毯にぶちまいて掃除をしていたら力尽きてしまいました。というわけで次回に続く。