NVPTX 後端使用者指南

簡介

為了支援 GPU 程式設計，NVPTX 後端支援 LLVM IR 的子集，以及用於表示 GPU 程式設計概念的已定義慣例集。本文檔概述了後端的一般用法，包括對所使用慣例和接受的 LLVM IR 集的描述。

注意

本文檔假設讀者已基本熟悉 CUDA 和 PTX 組譯語言。有關 CUDA Driver API 和 PTX 組譯語言的資訊，請參閱 CUDA 文件。

慣例

將函數標記為核心函數

在 PTX 中，有兩種函數類型：裝置函數，僅可由裝置程式碼呼叫；以及核心函數，可由主機程式碼呼叫。預設情況下，後端將發射裝置函數。ptx_kernel 呼叫慣例用於將函數宣告為核心函數。

以下範例顯示了 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 ptx_kernel 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
}

編譯後，PTX 核心函數可由主機端程式碼呼叫。

函數屬性

"nvvm.maxclusterrank"="<n>"

此屬性指定每個叢集的最大區塊數。必須為非零值。僅適用於 Hopper+。

"nvvm.minctasm"="<n>"

這表示對編譯器/驅動程式的提示/指令，要求將至少這麼多 CTA 放在一個 SM 上。

"nvvm.maxnreg"="<n>"

此屬性表示核心函數要使用的最大暫存器數量。

"nvvm.maxntid"="<x>[,<y>[,<z>]]"

此屬性宣告執行緒區塊 (CTA) 中的最大執行緒數。最大執行緒數是每個維度中最大範圍的乘積。超過最大執行緒數會導致執行階段錯誤或核心啟動失敗。

"nvvm.reqntid"="<x>[,<y>[,<z>]]"

此屬性宣告執行緒區塊 (CTA) 中的確切執行緒數。執行緒數是每個維度值的乘積。在啟動時指定不同的 CTA 維度將導致執行階段錯誤或核心啟動失敗。

"nvvm.cluster_dim"="<x>[,<y>[,<z>]]"

此屬性宣告叢集中的執行緒區塊 (CTA) 數。CTA 總數是每個維度中 CTA 數量的乘積。在啟動時指定不同的叢集維度將導致執行階段錯誤或核心啟動失敗。僅適用於 Hopper+。

位址空間

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 Driver API 按名稱將資料複製到其中。

請注意，由於位址空間 0 是通用空間，因此在位址空間 0 中擁有全域變數是非法的。位址空間 0 是 LLVM 中的預設位址空間，因此全域變數必須使用 addrspace(N) 註解。

三元組

NVPTX 目標使用模組三元組來選擇 32/64 位元程式碼產生以及要使用的驅動程式-編譯器介面。三元組架構可以是 nvptx (32 位元 PTX) 或 nvptx64 (64 位元 PTX) 之一。作業系統應為 cuda 或 nvcl 之一，這決定了產生的程式碼用於與驅動程式通訊的介面。大多數使用者會想要使用 cuda 作為作業系統，這使得產生的 PTX 與 CUDA Driver API 相容。

範例：適用於 CUDA Driver API 的 32 位元 PTX：nvptx-nvidia-cuda

範例：適用於 CUDA Driver API 的 64 位元 PTX：nvptx64-nvidia-cuda

NVPTX 內建函數

讀取 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 指定的一組執行緒中選取一個預測的活動領導執行緒。如果執行中的執行緒不在 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 proxy 執行的記憶體存取之間建立順序。nvvm.fence.proxy.tensormap_generic.release 可以形成一個發布序列，該序列與包含 nvvm.fence.proxy.tensormap_generic.acquire proxy 柵欄的獲取序列同步。下表描述了 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)，在該範圍上提供跨 proxy 的記憶體存取順序保證。size 運算元唯一支援的值是 128，並且必須是立即值。通用定址無條件使用，且運算元 addr 指定的位址必須落在 .global 狀態空間內。否則，行為未定義。如需更多資訊，請參閱 PTX ISA。

位址空間內建函數

‘llvm.nvvm.isspacep.*’ 內建函數

語法：

declare i1 @llvm.nvvm.isspacep.const(ptr %p)
declare i1 @llvm.nvvm.isspacep.global(ptr %p)
declare i1 @llvm.nvvm.isspacep.local(ptr %p)
declare i1 @llvm.nvvm.isspacep.shared(ptr %p)
declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr %p)

概述：

‘llvm.nvvm.isspacep.*’ 內建函數確定提供的通用指標是否引用落在特定位址空間內的記憶體。

語意：

如果通用位址空間中給定的指標引用落在內建函數的狀態空間內的記憶體（因此可以安全地位址空間轉換到此空間），則傳回 1，否則傳回 0。

算術內建函數

‘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 為 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 以產生傳回值。

位元操作內建函數

‘llvm.nvvm.fshl.clamp.*’ 內建函數

語法：

declare i32 @llvm.nvvm.fshl.clamp.i32(i32 %hi, i32 %lo, i32 %n)

概述：

‘llvm.nvvm.fshl.clamp’ 系列內建函數執行鉗位式漏斗左移。這些內建函數與 ‘llvm.fshl’ 非常相似，只是移位量鉗位在整數寬度（而不是模數）。目前僅支援 i32。

