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>"

这表示对编译器/驱动程序的提示/指令，要求在单个 SM 上至少放置这么多 CTA。

"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.nvda.net.cn/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.nvda.net.cn/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.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 为真，则选择最高有效字节，否则选择最低有效字节。 然后将这些字节扩展为 32 位。 对于 ‘llvm.nvvm.idp2a.[us].u’ 变体，使用零扩展，而对于 ‘llvm.nvvm.idp2a.[us].s’，则使用符号扩展。 这两个元素向量的点积被添加到 %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’ 系列内建函数执行 clamped funnel shift left (钳位漏斗左移)。 这些内建函数与 ‘llvm.fshl’ 非常相似，只是移位量被钳位在整数宽度处（而不是对其取模）。 目前，仅支持 i32。

语义：

‘llvm.nvvm.fshl.clamp’ 系列内建函数执行 clamped funnel shift left：前两个值连接为 { %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’ 系列内建函数执行 clamped funnel shift right (钳位漏斗右移)。 这些内建函数与 ‘llvm.fshr’ 非常相似，只是移位量被钳位在整数宽度处（而不是对其取模）。 目前，仅支持 i32。

语义：

‘llvm.nvvm.fshr.clamp’ 系列内建函数执行 clamped funnel shift right：前两个值连接为 { %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’ 系列内建函数识别前导 1 的位位置，返回其与最高有效位或最低有效位的偏移量。

语义：

‘llvm.nvvm.flo.u’ 系列内建函数返回最高有效位 1 的位位置。 如果 %shiftamt 为真，则结果是将找到的位左移到最高有效位位置所需的移位量，否则结果是将找到的位右移到最低有效位位置所需的移位量。 如果未找到 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 为真，则结果是将找到的位左移到最高有效位位置所需的移位量，否则结果是将找到的位右移到最低有效位位置所需的移位量。 如果未找到 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.nvda.net.cn/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.nvda.net.cn/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.nvda.net.cn/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.nvda.net.cn/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’ 修饰符访问。

• 预取到共享内存位置不执行任何操作。

• 预取到统一缓存需要通用地址，如果地址映射到 const、local 或共享内存位置，则不会发生任何操作。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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.*’ 对指定缓存级别中地址范围 [a..a+size) 应用由 .level::eviction_priority 限定符指定的缓存驱逐优先级。如果未指定状态空间，则使用通用寻址。如果指定的地址不在 .global 状态空间的地址窗口内，则行为未定义。操作数 size 是一个整数常量，用于指定要应用优先级的指定缓存级别中的数据量（以字节为单位）。size 操作数唯一支持的值为 128。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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* intrinsic 的效果类似于非原子、非易失的 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’ intrinsic 对应于 PTX 指令的 cp.async.bulk.tensor.[1-5]d.* 集。这些指令启动从全局内存到 shared::cluster 内存（由 g2s 前缀指示）的张量数据的异步复制，采用 tile 模式。在 tile 模式下，源张量的多维布局在目标位置被保留。张量数据的维度范围为 1d 到 5d，坐标由 i32 %d0 ... i32 %d4 参数指定。

• 这些内建函数的最后两个参数是布尔标志，指示对 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.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 cp.async.bulk.tensor.[1-5]d.* 集。这些指令启动从全局内存到 shared::cluster 内存（由 g2s 前缀指示）的张量数据的异步复制，采用 im2col 模式。在 im2col 模式下，源张量的某些维度被展开为目标位置的单维列。在这种模式下，张量必须至少是三维的。除了张量坐标外，还指定了 im2col 偏移量（由 i16 im2col0...i16 %im2col2 表示）。im2col 偏移量的数量比张量运算的维度数量少两个。这些 intrinsic 的最后两个参数是布尔标志，其功能与上面 tile 模式 intrinsic 中描述的功能相同。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 cp.async.bulk.tensor.[1-5]d.* 集。这些指令启动从 shared::cta 到全局内存（由 s2g 前缀指示）的张量数据的异步复制，采用 tile 模式。张量数据的维度范围为 1d 到 5d，坐标由 i32 %d0 ... i32 %d4 参数指定。

• 这些内建函数的最后一个参数是布尔标志，指示对 cache_hint 的支持。 此标志参数必须是编译时常量。 设置后，它表示有效的 cache_hint (i64 %ch) 并生成 PTX 指令的 .L2::cache_hint 变体。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 cp.async.bulk.tensor.[1-5]d.* 集。这些指令启动从 shared::cta 到全局内存（由 s2g 前缀指示）的张量数据的异步复制，采用 im2col 模式。在这种模式下，张量必须至少是三维的。与 g2s 变体不同，这些 intrinsic 没有 im2col_offsets。这些 intrinsic 的最后一个参数是布尔标志，其功能与上面 s2g.tile 模式 intrinsic 中描述的功能相同。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 cp.async.bulk.prefetch.tensor.[1-5]d.L2.global* 集。这些指令启动从全局内存到 L2 缓存的张量数据的异步预取。在 tile 模式下，源张量的多维布局在目标位置被保留。张量数据的维度范围为 1d 到 5d，坐标由 i32 %d0 ... i32 %d4 参数指定。

• 这些内建函数的最后一个参数是布尔标志，指示对 cache_hint 的支持。 此标志参数必须是编译时常量。 设置后，它表示有效的 cache_hint (i64 %ch) 并生成 PTX 指令的 .L2::cache_hint 变体。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 cp.async.bulk.prefetch.tensor.[1-5]d.L2.global* 集。这些指令启动从全局内存到 L2 缓存的张量数据的异步预取。在 im2col 模式下，源张量的某些维度被展开为目标位置的单维列。在这种模式下，张量必须至少是三维的。除了张量坐标外，还指定了 im2col 偏移量（由 i16 im2col0...i16 %im2col2 表示）。im2col 偏移量的数量比张量运算的维度数量少两个。这些 intrinsic 的最后一个参数是布尔标志，其功能与上面 tile 模式 intrinsic 中描述的功能相同。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 cp.reduce.async.bulk.tensor.[1-5]d.* 集。这些指令启动全局内存中的张量数据与 shared{::cta} 内存中的张量数据之间的异步归约运算，采用 tile 模式。张量数据的维度范围为 1d 到 5d，坐标由 i32 %d0 ... i32 %d4 参数指定。支持的归约运算包括 {add, min, max, inc, dec, and, or, xor}，如 tile.1d intrinsic 中所述。

• 这些内建函数的最后一个参数是布尔标志，指示对 cache_hint 的支持。 此标志参数必须是编译时常量。 设置后，它表示有效的 cache_hint (i64 %ch) 并生成 PTX 指令的 .L2::cache_hint 变体。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 cp.reduce.async.bulk.tensor.[3-5]d.* 集。这些指令启动全局内存中的张量数据与 shared{::cta} 内存中的张量数据之间的异步归约运算，采用 im2col 模式。在这种模式下，张量必须至少是三维的。支持的归约运算与 tile 模式中的相同。这些 intrinsic 的最后一个参数是布尔标志，其功能与上面 tile 模式 intrinsic 中描述的功能相同。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor。

Warp Group Intrinsics

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

语法：

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

概述：

‘@llvm.nvvm.wgmma.fence.sync.aligned’ intrinsic 生成 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.nvda.net.cn/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’ intrinsic 生成 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.nvda.net.cn/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’ intrinsic 生成 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.nvda.net.cn/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 intrinsic 允许运行时定义的依赖网格和先决条件网格以以下方式控制执行

griddepcontrol.launch_dependents intrinsic 发出信号，表明可以在当前网格完成之前调度依赖项。 intrinsic 可以由当前 CTA 中的多个线程调用，并且重复调用 intrinsic 除了第一次调用的副作用外，不会产生额外的副作用。

griddepcontrol.wait intrinsic 导致执行线程等待，直到所有正在运行的先决条件网格完成，并且来自先决条件网格的所有内存操作都已执行并对当前网格可见。

有关更多信息，请参阅 PTX ISA。

TCGEN05 系列 Intrinsics

llvm.nvvm.tcgen05.* intrinsic 建模了 PTX 公开的 TCGEN05 系列指令。这些 intrinsic 使用“张量内存”（以下简称 tmem）。 NVPTX 使用 addrspace(6) 表示此内存，并且始终为 32 位。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/cuda/parallel-thread-execution/#tensor-memory。

张量内存指针只能与 tcgen05 intrinsic 一起使用。提供了专门的加载/存储指令 (tcgen05.ld/st) 来处理张量内存。

有关张量内存加载/存储指令的更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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.*’ intrinsic 对应于 PTX 指令的 tcgen05.alloc.cta_group*.sync.aligned.b32 系列。 tcgen05.alloc 是一个潜在的阻塞指令，它在张量内存中动态分配指定数量的列，并将已分配的张量内存的地址写入 %dst 指定位置的共享内存中。 32 位操作数 %ncols 指定要分配的列数，并且它必须是 2 的幂。 .shared 变体显式使用共享内存地址空间作为 %dst 操作数。 .cg1 和 .cg2 变体分别生成指令的 cta_group::1 和 cta_group::2 变体。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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.*’ intrinsic 对应于 PTX 指令的 tcgen05.dealloc.* 集。 tcgen05.dealloc 指令释放由张量内存地址 %tmem_addr 指定的张量内存。操作数 %tmem_addr 必须指向先前的张量内存分配。 32 位操作数 %ncols 指定要释放分配的列数。 .cg1 和 .cg2 变体分别生成指令的 cta_group::1 和 cta_group::2 变体。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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.*’ intrinsic 对应于 PTX 指令的 tcgen05.relinquish_alloc_permit.* 集。此指令指定执行线程的 CTA 正在放弃分配张量内存的权利。因此，在任何组成线程执行 tcgen05.relinquish_alloc_permit 之后，CTA 执行 tcgen05.alloc 是非法的。 .cg1 和 .cg2 变体分别生成指令的 cta_group::1 和 cta_group::2 类型。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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.*’ intrinsic 对应于 PTX 指令的 tcgen05.commit.{cg1/cg2}.mbarrier::arrive::one.* 集。 tcgen05.commit 是一个异步指令，它使 mbarrier 对象 (%mbar) 跟踪所有先前的异步 tcgen05 操作的完成情况。 .mc 变体允许在集群中多个 CTA（由 %mc 指定）的 mbarrier 对象上发出信号。 .cg1 和 .cg2 变体分别生成指令的 cta_group::1 和 cta_group::2 类型。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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’ intrinsic 对应于 PTX 指令的 tcgen05.wait::{ld/st}.sync.aligned 对。 tcgen05.wait::ld 导致执行线程阻塞，直到执行线程发出的所有先前的 tcgen05.ld 操作完成。 tcgen05.wait::st 导致执行线程阻塞，直到执行线程发出的所有先前的 tcgen05.st 操作完成。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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.*’ intrinsic 对应于 PTX 指令的 tcgen05.fence::{before/after}_thread_sync 对。这些指令充当异步 tcgen05 操作的代码移动栅栏。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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}’ intrinsic 对应于 PTX 指令的 tcgen05.shift.{cg1/cg2}。 tcgen05.shift 是一个异步指令，它启动跨所有行（最后一行除外）向下移动 32 字节元素的操作，移动一行。地址操作数 %tmem_addr 指定张量内存中矩阵的基地址，该矩阵的行必须向下移动。

有关更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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}’ intrinsic 对应于 PTX 指令的 tcgen05.cp.* 系列。 tcgen05.cp 指令启动从共享内存到 %tmem_addr 在张量内存中指定位置的异步复制操作。 64 位寄存器操作数 %sdesc 是矩阵描述符，表示需要复制的共享内存中的源矩阵。

