NVPTXバックエンドのユーザーガイド

はじめに

GPUプログラミングをサポートするために、NVPTXバックエンドはLLVM IRのサブセットと、GPUプログラミングの概念を表すために使用される定義済みの規約セットをサポートしています。このドキュメントでは、使用される規約と受け入れられるLLVM IRのセットの説明を含め、バックエンドの一般的な使用方法の概要を説明します。

注記

このドキュメントでは、CUDAとPTXアセンブリ言語の基本的な知識があることを前提としています。CUDAドライバAPIとPTXアセンブリ言語に関する情報は、CUDAドキュメントにあります。

規約

関数をカーネルとしてマークする

PTXには、デバイスコードからのみ呼び出し可能な*デバイス関数*と、ホストコードから呼び出し可能な*カーネル関数*の2種類の関数があります。デフォルトでは、バックエンドはデバイス関数を生成します。関数をカーネル関数として宣言するには、メタデータを使用します。このメタデータは、nvvm.annotationsという名前のメタデータオブジェクトに添付され、次の形式です。

!0 = !{<function-ref>, metadata !"kernel", i32 1}

最初のパラメータは、カーネル関数への参照です。次の例は、LLVM IRでデバイス関数を呼び出すカーネル関数を示しています。@my_kernel関数はホストコードから呼び出すことができますが、@my_fmad関数は呼び出すことができません。

define float @my_fmad(float %x, float %y, float %z) {
  %mul = fmul float %x, %y
  %add = fadd float %mul, %z
  ret float %add
}

define void @my_kernel(ptr %ptr) {
  %val = load float, ptr %ptr
  %ret = call float @my_fmad(float %val, float %val, float %val)
  store float %ret, ptr %ptr
  ret void
}

!nvvm.annotations = !{!1}
!1 = !{ptr @my_kernel, !"kernel", i32 1}

コンパイルされると、PTXカーネル関数はホスト側のコードから呼び出すことができます。

アドレス空間

NVPTXバックエンドは、次のアドレス空間マッピングを使用します。

アドレス空間

メモリ空間

0

汎用

1

グローバル

2

内部使用

3

共有

4

定数

5

ローカル

すべてのグローバル変数とポインタ型は、これらのアドレス空間のいずれかに割り当てられ、0がデフォルトのアドレス空間です。汎用アドレス空間と非汎用アドレス空間の間でポインタを変換するために使用できる組み込み関数が用意されています。

例として、次のIRは、グローバルデバイスメモリに存在する配列@gを定義します。

@g = internal addrspace(1) global [4 x i32] [ i32 0, i32 1, i32 2, i32 3 ]

LLVM IR関数は、この配列を読み書きでき、ホスト側のコードは、CUDAドライバAPIを使用して名前でデータをコピーできます。

アドレス空間0は汎用空間であるため、アドレス空間0にグローバル変数を持つことは不正です。アドレス空間0はLLVMのデフォルトのアドレス空間であるため、グローバル変数にはaddrspace(N)アノテーションが*必須*です。

トリプル

NVPTXターゲットは、モジュールのトリプルを使用して、32/64ビットコード生成と使用するドライバコンパイラインターフェースを選択します。トリプルのアーキテクチャは、nvptx（32ビットPTX）またはnvptx64（64ビットPTX）のいずれかです。オペレーティングシステムは、cudaまたはnvclのいずれかである必要があり、これにより、生成されたコードがドライバと通信するために使用するインターフェースが決まります。ほとんどのユーザーは、オペレーティングシステムとしてcudaを使用することをお勧めします。これにより、生成されたPTXはCUDAドライバAPIと互換性があります。

例: CUDA Driver API 用の 32 ビット PTX: nvptx-nvidia-cuda

例: CUDA Driver API 用の 64 ビット PTX: nvptx64-nvidia-cuda

NVPTX組み込み関数

アドレス空間変換

‘llvm.nvvm.ptr.*.to.gen’ 組み込み関数

構文:

これらはオーバーロードされた組み込み関数です。これらは任意のポインタ型で使用できます。

declare ptr @llvm.nvvm.ptr.global.to.gen.p0.p1(ptr addrspace(1))
declare ptr @llvm.nvvm.ptr.shared.to.gen.p0.p3(ptr addrspace(3))
declare ptr @llvm.nvvm.ptr.constant.to.gen.p0.p4(ptr addrspace(4))
declare ptr @llvm.nvvm.ptr.local.to.gen.p0.p5(ptr addrspace(5))
概要:

‘llvm.nvvm.ptr.*.to.gen’組み込み関数は、非汎用アドレス空間のポインタを汎用アドレス空間ポインタに変換します。