語意：

‘llvm.nvvm.fshl.clamp’ 系列內建函數執行鉗位式漏斗左移：前兩個值串連為 { %hi : %lo }（%hi 是寬值的最高有效位元），組合值向左移位，並提取最高有效位元以產生與原始引數大小相同的結果。移位量是 %n 值與整數類型位寬的最小值。

‘llvm.nvvm.fshr.clamp.*’ 內建函數

語法：

declare i32 @llvm.nvvm.fshr.clamp.i32(i32 %hi, i32 %lo, i32 %n)

概述：

‘llvm.nvvm.fshr.clamp’ 系列內建函數執行鉗位式漏斗右移。這些內建函數與 ‘llvm.fshr’ 非常相似，只是移位量鉗位在整數寬度（而不是模數）。目前僅支援 i32。

語意：

‘llvm.nvvm.fshr.clamp’ 系列內建函數執行鉗位式漏斗右移：前兩個值串連為 { %hi : %lo }（%hi 是寬值的最高有效位元），組合值向右移位，並提取最低有效位元以產生與原始引數大小相同的結果。移位量是 %n 值與整數類型位寬的最小值。

‘llvm.nvvm.flo.u.*’ 內建函數

語法：

declare i32 @llvm.nvvm.flo.u.i32(i32 %a, i1 %shiftamt)
declare i32 @llvm.nvvm.flo.u.i64(i64 %a, i1 %shiftamt)

概述：

‘llvm.nvvm.flo.u’ 系列內建函數識別前導一的位元位置，傳回其與最高有效位元或最低有效位元的偏移量。

語意：

‘llvm.nvvm.flo.u’ 系列內建函數傳回最高有效位元 1 的位元位置。如果 %shiftamt 為 true，則結果是將找到的位元左移到最高有效位元位置所需的移位量，否則結果是將找到的位元右移到最低有效位元位置所需的移位量。如果未找到 1 位元，則傳回 0xffffffff。

‘llvm.nvvm.flo.s.*’ 內建函數

語法：

declare i32 @llvm.nvvm.flo.s.i32(i32 %a, i1 %shiftamt)
declare i32 @llvm.nvvm.flo.s.i64(i64 %a, i1 %shiftamt)

概述：

‘llvm.nvvm.flo.s’ 系列內建函數識別前導非符號位元的位元位置，傳回其與最高有效位元或最低有效位元的偏移量。

語意：

‘llvm.nvvm.flo.s’ 系列內建函數傳回負輸入的最高有效位元 0 和非負輸入的最高有效位元 1 的位元位置。如果 %shiftamt 為 true，則結果是將找到的位元左移到最高有效位元位置所需的移位量，否則結果是將找到的位元右移到最低有效位元位置所需的移位量。如果未找到 1 位元，則傳回 0xffffffff。

TMA 系列內建函數

‘llvm.nvvm.cp.async.bulk.global.to.shared.cluster’

語法：

declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3) %dst, ptr addrspace(3) %mbar, ptr addrspace(1) %src, i32 %size, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch)

概述：

‘@llvm.nvvm.cp.async.bulk.global.to.shared.cluster’ 內建函數對應於 cp.async.bulk.shared::cluster.global.* 系列 PTX 指令。這些指令啟動從全域記憶體到 shared::cluster 記憶體的批量資料非同步複製。32 位元運算元 %size 指定要複製的記憶體量，並且必須是 16 的倍數。

• 這些內建函數的最後兩個引數是布林旗標，指示對 cache_hint 和/或 multicast 修飾符的支援。這些旗標引數必須是編譯時常數。後端會檢查這些旗標並適當地降低內建函數。

• 第 N 個引數（以 i1 %flag_ch 表示）設定後，表示有效的 cache_hint (i64 %ch) 並產生 PTX 指令的 .L2::cache_hint 變體。

• 第 [N-1] 個引數（以 i1 %flag_mc 表示）設定後，表示存在多播遮罩 (i16 %mc) 並產生具有 .multicast::cluster 修飾符的 PTX 指令。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk。

‘llvm.nvvm.cp.async.bulk.shared.cta.to.global’

語法：

declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.global(ptr addrspace(1) %dst, ptr addrspace(3) %src, i32 %size, i64 %ch, i1 %flag_ch)

概述：

‘@llvm.nvvm.cp.async.bulk.shared.cta.to.global’ 內建函數對應於 cp.async.bulk.global.shared::cta.* 系列 PTX 指令。這些指令啟動從 shared::cta 到全域記憶體的非同步複製。32 位元運算元 %size 指定要複製的記憶體量，並且必須是 16 的倍數。

• 這些內建函數的最後一個引數是一個布林旗標，指示對 cache_hint 的支援。此旗標引數必須是編譯時常數。設定後，表示有效的 cache_hint (i64 %ch) 並產生 PTX 指令的 .L2::cache_hint 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk。

‘llvm.nvvm.cp.async.bulk.shared.cta.to.cluster’

語法：

declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3) %dst, ptr addrspace(3) %mbar, ptr addrspace(3) %src, i32 %size)

概述：

