NVPTX 後端使用者指南¶
簡介¶
為了支援 GPU 程式設計，NVPTX 後端支援 LLVM IR 的一部分，以及用於表示 GPU 程式設計概念的一組定義的慣例。本文檔概述了後端的常規用法，包括對所用慣例和可接受的 LLVM IR 集的描述。
注意
本文檔假設您對 CUDA 和 PTX 組合語言有基本的了解。有關 CUDA 驅動程式 API 和 PTX 組合語言的資訊，請參閱 CUDA 文件。
慣例¶
將函數標記為核心函數¶
在 PTX 中，有兩種函數：*裝置函數*，只能由裝置程式碼呼叫，以及*核心函數*，可以由主機程式碼呼叫。預設情況下，後端會發出裝置函數。中繼資料用於將函數宣告為核心函數。此中繼資料附加到 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 驅動程式 API 的 32 位元 PTX：nvptx-nvidia-cuda
範例：適用於 CUDA 驅動程式 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 指定的一組執行緒中選出一個斷言為 true 的領導執行緒。如果正在執行的執行緒不在 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.* 是一種單向柵欄，用於在透過泛型 proxy<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 中選擇兩個位元組，如果 %is.hi 為真，則選擇最高有效位元組，否則選擇最低有效位元組。然後將這些位元組擴展為 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 工具包隨附一個稱為 libdevice 的 LLVM 位元碼程式庫，其中實作了許多常用的數學函數。這個程式庫可以用於任何使用 LLVM NVPTX 目標的編譯器的效能數學程式庫。這個程式庫位於 CUDA 工具包中的 nvvm/libdevice/ 底下，並且每個計算架構都有一個獨立的版本。
如需 libdevice 中實作的所有數學函數的清單，請參閱 libdevice 使用者指南。
為了適應可能會影響 libdevice 程式碼產生之各種與數學相關的編譯器旗標，程式庫程式碼取決於特殊的 LLVM IR 階段 (NVVMReflect) 來處理 LLVM IR 中的條件式編譯。這個階段會搜尋對 @__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 引數來避免潛在不相容的指令。
儲存
module.bc中的外部函數清單將
module.bc與libdevice.compute_XX.YY.bc連結將 (1) 中清單以外的所有函數內部化
移除所有未使用的內部函數
執行
NVVMReflect階段執行標準最佳化流程
注意
鏈接類型 linkonce 和 linkonce_odr 不適用於 libdevice 函式。可以鏈接兩個已使用不同反射變量鏈接到 libdevice 的 IR 模組。
由於 NVVMReflect pass 會將條件式替換為常數，因此它通常會留下如下形式的死碼
entry:
..
br i1 true, label %foo, label %bar
foo:
..
bar:
; Dead code
..
因此，建議在死碼消除之前的優化流程早期執行 NVVMReflect。
NVPTX TargetMachine 知道如何在 pass 管理器的開頭安排 NVVMReflect；只需在設定 pass 管理器時使用以下代碼，PassBuilder 就會使用 registerPassBuilderCallbacks 讓 NVPTXTargetMachine::registerPassBuilderCallbacks 將 pass 添加到 pass 管理器中
std::unique_ptr<TargetMachine> TM = ...;
PassBuilder PB(TM);
ModulePassManager MPM;
PB.parsePassPipeline(MPM, ...);
反射參數¶
libdevice 庫目前使用以下反射參數來控制代碼生成
標誌 |
說明 |
|---|---|
|
使用將次正規數刷新為零的優化代碼路徑 |
此標誌的值由「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 驅動程序 API。此 API 是 GPU 驅動程序的低級介面，允許將 PTX 代碼 JIT 編譯為原生 GPU 機器代碼。
初始化驅動程序 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 pass。有關更多資訊，請參閱 與 Libdevice 鏈接。
教學課程：簡單的計算核心¶
首先，讓我們看一下直接用 LLVM IR 編寫的簡單計算核心。該核心實現向量加法，其中每個線程從輸入向量 A 和 B 計算輸出向量 C 的一個元素。為了簡化操作，我們還假設只會啟動一個 CTA（線程塊），並且它是一維的。
核心¶
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 等效項 |
|---|---|
|
threadIdx.{x,y,z} |
|
blockIdx.{x,y,z} |
|
blockDim.{x,y,z} |
|
gridDim.{x,y,z} |
|
__syncthreads() |
位址空間¶
您可能已經注意到，LLVM IR 範例中的所有指標類型都有一個明確的位址空間說明符。什麼是位址空間 1？NVIDIA GPU 裝置（通常）具有四種類型的記憶體
全域：大型、晶片外記憶體
共用：小型、晶片上記憶體，由 CTA 中的所有執行緒共用
區域：每個執行緒的私有記憶體
常數：所有執行緒共用的唯讀記憶體
這些不同類型的記憶體在 LLVM IR 中表示為位址空間。NVPTX 程式碼產生器還使用了第五個位址空間，它對應於「通用」位址空間。這個位址空間可以表示任何其他位址空間中的位址（除了一些例外）。這允許使用者編寫可以使用相同指令載入/儲存記憶體的 IR 函數。提供了在通用和非通用位址空間之間轉換指標的內建函數。
如需更多資訊，請參閱 位址空間 和 NVPTX 內建函數。
核心中繼資料¶
在 PTX 中，函數可以是 核心 函數（可從主機程式呼叫），也可以是 裝置 函數（只能從 GPU 程式碼呼叫）。您可以將 核心 函數視為 GPU 程式中的進入點。若要將 LLVM IR 函數標記為 核心 函數，我們使用特殊的 LLVM 中繼資料。NVPTX 後端將尋找一個名為 nvvm.annotations 的命名中繼資料節點。這個命名中繼資料必須包含描述 IR 的中繼資料清單。就我們的目的而言，我們需要宣告一個中繼資料節點，將「核心」屬性賦予應作為 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 驅動程式 API 提供了一種方便的機制，可以將 PTX 載入到原生 GPU 裝置並進行 JIT 編譯，以及啟動核心。該 API 類似於 OpenCL。以下顯示了一個簡單的範例，說明如何載入和執行我們的向量加法程式碼。請注意，為了簡潔起見，此程式碼沒有執行太多錯誤檢查！
注意
您也可以使用 CUDA 工具包提供的 ptxas 工具，將 PTX 離線編譯為特定 GPU 架構的機器碼 (SASS)。這些二進制檔案可以像 PTX 一樣由 CUDA 驅動程式 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 = pow(A, B) 而不是 C = 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}
要編譯此核心，我們執行以下步驟
與 libdevice 連結
將除公開核心函式外的所有函式內部化
執行
NVVMReflect並將__CUDA_FTZ設定為 0最佳化連結的模組
產生模組的程式碼
這些步驟可以透過 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;
}