セマンティクス:

これらの組み込み関数は、ポインタ値を有効な汎用アドレス空間ポインタになるように変更します。

‘llvm.nvvm.ptr.gen.to.*’ 組み込み関数

構文:

これらはオーバーロードされた組み込み関数です。これらは任意のポインタ型で使用できます。

declare ptr addrspace(1) @llvm.nvvm.ptr.gen.to.global.p1.p0(ptr)
declare ptr addrspace(3) @llvm.nvvm.ptr.gen.to.shared.p3.p0(ptr)
declare ptr addrspace(4) @llvm.nvvm.ptr.gen.to.constant.p4.p0(ptr)
declare ptr addrspace(5) @llvm.nvvm.ptr.gen.to.local.p5.p0(ptr)
概要:

‘llvm.nvvm.ptr.gen.to.*’組み込み関数は、汎用アドレス空間のポインタをターゲットアドレス空間のポインタに変換します。これらの組み込み関数は、ポインタのターゲットアドレス空間のアドレス空間がわかっている場合にのみ役立ちます。アドレス空間変換組み込み関数を使用して、ある非汎用アドレス空間から別の非汎用アドレス空間にポインタを変換することはできません。

セマンティクス:

これらの組み込み関数は、ポインタ値をターゲットの非汎用アドレス空間で有効なポインタになるように変更します。

PTX特殊レジスタの読み取り

‘llvm.nvvm.read.ptx.sreg.*’

構文:
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()
概要:

‘@llvm.nvvm.read.ptx.sreg.*’組み込み関数は、PTX特殊レジスタ、特にカーネル起動境界へのアクセスを提供します。これらのレジスタは、次の方法でCUDA組み込み関数にマップされます。

CUDA組み込み関数

PTX特殊レジスタ組み込み関数

threadId

@llvm.nvvm.read.ptx.sreg.tid.*

blockIdx

@llvm.nvvm.read.ptx.sreg.ctaid.*

blockDim

@llvm.nvvm.read.ptx.sreg.ntid.*

gridDim

@llvm.nvvm.read.ptx.sreg.nctaid.*

バリア

‘llvm.nvvm.barrier0’

構文:
declare void @llvm.nvvm.barrier0()
概要:

‘@llvm.nvvm.barrier0()’組み込み関数は、PTXのbar.sync 0命令を生成します。これは、CUDAの__syncthreads()呼び出しと同じです。

スレッドの選出

‘llvm.nvvm.elect.sync’

構文:
declare {i32, i1} @llvm.nvvm.elect.sync(i32 %membermask)
概要:

‘@llvm.nvvm.elect.sync’ 組み込み関数は、elect.sync PTX 命令を生成します。この命令は、membermask で指定されたスレッドの集合から、述語でアクティブなリーダー スレッドを 1 つ選択します。実行中のスレッドが membermask に含まれていない場合、動作は未定義です。選択されたスレッドの laneid は、i32 戻り値にキャプチャされます。i1 戻り値は、リーダー スレッドの場合は True に、他のすべてのスレッドの場合は False に設定されます。リーダー スレッドの選択は決定論的に行われます。つまり、同じ membermask に対しては、毎回同じリーダー スレッドが選択されます。詳細については、PTX ISA を参照してください https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync。

メモリバリア/フェンス

‘llvm.nvvm.fence.proxy.tensormap_generic.*’

構文:
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.cta()
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.cluster()
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.gpu()
declare void @llvm.nvvm.fence.proxy.tensormap_generic.release.sys()

declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr %addr, i32 %size)
declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cluster(ptr %addr, i32 %size)
declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.gpu(ptr %addr, i32 %size)
declare void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.sys(ptr %addr, i32 %size)
概要:

@llvm.nvvm.fence.proxy.tensormap_generic.* は、ジェネリック プロキシ<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#proxies>_ を介して実行された以前のメモリアクセスと、tensormap プロキシを介して実行された後続のメモリアクセスの間に順序を確立するために使用される単方向フェンスです。 nvvm.fence.proxy.tensormap_generic.release は、nvvm.fence.proxy.tensormap_generic.acquire プロキシフェンスを含む取得シーケンスと同期するリリースシーケンスを形成できます。次の表は、LLVM組み込み関数とPTX命令間のマッピングを示しています。

NVVM 組み込み関数

PTX 命令

@llvm.nvvm.fence.proxy.tensormap_generic.release.*

fence.proxy.tensormap::generic.release.*

@llvm.nvvm.fence.proxy.tensormap_generic.acquire.*

fence.proxy.tensormap::generic.acquire.* [addr], size