‘@llvm.nvvm.cp.async.bulk.shared.cta.to.cluster’ 內建函數對應於 cp.async.bulk.shared::cluster.shared::cta.* PTX 指令。此指令啟動從 shared::cta 到 shared::cluster 記憶體的非同步複製。目的地必須位於叢集中不同 CTA 的共享記憶體中。32 位元運算元 %size 指定要複製的記憶體量，並且必須是 16 的倍數。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk。

‘llvm.nvvm.cp.async.bulk.prefetch.L2’

語法：

declare void @llvm.nvvm.cp.async.bulk.prefetch.L2(ptr addrspace(1) %src, i32 %size, i64 %ch, i1 %flag_ch)

概述：

‘@llvm.nvvm.cp.async.bulk.prefetch.L2’ 內建函數對應於 cp.async.bulk.prefetch.L2.* 系列 PTX 指令。這些指令啟動從全域記憶體到 L2 快取的批量資料非同步預取。32 位元運算元 %size 以位元組為單位指定要預取的記憶體量，並且必須是 16 的倍數。

• 這些內建函數的最後一個引數是布林旗標，指示對 cache_hint 的支援。這些旗標引數必須是編譯時常數。設定後，表示有效的 cache_hint (i64 %ch) 並產生 PTX 指令的 .L2::cache_hint 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch。

‘llvm.nvvm.prefetch.*’

語法：

declare void  @llvm.nvvm.prefetch.global.L1(ptr addrspace(1) %global_ptr)
declare void  @llvm.nvvm.prefetch.global.L2(ptr addrspace(1) %global_ptr)
declare void  @llvm.nvvm.prefetch.local.L1(ptr addrspace(5) %local_ptr)
declare void  @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)

declare void  @llvm.nvvm.prefetch.L1(ptr %ptr)
declare void  @llvm.nvvm.prefetch.L2(ptr %ptr)

declare void  @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
declare void  @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)

declare void  @llvm.nvvm.prefetchu.L1(ptr %ptr)

概述：