copy 操作的有效形状为：{128x256b, 4x256b, 128x128b, 64x128b_warpx2_02_13, 64x128b_warpx2_01_23, 32x128b_warpx4}。

形状 64x128b 和 32x128b 需要专用的多播限定符，这些限定符附加到相应的固有名称之后。

可选地，数据可以在复制操作期间从共享内存中的源格式解压缩到 Tensor 内存中的目标格式。目前，仅支持 .b8x16 作为目标格式。有效的源格式为 .b6x16_p32 和 .b4x16_p64。

当源格式为 .b6x16_p32 时，共享内存中连续的 16 个 6 位元素集后跟四个字节的填充 (_p32) 将被解压缩为 Tensor 内存中每个元素 8 位的 16 个元素 (.b8x16)。

当源格式为 .b4x16_p64 时，共享内存中连续的 16 个 4 位元素集后跟八个字节的填充 (_p64) 将被解压缩为 Tensor 内存中每个元素 8 位的 16 个元素 (.b8x16)。

有关解压缩方案的更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/cuda/parallel-thread-execution/#optional-decompression。

有关 tcgen05.cp 指令的更多信息，请参阅 PTX ISA https://docs.nvda.net.cn/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 异步地将数据从 Tensor 内存中由 32 位地址操作数 tmem_addr 指定的位置加载到目标寄存器中，跨 warp 中的所有线程集合。