アドレスオペランド addr とオペランド size は、プロキシ間のメモリアクセスの順序保証を提供するメモリ範囲 [addr, addr+size) を指定します。 size オペランドでサポートされている値は 128 のみであり、即値である必要があります。ジェネリックアドレッシングは無条件に使用され、オペランドaddrで指定されたアドレスは .global 状態空間に含まれている必要があります。そうでない場合、動作は未定義です。詳細については、PTX ISA を参照してください。

算術組み込み関数

‘llvm.nvvm.idp2a.[us].[us]’ 組み込み関数

構文:
declare i32 @llvm.nvvm.idp2a.s.s(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
declare i32 @llvm.nvvm.idp2a.s.u(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
declare i32 @llvm.nvvm.idp2a.u.s(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
declare i32 @llvm.nvvm.idp2a.u.u(i32 %a, i32 %b, i1 immarg %is.hi, i32 %c)
概要:

‘llvm.nvvm.idp2a.[us].[us]’ 組み込み関数は、2 要素ベクトル ドット積を実行し、加算を行います。これらは、dp2a PTX 命令に直接対応します。

セマンティクス:

%a の 32 ビット値は、2 つの 16 ビット値に分割され、32 ビットに拡張されます。 ‘llvm.nvvm.idp2a.u.[us]’ バリアントではゼロ拡張が使用され、’llvm.nvvm.idp2a.s.[us]’ バリアントでは符号拡張が使用されます。 %b から 2 バイトが選択されます。 %is.hi が true の場合、最上位バイトが選択され、そうでない場合は最下位バイトが選択されます。これらのバイトは 32 ビットに拡張されます。 ‘llvm.nvvm.idp2a.[us].u’ バリアントではゼロ拡張が使用され、’llvm.nvvm.idp2a.[us].s’ バリアントでは符号拡張が使用されます。これらの 2 要素ベクトルのドット積が %c に加算され、戻り値が生成されます。

‘llvm.nvvm.idp4a.[us].[us]’ 組み込み関数

構文:
declare i32 @llvm.nvvm.idp4a.s.s(i32 %a, i32 %b, i32 %c)
declare i32 @llvm.nvvm.idp4a.s.u(i32 %a, i32 %b, i32 %c)
declare i32 @llvm.nvvm.idp4a.u.s(i32 %a, i32 %b, i32 %c)
declare i32 @llvm.nvvm.idp4a.u.u(i32 %a, i32 %b, i32 %c)
概要:

‘llvm.nvvm.idp4a.[us].[us]’ 組み込み関数は、4 要素ベクトル ドット積を実行し、加算を行います。これらは、dp4a PTX 命令に直接対応します。

セマンティクス:

%a と %b の 4 バイトはそれぞれ 32 ビット整数に拡張され、2 つの <4 x i32> が形成されます。 %a では、’llvm.nvvm.idp4a.u.[us]’ バリアントではゼロ拡張が使用され、’llvm.nvvm.idp4a.s.[us]’ バリアントでは符号拡張が使用されます。 同様に、%b では、’llvm.nvvm.idp4a.[us].u’ バリアントではゼロ拡張が使用され、’llvm.nvvm.idp4a.[us].s’ バリアントでは符号拡張が使用されます。これらの 4 要素ベクトルのドット積が %c に加算され、戻り値が生成されます。

その他の組み込み関数

NVPTX 組み込み関数の完全なセットについては、LLVM ソースツリーの include/llvm/IR/IntrinsicsNVVM.td ファイルを参照してください。

Libdevice とのリンク

CUDA Toolkit には、多くの一般的な数学関数を実装した libdevice と呼ばれる LLVM ビットコードライブラリが付属しています。このライブラリは、LLVM NVPTX ターゲットを使用するコンパイラ向けの高性能数学ライブラリとして使用できます。ライブラリは CUDA Toolkit の nvvm/libdevice/ にあり、計算アーキテクチャごとに個別のバージョンがあります。

libdevice に実装されているすべての数学関数のリストについては、libdevice ユーザーガイド を参照してください。

libdevice コードのコード生成に影響を与える可能性のある、数学関連のさまざまなコンパイラフラグに対応するために、ライブラリコードは、LLVM IR 内の条件付きコンパイルを処理するための特別な LLVM IR パス (NVVMReflect) に依存しています。このパスは、@__nvvm_reflect 関数への呼び出しを探し、定義されたリフレクションパラメータに基づいてそれらを定数に置き換えます。このような条件付きコードは、多くの場合、次のパターンに従います。

float my_function(float a) {
  if (__nvvm_reflect("FASTMATH"))
    return my_function_fast(a);
  else
    return my_function_precise(a);
}

指定されていないすべてのリフレクションパラメータのデフォルト値はゼロです。

NVVMReflect パスは、リンクステージの直後、最適化パイプラインの初期段階で実行する必要があります。 internalize パスは、結果の PTX から未使用の数学関数を削除するためにも推奨されます。入力 IR モジュール module.bc の場合、次のコンパイルフローが推奨されます。

NVVMReflect パスは、最適化を行わなくても、デッドコードの削除を試みます。これにより、__CUDA_ARCH 引数を使用することで、すべての最適化レベルで互換性のない命令を回避できます。

  1. module.bc 内の外部関数のリストを保存する

  2. module.bc を libdevice.compute_XX.YY.bc とリンクする

  3. (1) のリストにないすべての関数を内部化する

  4. 未使用のすべての内部関数を削除する

  5. NVVMReflect パスを実行する

  6. 標準の最適化パイプラインを実行する

注記

linkonce および linkonce_odr リンクタイプは、libdevice 関数には適していません。異なるリフレクション変数を使用して libdevice にリンクされている 2 つの IR モジュールをリンクすることができます。

NVVMReflect パスは条件を定数に置き換えるため、多くの場合、次のような形式のデッドコードが残ります。

entry:
  ..
  br i1 true, label %foo, label %bar
foo:
  ..
bar:
  ; Dead code
  ..

したがって、デッドコードの削除を行う前に、最適化パイプラインの初期段階で NVVMReflect を実行することをお勧めします。

NVPTX TargetMachine は、パスマネージャーの先頭に NVVMReflect をスケジュールする方法を認識しています。パスマネージャーを設定する際に以下のコードを使用するだけで、PassBuilder は registerPassBuilderCallbacks を使用して、NVPTXTargetMachine::registerPassBuilderCallbacks にパスをパスマネージャーに追加させます。

std::unique_ptr<TargetMachine> TM = ...;
PassBuilder PB(TM);
ModulePassManager MPM;
PB.parsePassPipeline(MPM, ...);

リフレクションパラメータ

libdeviceライブラリは現在、コード生成を制御するために以下のリフレクションパラメータを使用しています。

フラグ

説明

__CUDA_FTZ=[0,1]

非正規数をゼロにフラッシュする最適化されたコードパスを使用します。

このフラグの値は、「nvvm-reflect-ftz」モジュールフラグによって決定されます。以下は、ftzフラグを1に設定します。

!llvm.module.flags = !{!0}
!0 = !{i32 4, !"nvvm-reflect-ftz", i32 1}

(i32 4 は、ここで設定された値が、リンクする別のモジュールの値をオーバーライドすることを示します。詳細は、LangRef <LangRef.html#module-flags-metadata> を参照してください。)

PTXの実行

GPUデバイス上でPTXアセンブリを実行する最も一般的な方法は、CUDA Driver APIを使用することです。このAPIは、GPUドライバへの低レベルインターフェースであり、PTXコードをネイティブGPUマシンコードにJITコンパイルすることができます。

Driver APIの初期化

CUdevice device;
CUcontext context;

// Initialize the driver API
cuInit(0);
// Get a handle to the first compute device
cuDeviceGet(&device, 0);
// Create a compute device context
cuCtxCreate(&context, 0, device);

PTX文字列をデバイスバイナリにJITコンパイルする

CUmodule module;
CUfunction function;

// JIT compile a null-terminated PTX string
cuModuleLoadData(&module, (void*)PTXString);

// Get a handle to the "myfunction" kernel function
cuModuleGetFunction(&function, module, "myfunction");

PTXアセンブリの実行に関する完全な例については、CUDAサンプルディストリビューションを参照してください。

一般的な問題

ptxasが未定義の関数: __nvvm_reflectを訴える

libdeviceとリンクする場合、NVVMReflect パスを使用する必要があります。詳細は、Libdeviceとのリンク を参照してください。

チュートリアル: 簡単なコンピュートカーネル

まず、LLVM IRで直接記述された単純なコンピュートカーネルを見てみましょう。カーネルはベクトル加算を実装し、各スレッドは入力ベクトルAとBから出力ベクトルCの1つの要素を計算します。これを簡単にするために、単一のCTA（スレッドブロック）のみが起動され、それが1次元であると仮定します。

カーネル

target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

; Intrinsic to read X component of thread ID
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind

define void @kernel(ptr addrspace(1) %A,
                    ptr addrspace(1) %B,
                    ptr addrspace(1) %C) {
entry:
  ; What is my ID?
  %id = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind

  ; Compute pointers into A, B, and C
  %ptrA = getelementptr float, ptr addrspace(1) %A, i32 %id
  %ptrB = getelementptr float, ptr addrspace(1) %B, i32 %id
  %ptrC = getelementptr float, ptr addrspace(1) %C, i32 %id

  ; Read A, B
  %valA = load float, ptr addrspace(1) %ptrA, align 4
  %valB = load float, ptr addrspace(1) %ptrB, align 4

  ; Compute C = A + B
  %valC = fadd float %valA, %valB

  ; Store back to C
  store float %valC, ptr addrspace(1) %ptrC, align 4

  ret void
}

!nvvm.annotations = !{!0}
!0 = !{ptr @kernel, !"kernel", i32 1}

LLVMの llc ツールを使用して、NVPTXコードジェネレーターを直接実行できます。

# llc -mcpu=sm_20 kernel.ll -o kernel.ptx

注記

32ビットコードを生成する場合は、モジュールデータレイアウト文字列の p:64:64:64 を p:32:32:32 に変更し、ターゲットトリプルとして nvptx-nvidia-cuda を使用します。

llc から得られる出力 (LLVM 3.4現在)

//
// Generated by LLVM NVPTX Back-End
//

.version 3.1
.target sm_20
.address_size 64

  // .globl kernel
                                        // @kernel
.visible .entry kernel(
  .param .u64 kernel_param_0,
  .param .u64 kernel_param_1,
  .param .u64 kernel_param_2
)
{
  .reg .f32   %f<4>;
  .reg .s32   %r<2>;
  .reg .s64   %rl<8>;

// %bb.0:                                // %entry
  ld.param.u64    %rl1, [kernel_param_0];
  mov.u32         %r1, %tid.x;
  mul.wide.s32    %rl2, %r1, 4;
  add.s64         %rl3, %rl1, %rl2;
  ld.param.u64    %rl4, [kernel_param_1];
  add.s64         %rl5, %rl4, %rl2;
  ld.param.u64    %rl6, [kernel_param_2];
  add.s64         %rl7, %rl6, %rl2;
  ld.global.f32   %f1, [%rl3];
  ld.global.f32   %f2, [%rl5];
  add.f32         %f3, %f1, %f2;
  st.global.f32   [%rl7], %f3;
  ret;
}

カーネルの分析

それでは、このカーネルを構成するLLVM IRを分析してみましょう。

データレイアウト

データレイアウト文字列は、一般的なデータ型のビット単位のサイズ、ABIアライメント、およびストレージサイズを決定します。NVPTXの場合、以下のいずれかを使用する必要があります。

32ビットPTX

target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

64ビットPTX

target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

ターゲット組み込み関数

この例では、@llvm.nvvm.read.ptx.sreg.tid.x 組み込み関数を使用して、現在のスレッドのIDのX成分を読み取ります。これは、PTXのレジスタ %tid.x の読み取りに対応します。NVPTXバックエンドは、多数の組み込み関数をサポートしています。以下に簡単なリストを示します。完全なリストについては、include/llvm/IR/IntrinsicsNVVM.td を参照してください。

組み込み関数

CUDA相当

i32 @llvm.nvvm.read.ptx.sreg.tid.{x,y,z}

threadIdx.{x,y,z}

i32 @llvm.nvvm.read.ptx.sreg.ctaid.{x,y,z}

blockIdx.{x,y,z}

i32 @llvm.nvvm.read.ptx.sreg.ntid.{x,y,z}

blockDim.{x,y,z}

i32 @llvm.nvvm.read.ptx.sreg.nctaid.{x,y,z}

gridDim.{x,y,z}

void @llvm.nvvm.barrier0()

__syncthreads()

アドレス空間

LLVM IRの例のすべてのポインタ型に、明示的なアドレス空間指定子が付いていたことに気付いたかもしれません。アドレス空間1とは何でしょうか？NVIDIA GPUデバイスには、（一般的に）4種類のメモリがあります。

  • グローバル: 大容量のオフチップメモリ

  • 共有: CTA内のすべてのスレッド間で共有される、小容量のオンチップメモリ

  • ローカル: スレッドごとのプライベートメモリ

  • 定数: すべてスレッド間で共有される読み取り専用メモリ

これらの異なるタイプのメモリは、LLVM IRではアドレス空間として表されます。「汎用」アドレス空間に対応する、NVPTXコードジェネレーターで使用される5番目のアドレス空間もあります。このアドレス空間は、（いくつかの例外を除いて）他のアドレス空間のアドレスを表すことができます。これにより、ユーザーは同じ命令を使用してメモリを読み込み/書き込みできるIR関数を作成できます。汎用アドレス空間と非汎用アドレス空間の間でポインタを変換するための組み込み関数が用意されています。

詳細は、アドレス空間 および NVPTX組み込み関数 を参照してください。

カーネルメタデータ

PTXでは、関数は カーネル 関数（ホストプログラムから呼び出し可能）または デバイス 関数（GPUコードからのみ呼び出し可能）のいずれかになります。カーネル 関数は、GPUプログラムのエントリポイントと考えることができます。LLVM IR関数を カーネル 関数としてマークするために、特別なLLVMメタデータを使用します。NVPTXバックエンドは、nvvm.annotations と呼ばれる名前付きメタデータノードを探します。この名前付きメタデータには、IRを記述するメタデータのリストが含まれている必要があります。ここでの目的のために、「kernel」属性を、PTX カーネル 関数として出力されるべきLLVM IR関数に割り当てるメタデータノードを宣言する必要があります。これらのメタデータノードは、次の形式を取ります。

!{<function ref>, metadata !"kernel", i32 1}

前の例では、次のようになります。

!nvvm.annotations = !{!0}
!0 = !{ptr @kernel, !"kernel", i32 1}

ここでは、nvvm.annotations に単一のメタデータ宣言があります。このメタデータは、@kernel 関数に kernel 属性でアノテーションを付けます。

カーネルの実行

LLVM IRからPTXを生成することは結構なことですが、実際のGPUデバイスでどのように実行するのでしょうか？CUDA Driver APIは、PTXをネイティブGPUデバイスにロードおよびJITコンパイルし、カーネルを起動するための便利なメカニズムを提供します。APIはOpenCLに似ています。ベクトル加算コードをロードして実行する方法を示す簡単な例を以下に示します。簡潔にするために、このコードはあまりエラーチェックを行いません！

注記

CUDA Toolkitによって提供される ptxas ツールを使用して、PTXを特定のGPUアーキテクチャのマシンコード（SASS）にオフラインコンパイルすることもできます。このようなバイナリは、PTXと同じ方法でCUDA Driver APIによってロードできます。これは、PTXカーネルをプリコンパイルすることで起動時間を短縮するのに役立ちます。

#include <iostream>
#include <fstream>
#include <cassert>
#include "cuda.h"


void checkCudaErrors(CUresult err) {
  assert(err == CUDA_SUCCESS);
}

/// main - Program entry point
int main(int argc, char **argv) {
  CUdevice    device;
  CUmodule    cudaModule;
  CUcontext   context;
  CUfunction  function;
  CUlinkState linker;
  int         devCount;

  // CUDA initialization
  checkCudaErrors(cuInit(0));
  checkCudaErrors(cuDeviceGetCount(&devCount));
  checkCudaErrors(cuDeviceGet(&device, 0));

  char name[128];
  checkCudaErrors(cuDeviceGetName(name, 128, device));
  std::cout << "Using CUDA Device [0]: " << name << "\n";

  int devMajor, devMinor;
  checkCudaErrors(cuDeviceComputeCapability(&devMajor, &devMinor, device));
  std::cout << "Device Compute Capability: "
            << devMajor << "." << devMinor << "\n";
  if (devMajor < 2) {
    std::cerr << "ERROR: Device 0 is not SM 2.0 or greater\n";
    return 1;
  }

  std::ifstream t("kernel.ptx");
  if (!t.is_open()) {
    std::cerr << "kernel.ptx not found\n";
    return 1;
  }
  std::string str((std::istreambuf_iterator<char>(t)),
                    std::istreambuf_iterator<char>());

  // Create driver context
  checkCudaErrors(cuCtxCreate(&context, 0, device));

  // Create module for object
  checkCudaErrors(cuModuleLoadDataEx(&cudaModule, str.c_str(), 0, 0, 0));

  // Get kernel function
  checkCudaErrors(cuModuleGetFunction(&function, cudaModule, "kernel"));

  // Device data
  CUdeviceptr devBufferA;
  CUdeviceptr devBufferB;
  CUdeviceptr devBufferC;

  checkCudaErrors(cuMemAlloc(&devBufferA, sizeof(float)*16));
  checkCudaErrors(cuMemAlloc(&devBufferB, sizeof(float)*16));
  checkCudaErrors(cuMemAlloc(&devBufferC, sizeof(float)*16));

  float* hostA = new float[16];
  float* hostB = new float[16];
  float* hostC = new float[16];

  // Populate input
  for (unsigned i = 0; i != 16; ++i) {
    hostA[i] = (float)i;
    hostB[i] = (float)(2*i);
    hostC[i] = 0.0f;
  }

  checkCudaErrors(cuMemcpyHtoD(devBufferA, &hostA[0], sizeof(float)*16));
  checkCudaErrors(cuMemcpyHtoD(devBufferB, &hostB[0], sizeof(float)*16));


  unsigned blockSizeX = 16;
  unsigned blockSizeY = 1;
  unsigned blockSizeZ = 1;
  unsigned gridSizeX  = 1;
  unsigned gridSizeY  = 1;
  unsigned gridSizeZ  = 1;

  // Kernel parameters
  void *KernelParams[] = { &devBufferA, &devBufferB, &devBufferC };

  std::cout << "Launching kernel\n";

  // Kernel launch
  checkCudaErrors(cuLaunchKernel(function, gridSizeX, gridSizeY, gridSizeZ,
                                 blockSizeX, blockSizeY, blockSizeZ,
                                 0, NULL, KernelParams, NULL));

  // Retrieve device data
  checkCudaErrors(cuMemcpyDtoH(&hostC[0], devBufferC, sizeof(float)*16));


  std::cout << "Results:\n";
  for (unsigned i = 0; i != 16; ++i) {
    std::cout << hostA[i] << " + " << hostB[i] << " = " << hostC[i] << "\n";
  }


  // Clean up after ourselves
  delete [] hostA;
  delete [] hostB;
  delete [] hostC;

  // Clean-up
  checkCudaErrors(cuMemFree(devBufferA));
  checkCudaErrors(cuMemFree(devBufferB));
  checkCudaErrors(cuMemFree(devBufferC));
  checkCudaErrors(cuModuleUnload(cudaModule));
  checkCudaErrors(cuCtxDestroy(context));

  return 0;
}

CUDAドライバとリンクし、cuda.hへのパスを指定する必要があります。

# clang++ sample.cpp -o sample -O2 -g -I/usr/local/cuda-5.5/include -lcuda

libcuda.so へのパスを指定する必要はありません。これは、CUDAツールキットではなく、ドライバによってシステムの場所にインストールされるためです。

すべてが計画通りに進めば、コンパイルされたプログラムを実行すると、次の出力が表示されます。

Using CUDA Device [0]: GeForce GTX 680
Device Compute Capability: 3.0
Launching kernel
Results:
0 + 0 = 0
1 + 2 = 3
2 + 4 = 6
3 + 6 = 9
4 + 8 = 12
5 + 10 = 15
6 + 12 = 18
7 + 14 = 21
8 + 16 = 24
9 + 18 = 27
10 + 20 = 30
11 + 22 = 33
12 + 24 = 36
13 + 26 = 39
14 + 28 = 42
15 + 30 = 45

注記

ハードウェアに基づいて、異なるデバイス識別子が表示される可能性があります。

チュートリアル: Libdeviceとのリンク

このチュートリアルでは、LLVM IRをlibdeviceライブラリとリンクする簡単な例を示します。前のチュートリアルと同じカーネルを使用しますが、C = A + B ではなく C = pow(A, B) を計算します。Libdeviceは、使用する __nv_powf 関数を提供します。

target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

; Intrinsic to read X component of thread ID
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind
; libdevice function
declare float @__nv_powf(float, float)

define void @kernel(ptr addrspace(1) %A,
                    ptr addrspace(1) %B,
                    ptr addrspace(1) %C) {
entry:
  ; What is my ID?
  %id = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind

  ; Compute pointers into A, B, and C
  %ptrA = getelementptr float, ptr addrspace(1) %A, i32 %id
  %ptrB = getelementptr float, ptr addrspace(1) %B, i32 %id
  %ptrC = getelementptr float, ptr addrspace(1) %C, i32 %id

  ; Read A, B
  %valA = load float, ptr addrspace(1) %ptrA, align 4
  %valB = load float, ptr addrspace(1) %ptrB, align 4

  ; Compute C = pow(A, B)
  %valC = call float @__nv_powf(float %valA, float %valB)

  ; Store back to C
  store float %valC, ptr addrspace(1) %ptrC, align 4

  ret void
}

!nvvm.annotations = !{!0}
!0 = !{ptr @kernel, !"kernel", i32 1}

このカーネルをコンパイルするには、次の手順を実行します。

  1. libdeviceとリンクする

  2. パブリックカーネル関数以外のすべてを内部化する

  3. NVVMReflect を実行し、__CUDA_FTZ を0に設定する

  4. リンクされたモジュールを最適化する

  5. モジュールのコード生成を行う

これらの手順は、LLVMの llvm-link、opt、および llc ツールによって実行できます。完全なコンパイラでは、適切なパスの構成を設定することにより、これらの手順を完全にプログラムで実行することもできます（Libdeviceとのリンク を参照）。

# llvm-link t2.bc libdevice.compute_20.10.bc -o t2.linked.bc
# opt -internalize -internalize-public-api-list=kernel -nvvm-reflect-list=__CUDA_FTZ=0 -nvvm-reflect -O3 t2.linked.bc -o t2.opt.bc
# llc -mcpu=sm_20 t2.opt.bc -o t2.ptx

注記

-nvvm-reflect-list=_CUDA_FTZ=0 は、未定義の変数はデフォルトでゼロになるため、厳密には必要ありません。ここでは、評価の目的で示されています。

これにより、次のPTXが得られます（抜粋）。

//
// Generated by LLVM NVPTX Back-End
//

.version 3.1
.target sm_20
.address_size 64

  // .globl kernel
                                        // @kernel
.visible .entry kernel(
  .param .u64 kernel_param_0,
  .param .u64 kernel_param_1,
  .param .u64 kernel_param_2
)
{
  .reg .pred  %p<30>;
  .reg .f32   %f<111>;
  .reg .s32   %r<21>;
  .reg .s64   %rl<8>;

// %bb.0:                                // %entry
  ld.param.u64  %rl2, [kernel_param_0];
  mov.u32   %r3, %tid.x;
  ld.param.u64  %rl3, [kernel_param_1];
  mul.wide.s32  %rl4, %r3, 4;
  add.s64   %rl5, %rl2, %rl4;
  ld.param.u64  %rl6, [kernel_param_2];
  add.s64   %rl7, %rl3, %rl4;
  add.s64   %rl1, %rl6, %rl4;
  ld.global.f32   %f1, [%rl5];
  ld.global.f32   %f2, [%rl7];
  setp.eq.f32 %p1, %f1, 0f3F800000;
  setp.eq.f32 %p2, %f2, 0f00000000;
  or.pred   %p3, %p1, %p2;
  @%p3 bra  BB0_1;
  bra.uni   BB0_2;
BB0_1:
  mov.f32   %f110, 0f3F800000;
  st.global.f32   [%rl1], %f110;
  ret;
BB0_2:                                  // %__nv_isnanf.exit.i
  abs.f32   %f4, %f1;
  setp.gtu.f32  %p4, %f4, 0f7F800000;
  @%p4 bra  BB0_4;
// %bb.3:                                // %__nv_isnanf.exit5.i
  abs.f32   %f5, %f2;
  setp.le.f32 %p5, %f5, 0f7F800000;
  @%p5 bra  BB0_5;
BB0_4:                                  // %.critedge1.i
  add.f32   %f110, %f1, %f2;
  st.global.f32   [%rl1], %f110;
  ret;
BB0_5:                                  // %__nv_isinff.exit.i

  ...

BB0_26:                                 // %__nv_truncf.exit.i.i.i.i.i
  mul.f32   %f90, %f107, 0f3FB8AA3B;
  cvt.rzi.f32.f32 %f91, %f90;
  mov.f32   %f92, 0fBF317200;
  fma.rn.f32  %f93, %f91, %f92, %f107;
  mov.f32   %f94, 0fB5BFBE8E;
  fma.rn.f32  %f95, %f91, %f94, %f93;
  mul.f32   %f89, %f95, 0f3FB8AA3B;
  // inline asm
  ex2.approx.ftz.f32 %f88,%f89;
  // inline asm
  add.f32   %f96, %f91, 0f00000000;
  ex2.approx.f32  %f97, %f96;
  mul.f32   %f98, %f88, %f97;
  setp.lt.f32 %p15, %f107, 0fC2D20000;
  selp.f32  %f99, 0f00000000, %f98, %p15;
  setp.gt.f32 %p16, %f107, 0f42D20000;
  selp.f32  %f110, 0f7F800000, %f99, %p16;
  setp.eq.f32 %p17, %f110, 0f7F800000;
  @%p17 bra   BB0_28;
// %bb.27:
  fma.rn.f32  %f110, %f110, %f108, %f110;
BB0_28:                                 // %__internal_accurate_powf.exit.i
  setp.lt.f32 %p18, %f1, 0f00000000;
  setp.eq.f32 %p19, %f3, 0f3F800000;
  and.pred    %p20, %p18, %p19;
  @!%p20 bra  BB0_30;
  bra.uni   BB0_29;
BB0_29:
  mov.b32    %r9, %f110;
  xor.b32   %r10, %r9, -2147483648;
  mov.b32    %f110, %r10;
BB0_30:                                 // %__nv_powf.exit
  st.global.f32   [%rl1], %f110;
  ret;
}