‘@llvm.nvvm.prefetch.*’ 和 ‘@llvm.nvvm.prefetchu.*’ 內建函數對應於 ‘prefetch.*;’ 和 ‘prefetchu.*’ 系列 PTX 指令。‘prefetch.*’ 指令將包含全域或區域記憶體位址空間中指定位址的快取行帶入指定的快取層級（L1 或 L2）。‘prefetchu.*`’ 指令將包含指定通用位址的快取行帶入指定的統一快取層級。如果未指定位址空間，則假定為通用位址。內建函數使用和逐出優先順序，可以透過 ‘.level::eviction_priority’ 修飾符存取。

• 預取到共享記憶體位置不執行任何操作。

• 預取到統一快取需要通用位址，如果位址映射到常數、區域或共享記憶體位置，則不會發生任何操作。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-prefetch-prefetchu。

‘llvm.nvvm.applypriority.*’

語法：

declare void  @llvm.nvvm.applypriority.global.L2.evict.normal(ptr addrspace(1) %global_ptr, i64 %size)
declare void  @llvm.nvvm.applypriority.L2.evict.normal(ptr %ptr, i64 %size)

概述：

‘@llvm.nvvm.applypriority.*’ 將由 .level::eviction_priority 限定詞指定的快取逐出優先權，套用至指定快取層級中的位址範圍 [a..a+size)。如果未指定狀態空間，則會使用通用定址。如果指定的位址未落在 .global 狀態空間的位址視窗內，則行為未定義。運算元大小是一個整數常數，用於指定要套用優先權的指定快取層級中的資料量（以位元組為單位）。大小運算元唯一支援的值為 128。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-applypriority。

‘llvm.nvvm.discard.*’

語法：

declare void  @llvm.nvvm.discard.global.L2(ptr addrspace(1) %global_ptr, i64 immarg)
declare void  @llvm.nvvm.discard.L2(ptr %ptr, i64 immarg)

概述：

@llvm.nvvm.discard.L2* 內建函數的效果與非原子、非揮發性的 llvm.memset 相同，後者會將 undef 寫入至目的地位址範圍 [%ptr, %ptr + immarg)。 %ptr 必須以 128 位元組對齊。在記憶體被其他值覆寫之前，後續從該位址範圍讀取可能會讀取到 undef。這些操作提示實作，L2 快取中的資料可以破壞性地捨棄，而無需寫回記憶體。運算元 immarg 是一個整數常數，用於指定要將 undef 寫入的位址範圍 [%ptr, %ptr + immarg) 的長度（以位元組為單位）。 immarg 運算元唯一支援的值為 128。如果使用通用定址，且指定的位址未落在全域記憶體 (addrspace(1)) 的位址視窗內，則行為未定義。

call void @llvm.nvvm.discard.L2(ptr %p, i64 128)  ;; writes `undef` to [p, p+128)
%a = load i64, ptr %p. ;; loads 8 bytes containing undef
%b = load i64, ptr %p  ;; loads 8 bytes containing undef
;; comparing %a and %b compares `undef` values!
%fa = freeze i64 %a  ;; freezes undef to stable bit-pattern
%fb = freeze i64 %b  ;; freezes undef to stable bit-pattern
;; %fa may compare different to %fb!

如需更多資訊，請參閱 CUDA C++ 捨棄文件 以及 PTX ISA 捨棄文件。

‘llvm.nvvm.cp.async.bulk.tensor.g2s.tile.[1-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(..., i32 %d0, i32 %d1, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.g2s.tile.[1-5]d’ 內建函數對應於 cp.async.bulk.tensor.[1-5]d.* PTX 指令集。這些指令啟動從全域記憶體到 shared::cluster 記憶體（由 g2s 字首指示）的張量資料非同步複製，採用 tile 模式。在 tile 模式下，來源張量的多維佈局會在目的地保留。張量資料的維度範圍從 1d 到 5d，座標由 i32 %d0 ... i32 %d4 引數指定。

• 這些內建函數的最後兩個引數是布林旗標，指示對 cache_hint 和/或 multicast 修飾符的支援。這些旗標引數必須是編譯時常數。後端會檢查這些旗標並適當地降低內建函數。

• 當設定第 N 個引數（以 i1 flag_ch 表示）時，表示有效的快取提示 (i64 %ch) 並產生 PTX 指令的 .L2::cache_hint 變體。

• 當設定第 [N-1] 個引數（以 i1 flag_mc 表示）時，表示存在多播遮罩 (i16 %mc) 並產生具有 .multicast::cluster 修飾符的 PTX 指令。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor。

‘llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.[3-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.[3-5]d’ 內建函數對應於 cp.async.bulk.tensor.[1-5]d.* PTX 指令集。這些指令啟動從全域記憶體到 shared::cluster 記憶體（由 g2s 字首指示）的張量資料非同步複製，採用 im2col 模式。在 im2col 模式下，來源張量的某些維度會展開到目的地的單一維度欄中。在此模式下，張量必須至少為三維。除了張量座標外，還指定了 im2col 偏移量（以 i16 im2col0...i16 %im2col2 表示）。im2col 偏移量的數量比張量運算的維度數量少兩個。這些內建函數的最後兩個引數是布林旗標，其功能與上述 tile 模式內建函數中描述的功能相同。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor。

‘llvm.nvvm.cp.async.bulk.tensor.s2g.tile.[1-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.2d(..., i32 %d0, i32 %d1, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.s2g.tile.[1-5]d’ 內建函數對應於 cp.async.bulk.tensor.[1-5]d.* PTX 指令集。這些指令啟動從 shared::cta 到全域記憶體（由 s2g 字首指示）的張量資料非同步複製，採用 tile 模式。張量資料的維度範圍從 1d 到 5d，座標由 i32 %d0 ... i32 %d4 引數指定。

• 這些內建函數的最後一個引數是一個布林旗標，指示對 cache_hint 的支援。此旗標引數必須是編譯時常數。設定後，表示有效的 cache_hint (i64 %ch) 並產生 PTX 指令的 .L2::cache_hint 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor。

‘llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.[3-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.[1-5]d’ 內建函數對應於 cp.async.bulk.tensor.[1-5]d.* PTX 指令集。這些指令啟動從 shared::cta 到全域記憶體（由 s2g 字首指示）的張量資料非同步複製，採用 im2col 模式。在此模式下，張量必須至少為三維。與 g2s 變體不同，這些內建函數沒有 im2col_offsets。這些內建函數的最後一個引數是布林旗標，其功能與上述 s2g.tile 模式內建函數中描述的功能相同。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor。

‘llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.[1-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.1d(ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.2d(..., i32 %d0, i32 %d1, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.[1-5]d’ 內建函數對應於 cp.async.bulk.prefetch.tensor.[1-5]d.L2.global* PTX 指令集。這些指令啟動從全域記憶體到 L2 快取的張量資料非同步預取。在 tile 模式下，來源張量的多維佈局會在目的地保留。張量資料的維度範圍從 1d 到 5d，座標由 i32 %d0 ... i32 %d4 引數指定。

• 這些內建函數的最後一個引數是一個布林旗標，指示對 cache_hint 的支援。此旗標引數必須是編譯時常數。設定後，表示有效的 cache_hint (i64 %ch) 並產生 PTX 指令的 .L2::cache_hint 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor。

‘llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.[3-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.3d(ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.prefetch.im2col.[3-5]d’ 內建函數對應於 cp.async.bulk.prefetch.tensor.[1-5]d.L2.global* PTX 指令集。這些指令啟動從全域記憶體到 L2 快取的張量資料非同步預取。在 im2col 模式下，來源張量的某些維度會展開到目的地的單一維度欄中。在此模式下，張量必須至少為三維。除了張量座標外，還指定了 im2col 偏移量（以 i16 im2col0...i16 %im2col2 表示）。im2col 偏移量的數量比張量運算的維度數量少兩個。這些內建函數的最後一個引數是布林旗標，其功能與上述 tile 模式內建函數中描述的功能相同。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor。

‘llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].tile.[1-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.add.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.min.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.max.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.inc.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.dec.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.and.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.or.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.xor.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)

declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.2d(..., i32 %d0, i32 %d1, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.[1-5]d’ 內建函數對應於 cp.reduce.async.bulk.tensor.[1-5]d.* PTX 指令集。這些指令使用 tile 模式，啟動全域記憶體中張量資料與 shared{::cta} 記憶體中張量資料的非同步縮減運算。張量資料的維度範圍從 1d 到 5d，座標由 i32 %d0 ... i32 %d4 引數指定。支援的縮減運算為 {add、min、max、inc、dec、and、or、xor}，如 tile.1d 內建函數中所述。

• 這些內建函數的最後一個引數是一個布林旗標，指示對 cache_hint 的支援。此旗標引數必須是編譯時常數。設定後，表示有效的 cache_hint (i64 %ch) 並產生 PTX 指令的 .L2::cache_hint 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor。

‘llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].im2col.[3-5]d’

語法：

declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag_ch)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)

概述：

‘@llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.[3-5]d’ 內建函數對應於 cp.reduce.async.bulk.tensor.[3-5]d.* PTX 指令集。這些指令使用 im2col 模式，啟動全域記憶體中張量資料與 shared{::cta} 記憶體中張量資料的非同步縮減運算。在此模式下，張量必須至少為三維。支援的縮減運算與 tile 模式中的運算相同。這些內建函數的最後一個引數是布林旗標，其功能與上述 tile 模式內建函數中描述的功能相同。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor。

Warp Group 內建函數

‘llvm.nvvm.wgmma.fence.sync.aligned’

語法：

declare void @llvm.nvvm.wgmma.fence.sync.aligned()

概述：

‘@llvm.nvvm.wgmma.fence.sync.aligned’ 內建函數產生 wgmma.fence.sync.aligned PTX 指令，該指令在先前對任何 warpgroup 暫存器的存取與後續由 wgmma.mma_async 指令對相同暫存器的存取之間建立順序。

wgmma.fence 指令必須由 warpgroup 的所有 warp 在以下位置發出

• 在 warpgroup 中的第一個 wgmma.mma_async 運算之前。

• 在 warpgroup 中執行緒的暫存器存取與任何存取相同暫存器的 wgmma.mma_async 指令之間，除非這些是多個相同形狀的 wgmma.mma_async 指令之間的累加器暫存器存取，在這種情況下，預設會提供排序保證。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-fence。

‘llvm.nvvm.wgmma.commit_group.sync.aligned’

語法：

declare void @llvm.nvvm.wgmma.commit_group.sync.aligned()

概述：

‘@llvm.nvvm.wgmma.commit_group.sync.aligned’ 內建函數產生 wgmma.commit_group.sync.aligned PTX 指令，該指令為每個 warpgroup 建立新的 wgmma-group，並將執行 warp 啟動但未提交至任何 wgmma-group 的所有先前 wgmma.mma_async 指令批次處理到新的 wgmma-group 中。如果沒有未提交的 wgmma mma_async 指令，則 wgmma.commit_group 會產生空的 wgmma-group。

執行緒可以使用 wgmma.wait_group 等待 wgmma-group 中所有 wgmma.mma_async 運算的完成。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-commit-group。

‘llvm.nvvm.wgmma.wait_group.sync.aligned’

語法：

declare void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64 immarg N)

概述：

‘@llvm.nvvm.wgmma.wait_group.sync.aligned’ 內建函數產生 wgmma.commit_group.sync.aligned N PTX 指令，這會導致執行緒等待，直到只有 N 個或更少個最新的 wgmma-group 處於擱置狀態，且執行緒提交的所有先前 wgmma-group 都已完成。例如，當 N 為 0 時，執行緒會等待所有先前的 wgmma-group 完成。運算元 N 是一個整數常數。

存取累加器暫存器或包含 wgmma.mma_async 指令的矩陣 A 片段的輸入暫存器，而未先執行等待包含該 wgmma.mma_async 指令的 wgmma-group 的 wgmma.wait_group 指令，則行為未定義。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#asynchronous-warpgroup-level-matrix-instructions-wgmma-wait-group。

‘llvm.nvvm.griddepcontrol.*’

語法：

declare void @llvm.nvvm.griddepcontrol.launch_dependents()
declare void @llvm.nvvm.griddepcontrol.wait()

概述：

griddepcontrol 內建函數允許由執行時期定義的相依網格和先決條件網格，以以下方式控制執行：

griddepcontrol.launch_dependents 內建函數發出訊號，表示相依項目可以在目前網格完成之前排程。內建函數可以由目前 CTA 中的多個執行緒調用，且重複調用內建函數除了第一次調用的副作用之外，不會產生其他副作用。

griddepcontrol.wait 內建函數會導致執行緒等待，直到所有正在執行的先決條件網格都已完成，且先決條件網格的所有記憶體運算都已執行並對目前網格可見。

如需更多資訊，請參閱 PTX ISA。

TCGEN05 系列內建函數

llvm.nvvm.tcgen05.* 內建函數模擬由 PTX 公開的 TCGEN05 系列指令。這些內建函數使用「張量記憶體」（以下簡稱 tmem）。NVPTX 使用 addrspace(6) 表示此記憶體，且始終為 32 位元。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory。

張量記憶體指標只能與 tcgen05 內建函數搭配使用。提供了專用的載入/儲存指令 (tcgen05.ld/st) 來與張量記憶體搭配使用。

請參閱 PTX ISA 以取得有關張量記憶體載入/儲存指令的更多資訊 https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-and-register-load-store-instructions。

‘llvm.nvvm.tcgen05.alloc’

語法：

declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %dst, i32 %ncols)
declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %dst, i32 %ncols)
declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %dst, i32 %ncols)
declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %dst, i32 %ncols)

概述：

‘@llvm.nvvm.tcgen05.alloc.*’ 內建函數對應於 tcgen05.alloc.cta_group*.sync.aligned.b32 PTX 指令系列。tcgen05.alloc 是一個可能封鎖的指令，它會動態配置張量記憶體中指定數量的欄，並將已配置張量記憶體的位址寫入至由 %dst 指定的位置的共享記憶體中。32 位元運算元 %ncols 指定要配置的欄數，且必須是 2 的冪。.shared 變體明確地將共享記憶體位址空間用於 %dst 運算元。.cg1 和 .cg2 變體分別產生指令的 cta_group::1 和 cta_group::2 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions。

‘llvm.nvvm.tcgen05.dealloc’

語法：

declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)

概述：

‘@llvm.nvvm.tcgen05.dealloc.*’ 內建函數對應於 tcgen05.dealloc.* PTX 指令集。tcgen05.dealloc 指令會解除配置由張量記憶體位址 %tmem_addr 指定的張量記憶體。運算元 %tmem_addr 必須指向先前的張量記憶體配置。32 位元運算元 %ncols 指定要解除配置的欄數。.cg1 和 .cg2 變體分別產生指令的 cta_group::1 和 cta_group::2 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions。

‘llvm.nvvm.tcgen05.relinq.alloc.permit’

語法：

declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1()
declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()

概述：

‘@llvm.nvvm.tcgen05.relinq.alloc.permit.*’ 內建函數對應於 tcgen05.relinquish_alloc_permit.* PTX 指令集。此指令指定執行緒的 CTA 放棄配置張量記憶體的權利。因此，在任何組成執行緒執行 tcgen05.relinquish_alloc_permit 之後，CTA 執行 tcgen05.alloc 是非法的。.cg1 和 .cg2 變體分別產生指令的 cta_group::1 和 cta_group::2 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions。

‘llvm.nvvm.tcgen05.commit’

語法：

declare void @llvm.nvvm.tcgen05.commit.{cg1,cg2}(ptr %mbar)
declare void @llvm.nvvm.tcgen05.commit.shared.{cg1,cg2}(ptr addrspace(3) %mbar)
declare void @llvm.nvvm.tcgen05.commit.mc.{cg1,cg2}(ptr %mbar, i16 %mc)
declare void @llvm.nvvm.tcgen05.commit.mc.shared.{cg1,cg2}(ptr addrspace(3) %mbar, i16 %mc)

概述：

‘@llvm.nvvm.tcgen05.commit.*’ 內建函數對應於 tcgen05.commit.{cg1/cg2}.mbarrier::arrive::one.* PTX 指令集。tcgen05.commit 是一個非同步指令，它使 mbarrier 物件 (%mbar) 追蹤所有先前的非同步 tcgen05 運算的完成情況。.mc 變體允許在叢集中多個 CTA（由 %mc 指定）的 mbarrier 物件上發出訊號。.cg1 和 .cg2 變體分別產生指令的 cta_group::1 和 cta_group::2 變體。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit。

‘llvm.nvvm.tcgen05.wait’

語法：

declare void @llvm.nvvm.tcgen05.wait.ld()
declare void @llvm.nvvm.tcgen05.wait.st()

概述：

‘@llvm.nvvm.tcgen05.wait.ld/st’ 內建函數對應於 tcgen05.wait::{ld/st}.sync.aligned PTX 指令對。tcgen05.wait::ld 會導致執行緒封鎖，直到執行緒發出的所有先前 tcgen05.ld 運算都已完成。tcgen05.wait::st 會導致執行緒封鎖，直到執行緒發出的所有先前 tcgen05.st 運算都已完成。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-wait。

‘llvm.nvvm.tcgen05.fence’

語法：

declare void @llvm.nvvm.tcgen05.fence.before.thread.sync()
declare void @llvm.nvvm.tcgen05.fence.after.thread.sync()

概述：

‘@llvm.nvvm.tcgen05.fence.*’ 內建函數對應於 tcgen05.fence::{before/after}_thread_sync PTX 指令對。這些指令充當非同步 tcgen05 運算的程式碼移動柵欄。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tensorcore-5th-generation-instructions-tcgen05-fence。

‘llvm.nvvm.tcgen05.shift’

語法：

declare void @llvm.nvvm.tcgen05.shift.down.cg1(ptr addrspace(6) %tmem_addr)
declare void @llvm.nvvm.tcgen05.shift.down.cg2(ptr addrspace(6) %tmem_addr)

概述：

‘@llvm.nvvm.tcgen05.shift.{cg1/cg2}’ 內建函數對應於 tcgen05.shift.{cg1/cg2} PTX 指令。tcgen05.shift 是一個非同步指令，它啟動跨所有列（最後一列除外）向下移動 32 位元組元素一個列的動作。位址運算元 %tmem_addr 指定張量記憶體中矩陣的基底位址，其列必須向下移動。

如需更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-shift。

‘llvm.nvvm.tcgen05.cp’

語法：

declare void @llvm.nvvm.tcgen05.cp.4x256b.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.128x256b.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.128x128b.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)

declare void @llvm.nvvm.tcgen05.cp.4x256b.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.128x256b.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.128x128b.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b6x16_p32.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)

declare void @llvm.nvvm.tcgen05.cp.4x256b.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.128x256b.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.128x128b.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)
declare void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b4x16_p64.{cg1,cg2}(ptr addrspace(6) %tmem_addr, i64 %sdesc)

概述：

‘@llvm.nvvm.tcgen05.cp.{shape}.{src_fmt}.{cg1/cg2}’ 內建函數對應於 tcgen05.cp.* PTX 指令系列。tcgen05.cp 指令啟動從共享記憶體到張量記憶體中由 %tmem_addr 指定的位置的非同步複製運算。64 位元暫存器運算元 %sdesc 是矩陣描述符，表示需要複製的共享記憶體中的來源矩陣。

複製運算的有效形狀為：{128x256b、4x256b、128x128b、64x128b_warpx2_02_13、64x128b_warpx2_01_23、32x128b_warpx4}。

形狀 64x128b 和 32x128b 需要專用的多播限定詞，這些限定詞會附加到對應的內建函數名稱。

或者，在複製操作期間，可以選擇將資料從共享記憶體中的來源格式解壓縮到張量記憶體中的目標格式。目前，僅支援 .b8x16 作為目標格式。有效的來源格式為 .b6x16_p32 和 .b4x16_p64。

當來源格式為 .b6x16_p32 時，共享記憶體中連續的 16 個 6 位元元素集合，後跟四個位元組的填充 (_p32)，將被解壓縮為張量記憶體中 16 個 8 位元元素 (.b8x16)。

當來源格式為 .b4x16_p64 時，共享記憶體中連續的 16 個 4 位元元素集合，後跟八個位元組的填充 (_p64)，將被解壓縮為張量記憶體中 16 個 8 位元元素 (.b8x16)。

有關解壓縮方案的更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#optional-decompression。

有關 tcgen05.cp 指令的更多資訊，請參閱 PTX ISA https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-cp。

‘llvm.nvvm.tcgen05.ld.*’

語法：

declare <n x i32> @llvm.nvvm.tcgen05.ld.<shape>.<num>(ptr addrspace(6) %tmem_addr, i1 %pack)

declare <n x i32> @llvm.nvvm.tcgen05.ld.16x32bx2.<num>(ptr addrspace(6) %tmem_addr, i64 %offset, i1 %pack)

概述：

這組 intrinsic 函式會跨 warp 中的所有執行緒，從張量記憶體中由 32 位元位址運算元 tmem_addr 指定的位置，非同步載入資料到目標暫存器中。

warp 中的所有執行緒都必須指定相同的 tmem_addr 值，該值必須是集體載入操作的基底位址。否則，行為將是未定義的。

shape 限定詞和 num 限定詞共同決定了從張量記憶體載入的資料 ('n') 的總維度。shape 限定詞表示資料的基礎維度。num 限定詞表示基礎維度上的重複因子，從而產生存取資料的總維度。

‘num’ 的允許值為 x1, x2, x4, x8, x16, x32, x64, x128。

第一個 intrinsic 函式中 ‘shape’ 的允許值為 16x64b, 16x128b, 16x256b, 32x32b。

第二個 intrinsic 函式中 ‘shape’ 的允許值為 16x32bx2。

intrinsic 函式的結果是一個向量，由一個或多個從 shape 和 num 導出的 32 位元暫存器組成，如下所示。

num/shape	16x32bx2/16x64b/32x32b	16x128b	16x256b
x1	1	2	4
x2	2	4	8
x4	4	8	16
x8	8	16	32
x16	16	32	64
x32	32	64	128
x64	64	128	NA
x128	128	NA	NA

最後一個引數 i1 %pack 是一個編譯時期常數，設定後表示在載入期間，相鄰的列會被封裝到單個 32 位元元素中。

更多資訊，請參閱 PTX ISA。

‘llvm.nvvm.tcgen05.st.*’

語法：

declare void @llvm.nvvm.tcgen05.st.<shape>.<num>(ptr addrspace(6) %tmem_addr, <n x i32> %args, i1 %unpack)

declare void @llvm.nvvm.tcgen05.st.16x32bx2.<num>(ptr addrspace(6) %tmem_addr, <n x i32> %args, i64 %offset, i1 %unpack)

概述：

這組 intrinsic 函式會跨 warp 中的所有執行緒，從來源向量非同步地將資料儲存到張量記憶體中由 32 位元位址運算元 ‘tmem_addr’ 指定的位置。

warp 中的所有執行緒都必須指定相同的 tmem_addr 值，該值必須是集體載入操作的基底位址。否則，行為將是未定義的。

shape 限定詞和 num 限定詞共同決定了從張量記憶體載入的資料 ('n') 的總維度。shape 限定詞表示資料的基礎維度。num 限定詞表示基礎維度上的重複因子，從而產生存取資料的總維度。

‘num’ 的允許值為 x1, x2, x4, x8, x16, x32, x64, x128。

第一個 intrinsic 函式中 ‘shape’ 的允許值為 16x64b, 16x128b, 16x256b, 32x32b。

第二個 intrinsic 函式中 ‘shape’ 的允許值為 16x32bx2。

args 引數是一個向量，由一個或多個從 shape 和 num 導出的 32 位元暫存器組成，如 tcgen05.ld 章節的表格中所列。

每個 shape 都支援 unpack 模式，允許將暫存器中的 32 位元元素解封裝成兩個 16 位元元素，並將它們儲存在相鄰的列中。可以透過將 %unpack 運算元設定為 1 來啟用 unpack 模式，設定為 0 則停用。

最後一個引數 i1 %unpack 是一個編譯時期常數，設定後表示暫存器中的 32 位元元素將被解封裝成兩個 16 位元元素，並儲存在相鄰的列中。

更多資訊，請參閱 PTX ISA。

其他 Intrinsic 函式

如需完整的 NVPTX intrinsic 函式集，請參閱 LLVM 原始碼樹中的 include/llvm/IR/IntrinsicsNVVM.td 檔案。

與 Libdevice 連結

CUDA 工具組隨附一個名為 libdevice 的 LLVM 位元碼函式庫，其中實作了許多常見的數學函式。對於任何使用 LLVM NVPTX 目標的編譯器，此函式庫都可以作為高效能數學函式庫使用。該函式庫可以在 CUDA 工具組的 nvvm/libdevice/ 下找到，並且每個運算架構都有一個單獨的版本。

如需 libdevice 中實作的所有數學函式列表，請參閱 libdevice 使用者指南。

為了適應各種可能影響 libdevice 代碼產生的數學相關編譯器標誌，該函式庫代碼依賴特殊的 LLVM IR pass (NVVMReflect) 來處理 LLVM IR 中的條件編譯。此 pass 會尋找對 @__nvvm_reflect 函式的呼叫，並根據定義的反射參數將其替換為常數。此類條件代碼通常遵循以下模式

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

所有未指定反射參數的預設值為零。

NVVMReflect pass 應在優化流程的早期執行，緊接在連結階段之後。internalize pass 也建議用於從產生的 PTX 中移除未使用的數學函式。對於輸入 IR 模組 module.bc，建議使用以下編譯流程

NVVMReflect pass 即使在沒有優化的情況下，也會嘗試移除無效代碼。這允許透過使用 __CUDA_ARCH 引數，在所有優化層級避免潛在的不相容指令。

1. 將外部函式列表儲存在 module.bc 中

2. 將 module.bc 與 libdevice.compute_XX.YY.bc 連結

3. 將 (1) 中列表以外的所有函式內部化

4. 消除所有未使用的內部函式

5. 執行 NVVMReflect pass

6. 執行標準優化流程

注意

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 函式庫目前使用以下反射參數來控制代碼產生

旗標	描述
__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 代碼即時編譯 (JIT) 為原生 GPU 機器碼。

初始化 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 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"

目標 Intrinsic 函式

在本範例中，我們使用 @llvm.nvvm.read.ptx.sreg.tid.x intrinsic 函式來讀取目前執行緒 ID 的 X 元件，這相當於讀取 PTX 中的暫存器 %tid.x。NVPTX 後端支援大量的 intrinsic 函式。下面顯示一個簡短列表；完整列表請參閱 include/llvm/IR/IntrinsicsNVVM.td。

Intrinsic 函式	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 裝置 (通常) 有四種類型的記憶體

• 全域：大型、晶片外記憶體

• 共享：小型、晶片上記憶體，在 CTA 中的所有執行緒之間共享

• 區域：每個執行緒的私有記憶體

• 常數：跨所有執行緒共享的唯讀記憶體

這些不同類型的記憶體在 LLVM IR 中表示為位址空間。NVPTX 代碼產生器還使用了第五個位址空間，對應於「通用」位址空間。此位址空間可以表示任何其他位址空間中的位址 (少數例外)。這允許使用者使用相同的指令編寫可以載入/儲存記憶體的 IR 函式。Intrinsic 函式用於在通用和非通用位址空間之間轉換指標。

有關更多資訊，請參閱位址空間 和 NVPTX Intrinsic 函式。

核心元數據

在 PTX 中，函式可以是 kernel 函式 (可從主機程式呼叫)，也可以是 device 函式 (僅可從 GPU 代碼呼叫)。您可以將 kernel 函式視為 GPU 程式中的進入點。為了將 LLVM IR 函式標記為 kernel 函式，我們使用特殊的 LLVM 元數據。NVPTX 後端將尋找名為 nvvm.annotations 的具名元數據節點。此具名元數據必須包含描述 IR 的元數據列表。就我們的目的而言，我們需要宣告一個元數據節點，將 “kernel” 屬性指派給應作為 PTX kernel 函式發出的 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 並將其即時編譯 (JIT) 為原生 GPU 裝置，並啟動核心。此 API 類似於 OpenCL。下面顯示了一個簡單的範例，說明如何載入和執行我們的向量加法代碼。請注意，為簡潔起見，此代碼未執行太多錯誤檢查！

注意

您也可以使用 CUDA 工具組提供的 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 = 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}

若要編譯此核心，我們執行以下步驟

1. 與 libdevice 連結

2. 將除公開核心函式之外的所有內容內部化

3. 執行 NVVMReflect 並將 __CUDA_FTZ 設定為 0

4. 優化連結的模組

5. 代碼產生模組

這些步驟可以由 LLVM llvm-link、opt 和 llc 工具執行。在完整的編譯器中，這些步驟也可以透過設定適當的 pass 配置以程式化的方式完全執行 (請參閱 與 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;
}