warp 中的所有线程必须指定相同的 tmem_addr 值，该值必须是集合加载操作的基地址。否则，行为是未定义的。

shape 限定符和 num 限定符共同决定了从 Tensor 内存加载的数据的总维度（‘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/形状	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 异步地将数据从源向量存储到 Tensor 内存中由 32 位地址操作数 ‘tmem_addr` 指定的位置，跨 warp 中的所有线程集合。

warp 中的所有线程必须指定相同的 tmem_addr 值，该值必须是集合加载操作的基地址。否则，行为是未定义的。

shape 限定符和 num 限定符共同决定了从 Tensor 内存加载的数据的总维度（‘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 部分中列出的表格所示。

每个形状都支持 unpack 模式，以允许寄存器中的 32 位元素被解包成两个 16 位元素并将它们存储在相邻的列中。unpack 模式可以通过将 %unpack 操作数设置为 1 来启用，通过将其设置为 0 来禁用。

最后一个参数 i1 %unpack 是一个编译时常量，当设置时，表示寄存器中的 32 位元素将被解包成两个 16 位元素并存储在相邻的列中。

有关更多信息，请参阅 PTX ISA。

其他 Intrinsics

有关 NVPTX intrinsic 的完整集合，请参阅 LLVM 源代码树中的 include/llvm/IR/IntrinsicsNVVM.td 文件。

与 Libdevice 链接

CUDA 工具包附带一个名为 libdevice 的 LLVM bitcode 库，该库实现了许多常见的数学函数。此库可以用作任何使用 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"

目标 Intrinsics

在此示例中，我们使用 @llvm.nvvm.read.ptx.sreg.tid.x intrinsic 读取当前线程 ID 的 X 分量，这对应于读取 PTX 中的寄存器 %tid.x。NVPTX 后端支持大量的 intrinsics。下面显示了一个简短列表；请参阅 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 Intrinsics。

内核元数据

在 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. Codegen 模块

这些步骤可以通过 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;
}