NVVM IR 规范
基于LLVM IR的NVVM编译器（中间表示）参考指南。
1. 简介
NVVM IR是一种基于LLVM IR的编译器中间表示。NVVM IR旨在表示GPU计算内核（例如CUDA内核）。高级语言前端（如CUDA C编译器前端）可以生成NVVM IR。NVVM编译器（基于LLVM）会从NVVM IR生成PTX代码。
NVVM IR和NVVM编译器基本与所使用的源语言无关。由于DCI（驱动程序/编译器接口）的差异，NVVM编译器的PTX代码生成部分需要了解源语言。
NVVM IR是一种二进制格式，基于LLVM IR位码格式的子集。本文档仅使用人类可读的形式来描述NVVM IR。
从技术角度来说，NVVM IR是带有特定规则、限制和约定的LLVM IR，外加一组支持的内置函数。用NVVM IR指定的程序始终是合法的LLVM程序，但合法的LLVM程序未必是合法的NVVM程序。
NVVM IR 有三个级别的支持。
支持：该功能已完全支持。大多数IR功能应属于此类别。
接受但忽略：NVVM编译器会接受此IR特性，但会忽略所需的语义。这适用于某些在GPU上没有实际意义且可被忽略的IR特性。调用约定标记就是一个例子。
非法且不受支持：指定的语义不受支持，例如
fence指令。未来版本的NVVM可能会支持或接受并忽略当前版本中非法的IR。
本文档描述了NVVM IR 2.0版本和NVVM调试元数据3.1版本（参见Source Level Debugging Support）。NVVM IR 2.0版本与之前的1.11版本不兼容。将NVVM IR 1.11版本与2.0版本链接会导致编译器错误。
NVVM IR可以支持两种方言。LLVM 7方言基于LLVM 7.0.1版本。现代方言则基于LLVM较新的公开发行版本(LLVM 18.1.8)。现代方言仅支持Blackwell及之后的架构(计算能力compute_100或更高)。关于IR的完整语义，本文档读者应参考官方LLVM语言参考手册版本7或版本18。当两种NVVM IR方言存在重要差异时，本文档会添加相应注释说明。
2. 标识符
命名全局标识符的名称必须符合以下格式：
@[a-zA-Z$_][a-zA-Z$_0-9]*
注意，它不能包含 . 字符。
[@%]llvm.nvvm.* 和 [@%]nvvm.* 是保留字。
3. 高层架构
3.1. 连接类型
支持：
privateinternalavailable_externallylinkonceweakcommonlinkonce_odrweak_odrexternal
不支持所有其他链接类型。
有关链接类型如何转换为PTX的详细信息，请参阅NVVM ABI for PTX。
3.2. 调用约定
所有LLVM调用约定标记均被接受并忽略。函数和调用根据PTX调用约定生成。
3.2.1. 规则与限制
当传递宽度小于32位的参数时，应设置
zeroext/signext参数属性。如果未设置，将默认采用zeroext。当返回值的宽度小于32位时，应设置
zeroext/signext参数属性。如果未设置，将默认采用zeroext。通过值传递的聚合或向量类型参数，可以通过设置
byval属性的指针来传递（下文称为by-pointer-byval情况）。如果类型需要非自然对齐（自然对齐是根据数据布局部分推断出的聚合类型对齐方式），则必须设置align属性。如果一个函数有一个按值直接传递的聚合或向量类型参数，且该类型具有非自然对齐要求，则必须通过全局属性注解<
align, 对齐值>来标注对齐方式，其中对齐值是一个32位整数，其高16位表示参数位置（从1开始计数），低16位表示对齐值。如果函数的返回类型是一个聚合体或具有非自然对齐方式的向量，那么必须通过全局属性注解<
align, alignment>来标注对齐要求，其中高16位为0，低16位表示对齐值。不需要为函数添加<
align, alignment>注解。如果添加了注解，则必须与自然对齐方式或在by-pointer-byval情况下的align属性保持一致。-
对于返回值或某个参数具有非自然对齐要求的间接函数调用指令，且该对齐要求未在
by-pointer-byval情况下通过对齐方式表达时，调用指令必须附加类型为callalign的元数据。该元数据包含一系列i32字段，每个字段表示一个非自然对齐要求。i32字段的高16位表示参数位置（0表示返回值，1表示第一个参数，以此类推），低16位表示对齐值。这些i32字段必须按升序排列。例如：
%call = call %struct.S %fp1(%struct.S* byval align 8 %arg1p, %struct.S %arg2),!callalign !10 !10 = !{i32 8, i32 520};
其他参数或返回值并不强制要求包含
i32元数据字段。如果存在该字段，其对齐方式必须与自然对齐方式或by-pointer-byval case中的align属性保持一致。直接调用指令不需要附加
callalign元数据。如果附加了该元数据，则对齐方式必须与自然对齐方式或by-pointer-byval情况下的对齐方式匹配。间接调用指令中缺少元数据意味着在
by-pointer-byval情况下使用自然对齐或align属性。
3.3. 可见性样式
所有样式——默认、隐藏和受保护的——都被接受并忽略。
3.4. DLL存储类
不支持。
3.5. 线程本地存储模型
不支持。
3.6. 运行时抢占说明符
不支持。
3.7. 结构体类型
支持。
3.8. 非整型指针类型
不支持。
3.9. Comdats
不支持。
3.10. source_filename
已接受并忽略。
3.11. 全局变量
一个全局变量（非内置全局变量）可以选择性地声明为驻留在以下地址空间之一：
globalsharedconstant
如果没有显式指定地址空间，则假定全局变量位于global地址空间中，并具有通用地址值。详情请参阅Address Space。
thread_local 变量不受支持。
不允许有显式部分（元数据部分除外）。
不支持shared变量的初始化。请使用未定义初始化。
3.12. 函数
以下功能在函数中不受支持：
对齐
显式部分
垃圾回收器名称
前缀数据
序言
个性
3.13. 别名
仅支持作为非内核函数的别名。
3.14. Ifuncs
不支持。
3.15. 命名元数据
除以下情况外，已接受并忽略：
!nvvm.annotations: 参见 Global Property Annotation!nvvmir.version!llvm.dbg.cu!llvm.module.flags
NVVM IR版本通过名为!nvvmir.version的命名元数据指定。该!nvvmir.version命名元数据可以包含一个元数据节点，用于表示该模块的NVVM IR版本。如果多个此类模块被链接在一起，链接后模块中的命名元数据可能包含多个元数据节点，每个节点都包含一个版本号。表示NVVM IR版本的元数据节点可采用以下任一形式：
-
它可能包含两个i32值——第一个表示NVVM IR的主版本号，第二个表示次版本号。如果未指定，则默认版本号为1.0，可以表示为：
!nvvmir.version = !{!0} !0 = !{i32 1, i32 0} -
它可能包含四个i32值——前两个分别表示NVVM IR的主版本号和次版本号。第三个值表示NVVM IR调试元数据的主版本号，第四个值表示对应的次版本号。如果未指定，则假定版本号为1.0，可以表示为：
!nvvmir.version = !{!0} !0 = !{i32 1, i32 0, i32 1, i32 0}
本文档描述的NVVM IR版本为2.0。本文档描述的NVVM IR调试元数据版本为3.1。
3.16. 参数属性
除以下情况外均支持：
已接受并忽略：
inregnest
不支持所有其他参数属性。
查看调用约定了解这些属性的用法。
3.17. 垃圾回收器策略名称
不支持。
3.18. 前缀数据
不支持。
3.19. 序章数据
不支持。
3.20. 属性分组
支持。支持的属性集等同于使用该属性组时接受的属性集。
3.21. 函数属性
支持：
allocsizealwaysinlinecoldconvergentinaccessiblememonlyinaccessiblemem_or_argmemonlyinlinehintminsizeno-jump-tablesnoduplicatenoinlinenoreturnnorecursenounwind"null-pointer-is-valid"optforfuzzingoptnoneoptsizereadnonereadonlywriteonlyargmemonlyspeculatablestrictfp
不支持所有其他函数属性。
3.22. 全局属性
不支持。
3.23. 操作数绑定集
不支持。
3.24. 模块级内联汇编
支持。
3.25. 数据布局
仅支持以下数据布局：
-
64位
e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64
以下数据布局已被弃用，将在未来的版本中移除。
-
32位
e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64
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位
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
3.26. 目标三元组
仅支持以下目标三元组，其中*可以是任意名称：
64位:
nvptx64-*-cuda
以下目标三元组已弃用，将在未来版本中移除：
32位:
nvptx-*-cuda
3.27. 指针别名规则
支持。
3.28. 易失性内存访问
支持。请注意，在代码生成时：ld.volatile 和 st.volatile 会被生成。
3.29. 并发操作的内存模型
不适用。NVVM IR程序中的线程必须使用原子操作或屏障同步来进行通信。
3.30. 原子内存顺序约束
不支持原子加载和存储操作。除32位或64位操作数外，其他位宽操作数的原子操作也不支持。
3.31. 快速数学运算标志
支持。
3.32. 使用列表顺序指令
不支持。
4. 类型系统
除以下情况外均支持：
不支持浮点类型
fp128、x86_fp80和ppc_fp128。不支持
x86_mmx类型。不支持
token类型。不支持
non-integral pointer类型。
5. 常量
支持以下情况除外：
Token constants不支持。blockaddress(@function, %block)不支持。对于用作全局变量
@g1初始化器的常量表达式，如果该常量表达式包含全局标识符@g2，那么当该表达式可简化为bitcast+offset形式时（其中offset是一个整数，包括0），该常量表达式即受支持。
6. 其他数值
6.1. 内联汇编表达式
支持PTX指令的内联汇编器，具有以下支持的约束条件：
约束条件 |
类型 |
|---|---|
c |
i8 |
h |
i16 |
r |
i32 |
l |
i64 |
f |
f32 |
d |
f64 |
内联汇编元数据 !srcloc 会被接受但忽略。
不支持内联汇编方言 inteldialect。
7. 元数据
7.1. 元数据节点与元数据字符串
支持。
NVVM编译器能够识别以下元数据：
专用元数据节点
llvm.loop.unroll.countllvm.loop.unroll.disablellvm.loop.unroll.fullcallalign(参见调用约定的规则与限制)
支持并验证模块标志元数据（llvm.module.flags），但元数据值将被忽略。
所有其他元数据会被接受但忽略。
8. ThinLTO 摘要
不支持。
9. 内置全局变量
支持全局变量
llvm.used。支持全局变量
llvm.compiler.used不支持全局变量
llvm.global_ctors不支持全局变量
llvm.global_dtors
10. 使用说明
10.1. 终止指令
支持：
retbrswitchunreachable
不支持所有其他终止指令。
10.2. 二进制运算
支持：
addfaddsubfsubmulfmuludivsdivfdivuremsremfrem
10.3. 位运算二进制操作
支持：
shllshrashrandorxor
10.4. 向量运算
支持：
extractelementinsertelementshufflevector
10.5. 聚合操作
支持：
extractvalueinsertvalue
10.6. 内存访问与寻址操作
10.6.1. alloca指令
alloca指令返回一个指向本地地址空间的通用指针。不支持inalloca属性。支持的最大对齐度为2^23。仅当num为0时才支持addrspace(说明符。
10.6.2. 加载指令
load atomic 不支持。
10.6.3. 存储指令
store atomic 不支持。
10.6.4. fence指令
不支持。请改用NVVM内置函数。
10.6.5. cmpxchg指令
支持 i32、i64 和 i128 类型，但有以下限制：
指针必须是全局指针、共享指针，或指向全局地址空间或共享地址空间的通用指针。
接受并忽略
weak标记和failure ordering。i128类型仅在compute_90及以上版本中受支持。
10.6.6. atomicrmw 指令
仅支持以下操作：
xchgaddsubandorxormaxminumaxumin
所有其他操作均不支持。
这些操作支持i32和i64类型。此外，xchg操作在compute_90及以上版本中还支持i128。
指针操作数必须是全局指针、共享指针，或指向global地址空间或shared地址空间的通用指针。
10.6.7. getelementptr指令
支持。
10.7. 转换操作
支持：
trunc .. tozext .. tosext .. tofptrunc .. tofpext .. tofptoui .. tofptosi .. touitofp .. tositofp .. toptrtoint .. tointtoptr .. toaddrspacecast .. to-
bitcast .. to关于
bitcast的特殊用例，请参阅Conversion。
10.8. 其他操作
支持：
icmpfcmpphiselectva_argcall(其他规则和限制请参阅Calling Conventions)
所有其他操作均不支持。
11. 支持的内置函数
11.1. 支持的变长参数处理内置函数
llvm.va_startllvm.va_endllvm.va_copy
11.2. 支持的C/C++标准库内置函数
-
llvm.copysign这仅在现代NVVM IR方言中受支持。
-
llvm.memcpy请注意，常量地址空间不能作为目标地址使用，因为它是只读的。
-
llvm.memmove请注意，由于常量地址空间是只读的，因此不能使用它。
-
llvm.memset请注意，由于常量地址空间是只读的，因此无法使用。
-
llvm.sqrt支持float/double类型及其向量类型。映射到PTX指令
sqrt.rn.f32和sqrt.rn.f64。 -
llvm.fma支持float/double类型及其向量类型。映射到PTX指令
fma.rn.f32和fma.rn.f64
11.3. 支持的位操作内置函数
-
llvm.bitreverse支持
i8,i16,i32和i64类型。 -
llvm.bswap支持
i16,i32, 和i64类型。 -
llvm.ctpop支持
i8,i16,i32,i64以及这些类型的向量。 -
llvm.ctlz支持
i8,i16,i32,i64类型及其向量。 -
llvm.cttz支持
i8,i16,i32,i64类型及其向量类型。 -
llvm.fshl支持
i8,i16,i32, 和i64类型。 -
llvm.fshr支持
i8,i16,i32, 和i64类型。
11.4. 支持的专用算术内建函数
llvm.fmuladd
11.5. 支持带溢出检测的算术内置函数
支持 i16、i32 和 i64。
11.6. 支持的半精度浮点内置函数
llvm.convert.to.fp16llvm.convert.from.fp16
11.7. 支持的调试器内置函数
llvm.dbg.addrllvm.dbg.declarellvm.dbg.value
11.8. 支持的内存使用标记
llvm.lifetime.startllvm.lifetime.endllvm.invariant.startllvm.invariant.end
11.9. 支持的通用内置函数
-
llvm.var.annotation接受并忽略。
-
llvm.ptr.annotation接受并忽略。
-
llvm.annotation接受并忽略。
llvm.trapllvm.expectllvm.assumellvm.donothingllvm.sideeffect
12. 地址空间
12.1. 地址空间
NVVM IR预定义了一组内存地址空间，其语义与CUDA C/C++、OpenCL C和PTX中定义的类似。以下未列出的任何地址空间均不受支持。
名称 |
地址空间编号 |
语义/示例 |
|---|---|---|
代码 |
0 |
函数，代码
|
generic |
0 |
仅可用于限定指针所指向的对象
|
全局 |
1 |
|
共享 |
3 |
|
常量 |
4 |
|
本地 |
5 |
|
<保留> |
2, 101及以上 |
每个非内置的全局变量都可以声明驻留在特定的非零地址空间中，该地址空间只能是以下之一：global、shared 或 constant。
如果声明一个非内置全局变量时未指定地址空间编号或使用地址空间编号0，则该全局变量位于global地址空间中，且该全局变量的指针持有通用指针值。
预定义的NVVM内存空间是语言前端用来模拟源语言中内存空间所必需的。例如，
// CUDA C/C++
__constant__ int c;
__device__ int g;
; NVVM IR
@c = addrspace(4) global i32 0, align 4
@g = addrspace(1) global [2 x i32] zeroinitializer, align 4
地址空间编号2和101及以上保留仅供NVVM编译器内部使用。任何语言前端都不应直接生成使用这些地址空间的代码。
12.2. 通用指针与非通用指针
12.2.1. 通用指针与非通用指针
NVVM IR中存在通用指针和非通用指针。通用指针可以指向任何地址空间中的内存。非通用指针则指向特定地址空间中的内存。
在NVVM IR中，通用指针具有地址空间为generic的指针类型，而非通用指针则具有非通用地址空间的指针类型。
请注意，通用地址空间的地址空间编号为0——这是NVVM IR和LLVM IR中的默认值。代码地址空间的地址空间编号同样为0。函数指针通过地址空间code（addrspace(0)）进行限定。
支持通过通用指针进行加载/存储，也支持通过非通用指针进行加载/存储。不支持通过函数指针进行加载/存储
@a = addrspace(1) global i32 0, align 4 ; 'global' addrspace, @a holds a specific value
@b = global i32 0, align 4 ; 'global' addrspace, @b holds a generic value
@c = addrspace(4) global i32 0, align 4 ; 'constant' addrspace, @c holds a specific value
... = load i32 addrspace(1)* @a, align 4 ; Correct
... = load i32* @a, align 4 ; Wrong
... = load i32* @b, align 4 ; Correct
... = load i32 addrspace(1)* @b, align 4 ; Wrong
... = load i32 addrspace(4)* @c, align4 ; Correct
... = load i32* @c, align 4 ; Wrong
12.2.2. 转换
指向特定对象的通用指针的位值可能与指向同一对象的特定指针的位值不同。
应使用addrspacecast IR指令来执行跨地址空间的指针转换（从通用到非通用或从非通用到通用）。不支持将非通用指针转换为不同的非通用指针。如果通用指针未指向目标非通用地址空间中的对象，则从通用指针到非通用指针的转换是未定义的。
inttoptr 和 ptrtoint 是受支持的。当两个操作数大小相同时，inttoptr 和 ptrtoint 是保留值的指令。通常，使用 ptrtoint 和 inttoptr 来实现地址空间转换是未定义的行为。
以下内置函数可用于查询参数指针是否源自具有grid_constant属性的内核函数参数的地址：
i1 @llvm.nvvm.isspacep.grid_const(i8*)
以下内置函数可用于查询输入通用指针是否源自共享地址空间中分配的变量地址，该变量位于与调用线程所属父CTA同一集群的CTA中。此内置函数仅支持Hopper及以上架构。
i1 @llvm.nvvm.isspacep.cluster_shared(i8*)
以下内置函数可用于查询通用指针是否可以安全转换为特定的非通用地址空间：
i1 @llvm.nvvm.isspacep.const(i8*)i1 @llvm.nvvm.isspacep.global(i8*)i1 @llvm.nvvm.isspacep.local(i8*)i1 @llvm.nvvm.isspacep.shared(i8*)
bitcast 支持对指针进行操作，尽管LLVM IR禁止使用 bitcast 来改变指针的地址空间。
12.2.3. 两个不同特定地址空间之间无别名
两个不同的特定地址空间不会重叠。NVVM编译器假设通过指向不同地址空间的非通用指针进行的两次内存访问不存在别名问题。
12.3. alloca指令
alloca指令返回一个仅指向local地址空间的通用指针。
13. 全局属性注解
13.1. 概述
NVVM使用命名元数据(Named Metadata)来为IR对象添加属性标注，这些属性无法直接在IR中表示。NVVM IR生成器可以利用命名元数据为IR添加属性标注，这些标注可由NVVM编译器进行处理。
13.2. 属性表示
For each translation unit (that is, per bitcode file), there is a named metadata called nvvm.annotations.
该命名元数据包含一个MDNodes列表。
每个MDNode的第一个操作数是一个实体，该节点使用剩余的操作数对其进行注解。
多个MDNode可能为同一实体提供注解，在这种情况下，它们的第一个操作数将是相同的。
MDNode的剩余操作数按<属性名称, 值>的顺序组织。
属性名称操作数为MDString，而值为
i32。从带注释的实体后的操作数开始，每隔一个操作数指定一个属性。
-
属性后的操作数是其值。
以下是一个示例。
!nvvm.annotations = !{!12, !13} !12 = !{void (i32, i32)* @_Z6kernelii, !"kernel", i32 1} !13 = !{void ()* @_Z7kernel2v, !"kernel", i32 1, !"maxntidx", i32 16}
如果两个bitcode文件被链接且都包含名为nvvm.annotations的元数据，则链接后的文件将合并为一个单一的命名元数据。如果两个文件都为同一实体foo定义了属性，链接后的文件将包含两个为foo定义属性的MDNodes。若两个文件对同一实体存在冲突属性，则属于非法操作。
13.3. 支持的属性
属性名称 |
标注位置 |
描述 |
|---|---|---|
|
内核函数 |
从任何启动中预期的最大CTA大小。 |
|
内核函数 |
任何启动所需的最小预期CTA大小。 |
|
内核函数 |
支持Hopper+的集群维度配置。如果任一维度指定为0，则所有维度都必须指定为0。 |
|
内核函数 |
每个集群的最大块数。必须为非零值。仅支持Hopper及以上架构。 |
|
内核函数 |
给编译器/驱动程序的提示/指令，要求它在每个SM上至少放置这么多CTA。 |
|
内核函数 |
该参数是一个元数据节点，包含一个整数列表，其中每个整数n表示第n个参数具有grid_constant注解（从1开始编号）。参数类型必须是指针类型且设置了byval属性。对参数所指向的内存进行写入操作将导致未定义行为。此属性仅支持Volta及以上架构。 |
|
函数 |
函数的最大寄存器数量。 |
|
函数 |
表示此函数是一个内核函数。 |
|
函数 |
表示如果第n个参数类型的对齐方式不是自然对齐，则32位值的低16位中的值包含其对齐方式。n由该值的高16位指定。对于返回类型，n为0。 |
|
全局变量 |
表示该变量是一个纹理。 |
|
全局变量 |
表示该变量是一个表面。 |
|
全局变量 |
表示该变量是由UVM管理的变量。 |
14. 纹理与表面
14.1. 纹理变量与表面变量
纹理或表面变量可以声明/定义为全局变量，类型为i64，并在global地址空间中使用texture或surface注解。
纹理或表面变量必须有一个名称，该名称必须遵循标识符命名规范。
对纹理或表面变量的地址进行存储或加载操作是非法的。纹理或表面变量仅允许以下用途：
在元数据节点中
作为内置函数参数，如下所示
在
llvm.used全局变量中
14.2. 访问纹理内存或表面内存
纹理内存和表面内存可以通过纹理或表面句柄进行访问。NVVM提供以下内置函数，用于从纹理或表面变量获取纹理或表面句柄。
delcare i64 %llvm.nvvm.texsurf.handle.p1i64(metadata, i64 addrspace(1)*)
该内置函数的第一个参数是保存纹理或表面变量的元数据。此类元数据可能仅包含一个纹理或一个表面变量。该内置函数的第二个参数是纹理或表面变量本身。该函数返回一个i64类型的句柄。
从内部调用返回的句柄值可用作PTX内联汇编中的操作数（受限于l约束），以访问纹理或表面内存。
15. NVVM专用内置函数
15.1. 原子操作
除了原子指令外，还支持以下额外的原子内置函数。
declare float @llvm.nvvm.atomic.load.add.f32.p0f32(float* address, float val)
declare float @llvm.nvvm.atomic.load.add.f32.p1f32(float addrspace(1)* address, float val)
declare float @llvm.nvvm.atomic.load.add.f32.p3f32(float addrspace(3)* address, float val)
declare double @llvm.nvvm.atomic.load.add.f64.p0f64(double* address, double val)
declare double @llvm.nvvm.atomic.load.add.f64.p1f64(double addrspace(1)* address, double val)
declare double @llvm.nvvm.atomic.load.add.f64.p3f64(double addrspace(3)* address, double val)
读取位于地址address的单/双精度浮点值old，计算old+val，并将结果存储回同一内存地址。这些操作在一个原子事务中完成。该函数返回old。
declare i32 @llvm.nvvm.atomic.load.inc.32.p0i32(i32* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.inc.32.p1i32(i32 addrspace(1)* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.inc.32.p3i32(i32 addrspace(3)* address, i32 val)
读取位于地址address的32位字old，计算((old >= val) ? 0 : (old+1))，并将结果存储回同一内存地址。这三个操作在一个原子事务中完成。该函数返回old。
declare i32 @llvm.nvvm.atomic.load.dec.32.p0i32(i32* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.dec.32.p1i32(i32 addrspace(1)* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.dec.32.p3i32(i32 addrspace(3)* address, i32 val)
读取位于地址address的32位字old，计算(((old == 0) | (old > val)) ? val : (old-1) )，并将结果存储回同一内存地址。这三个操作在一个原子事务中完成。该函数返回old。
15.2. 屏障与内存栅栏
declare void @llvm.nvvm.barrier0()
等待线程块中的所有线程都到达此点，并且这些线程在llvm.nvvm.barrier0()之前对全局和共享内存的所有访问对该块中的所有线程可见。
declare i32 @llvm.nvvm.barrier0.popc(i32)
与llvm.nvvm.barrier0()功能相同，但额外增加了评估块内所有线程的谓词条件并返回谓词结果为非零的线程数量的特性。
declare i32 @llvm.nvvm.barrier0.and(i32)
与llvm.nvvm.barrier0()功能相同，但额外具备一项特性：它会评估块内所有线程的谓词条件，当且仅当所有线程的谓词评估结果均为非零值时，该函数才会返回非零值。
declare i32 @llvm.nvvm.barrier0.or(i32)
与llvm.nvvm.barrier0()功能相同，但额外具备以下特性：它会评估块内所有线程的谓词条件，当且仅当任一线程的谓词评估结果非零时返回非零值。
declare void @llvm.nvvm.cluster.barrier(i32 %flags)
在同一集群内的线程间同步和通信。此内建函数仅支持Hopper+架构。%flags根据以下表格进行编码：
%flags 位 |
含义 |
|---|---|
31-8 |
保留 |
7-4 |
内存排序（参见下文集群屏障内存排序编码） |
3-0 |
操作模式（参见下方的集群屏障操作模式编码） |
集群屏障操作模式编码
编码 |
模式 |
描述 |
|---|---|---|
0 |
到达 |
抵达集群屏障 |
1 |
等待 |
在集群屏障处等待 |
2-15 |
保留 |
保留 |
集群屏障内存排序编码
编码 |
模式 |
描述 |
|---|---|---|
0 |
Default |
在执行入口到达前请求的所有同步内存访问操作都会完成，并且在等待后对集群中的所有入口可见。 |
1 |
宽松模式 |
在执行入口到达前请求的所有先前隔离内存访问均已完成，并且在等待后对集群中的所有入口可见。此排序仅在操作模式为Arrive时支持。 |
2-15 |
保留 |
保留 |
declare void @llvm.nvvm.membar.cta()
这是一个线程块级别的内存屏障。该内置函数已弃用。请改用带有标志参数的nvvm.membar。
declare void @llvm.nvvm.membar.gl()
这是一个设备级的内存屏障。该内置函数已弃用。请改用带有标志参数的nvvm.membar。
declare void @llvm.nvvm.membar.sys()
这是一个系统级的内存屏障。该内置函数已弃用。请改用带有标志参数的nvvm.membar。
declare void @llvm.nvvm.membar(i32 %flags)
等待该线程请求的所有先前内存访问在以下membar模式定义的内存屏障级别执行。内存屏障仅强制垂直排序。它不保证与其他线程的执行同步。对于水平同步，应改用屏障，或与membar结合使用。
%flags 根据以下表格进行编码：
%flags 位 |
含义 |
|---|---|
31-4 |
保留 |
3-0 |
内存屏障模式（参见内存屏障模式编码。） |
内存屏障模式编码
编码方式 |
模式 |
描述 |
|---|---|---|
0 |
GLOBAL |
全局级别的内存屏障 |
1 |
CTA |
CTA级别的内存屏障 |
2 |
SYSTEM |
系统级内存屏障 |
3 |
保留 |
保留 |
4 |
CLUSTER |
在集群级别的内存屏障，仅适用于Hopper及以上架构 |
5-15 |
保留 |
保留 |
15.3. 地址空间转换
注意
注意：请使用addrspacecast IR指令进行地址空间转换。
15.4. 特殊寄存器
提供以下内置函数以支持读取特殊PTX寄存器：
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()
15.5. 纹理/表面访问
提供以下内置函数用于将全局纹理/表面变量转换为纹理/表面句柄。
declare i64 %llvm.nvvm.texsurf.handle.p1i64(metadata, i64 addrspace(1)*)
详情请参阅访问纹理内存或表面内存。
以下IR定义适用于本节中的所有内置函数：
type %float4 = { float, float, float, float }
type %long2 = { i64, i64 }
type %int4 = { i32, i32, i32, i32 }
type %int2 = { i32, i32 }
type %short4 = { i16, i16, i16, i16 }
type %short2 = { i16, i16 }
15.5.1. 纹理读取
对一维纹理进行采样：
%float4 @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %tex, i32 %x)
%float4 @llvm.nvvm.tex.unified.1d.v4f32.f32(i64 %tex, float %x)
%float4 @llvm.nvvm.tex.unified.1d.level.v4f32.f32(i64 %tex, float %x,
float %level)
%float4 @llvm.nvvm.tex.unified.1d.grad.v4f32.f32(i64 %tex, float %x,
float %dPdx,
float %dPdy)
%int4 @llvm.nvvm.tex.unified.1d.v4s32.s32(i64 %tex, i32 %x)
%int4 @llvm.nvvm.tex.unified.1d.v4s32.f32(i64 %tex, float %x)
%int4 @llvm.nvvm.tex.unified.1d.level.v4s32.f32(i64 %tex, float %x,
float %level)
%int4 @llvm.nvvm.tex.unified.1d.grad.v4s32.f32(i64 %tex, float %x,
float %dPdx,
float %dPdy)
%int4 @llvm.nvvm.tex.unified.1d.v4u32.s32(i64 %tex, i32 %x)
%int4 @llvm.nvvm.tex.unified.1d.v4u32.f32(i64 %tex, float %x)
%int4 @llvm.nvvm.tex.unified.1d.level.v4u32.f32(i64 %tex, float %x,
float %level)
%int4 @llvm.nvvm.tex.unified.1d.grad.v4u32.f32(i64 %tex, float %x,
float %dPdx,
float %dPdy)
对一维纹理数组进行采样：
%float4 @llvm.nvvm.tex.unified.1d.array.v4f32.s32(i64 %tex, i32 %idx, i32 %x)
%float4 @llvm.nvvm.tex.unified.1d.array.v4f32.f32(i64 %tex, i32 %idx, float %x)
%float4 @llvm.nvvm.tex.unified.1d.array.level.v4f32.f32(i64 %tex, i32 %idx,
float %x,
float %level)
%float4 @llvm.nvvm.tex.unified.1d.array.grad.v4f32.f32(i64 %tex, i32 %idx,
float %x,
float %dPdx,
float %dPdy)
%int4 @llvm.nvvm.tex.unified.1d.array.v4s32.s32(i64 %tex, i32 %idx, i32 %x)
%int4 @llvm.nvvm.tex.unified.1d.array.v4s32.f32(i64 %tex, i32 %idx, float %x)
%int4 @llvm.nvvm.tex.unified.1d.array.level.v4s32.f32(i64 %tex, i32 %idx,
float %x,
float %level)
%int4 @llvm.nvvm.tex.unified.1d.array.grad.v4s32.f32(i64 %tex, i32 %idx,
float %x,
float %dPdx,
float %dPdy)
%int4 @llvm.nvvm.tex.unified.1d.array.v4u32.s32(i64 %tex, i32 %idx, i32 %x)
%int4 @llvm.nvvm.tex.unified.1d.array.v4u32.f32(i64 %tex, i32 %idx, float %x)
%int4 @llvm.nvvm.tex.unified.1d.array.level.v4u32.f32(i64 %tex, i32 %idx,
float %x,
float %level)
%int4 @llvm.nvvm.tex.unified.1d.array.grad.v4u32.f32(i64 %tex, i32 %idx,
float %x,
float %dPdx,
float %dPdy)
对2D纹理进行采样:
%float4 @llvm.nvvm.tex.unified.2d.v4f32.s32(i64 %tex, i32 %x, i32 %y)
%float4 @llvm.nvvm.tex.unified.2d.v4f32.f32(i64 %tex, float %x, float %y)
%float4 @llvm.nvvm.tex.unified.2d.level.v4f32.f32(i64 %tex, float %x, float %y,
float %level)
%float4 @llvm.nvvm.tex.unified.2d.grad.v4f32.f32(i64 %tex, float %x, float %y,
float %dPdx_x, float %dPdx_y,
float %dPdy_x, float %dPdy_y)
%int4 @llvm.nvvm.tex.unified.2d.v4s32.s32(i64 %tex, i32 %x, i32 %y)
%int4 @llvm.nvvm.tex.unified.2d.v4s32.f32(i64 %tex, float %x, float %y,)
%int4 @llvm.nvvm.tex.unified.2d.level.v4s32.f32(i64 %tex, float %x, float %y,
float %level)
%int4 @llvm.nvvm.tex.unified.2d.grad.v4s32.f32(i64 %tex, float %x, float %y,
float %dPdx_x, float %dPdx_y,
float %dPdy_x, float %dPdy_y)
%int4 @llvm.nvvm.tex.unified.2d.v4u32.s32(i64 %tex, i32 %x i32 %y)
%int4 @llvm.nvvm.tex.unified.2d.v4u32.f32(i64 %tex, float %x float %y)
%int4 @llvm.nvvm.tex.unified.2d.level.v4u32.f32(i64 %tex, float %x, float %y,
float %level)
%int4 @llvm.nvvm.tex.unified.2d.grad.v4u32.f32(i64 %tex, float %x, float %y,
float %dPdx_x, float %dPdx_y,
float %dPdy_x, float %dPdy_y)
对2D纹理数组进行采样：
%float4 @llvm.nvvm.tex.unified.2d.array.v4f32.s32(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%float4 @llvm.nvvm.tex.unified.2d.array.v4f32.f32(i64 %tex, i32 %idx,
float %x, float %y)
%float4 @llvm.nvvm.tex.unified.2d.array.level.v4f32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %level)
%float4 @llvm.nvvm.tex.unified.2d.array.grad.v4f32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %dPdx_x,
float %dPdx_y,
float %dPdy_x,
float %dPdy_y)
%int4 @llvm.nvvm.tex.unified.2d.array.v4s32.s32(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%int4 @llvm.nvvm.tex.unified.2d.array.v4s32.f32(i64 %tex, i32 %idx,
float %x, float %y)
%int4 @llvm.nvvm.tex.unified.2d.array.level.v4s32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %level)
%int4 @llvm.nvvm.tex.unified.2d.array.grad.v4s32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %dPdx_x,
float %dPdx_y,
float %dPdy_x,
float %dPdy_y)
%int4 @llvm.nvvm.tex.unified.2d.array.v4u32.s32(i64 %tex, i32 %idx,
i32 %x i32 %y)
%int4 @llvm.nvvm.tex.unified.2d.array.v4u32.f32(i64 %tex, i32 %idx,
float %x float %y)
%int4 @llvm.nvvm.tex.unified.2d.array.level.v4u32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %level)
%int4 @llvm.nvvm.tex.unified.2d.array.grad.v4u32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %dPdx_x,
float %dPdx_y,
float %dPdy_x,
float %dPdy_y)
对3D纹理进行采样：
%float4 @llvm.nvvm.tex.unified.3d.v4f32.s32(i64 %tex, i32 %x, i32 %y, i32 %z)
%float4 @llvm.nvvm.tex.unified.3d.v4f32.f32(i64 %tex, float %x, float %y,
float %z)
%float4 @llvm.nvvm.tex.unified.3d.level.v4f32.f32(i64 %tex,float %x, float %y,
float %z, float %level)
%float4 @llvm.nvvm.tex.unified.3d.grad.v4f32.f32(i64 %tex, float %x, float %y,
float %z, float %dPdx_x,
float %dPdx_y, float %dPdx_z,
float %dPdy_x, float %dPdy_y,
float %dPdy_z)
%int4 @llvm.nvvm.tex.unified.3d.v4s32.s32(i64 %tex, i32 %x, i32 %y, i32 %z)
%int4 @llvm.nvvm.tex.unified.3d.v4s32.f32(i64 %tex, float %x, float %y,
float %z)
%int4 @llvm.nvvm.tex.unified.3d.level.v4s32.f32(i64 %tex, float %x, float %y,
float %z, float %level)
%int4 @llvm.nvvm.tex.unified.3d.grad.v4s32.f32(i64 %tex, float %x, float %y,
float %z, float %dPdx_x,
float %dPdx_y, float %dPdx_z,
float %dPdy_x, float %dPdy_y,
float %dPdy_z)
%int4 @llvm.nvvm.tex.unified.3d.v4u32.s32(i64 %tex, i32 %x i32 %y, i32 %z)
%int4 @llvm.nvvm.tex.unified.3d.v4u32.f32(i64 %tex, float %x, float %y,
float %z)
%int4 @llvm.nvvm.tex.unified.3d.level.v4u32.f32(i64 %tex, float %x, float %y,
float %z, float %level)
%int4 @llvm.nvvm.tex.unified.3d.grad.v4u32.f32(i64 %tex, float %x, float %y,
float %z, float %dPdx_x,
float %dPdx_y, float %dPdx_z,
float %dPdy_x, float %dPdy_y,
float %dPdy_z)
对立方体纹理进行采样：
%float4 @llvm.nvvm.tex.unified.cube.v4f32.f32(i64 %tex, float %x, float %y,
float %z)
%float4 @llvm.nvvm.tex.unified.cube.level.v4f32.f32(i64 %tex,float %x, float %y,
float %z, float %level)
%int4 @llvm.nvvm.tex.unified.cube.v4s32.f32(i64 %tex, float %x, float %y,
float %z)
%int4 @llvm.nvvm.tex.unified.cube.level.v4s32.f32(i64 %tex, float %x, float %y,
float %z, float %level)
%int4 @llvm.nvvm.tex.unified.cube.v4u32.f32(i64 %tex, float %x, float %y,
float %z)
%int4 @llvm.nvvm.tex.unified.cube.level.v4u32.f32(i64 %tex, float %x, float %y,
float %z, float %level)
对立方体纹理数组进行采样：
%float4 @llvm.nvvm.tex.unified.cube.array.v4f32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %z)
%float4 @llvm.nvvm.tex.unified.cube.array.level.v4f32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %z,
float %level)
%int4 @llvm.nvvm.tex.unified.cube.array.v4s32.f32(i64 %tex, i32 %idx, float %x,
float %y, float %z)
%int4 @llvm.nvvm.tex.unified.cube.array.level.v4s32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %z, float %level)
%int4 @llvm.nvvm.tex.unified.cube.array.v4u32.f32(i64 %tex, i32 %idx, float %x,
float %y, float %z)
%int4 @llvm.nvvm.tex.unified.cube.array.level.v4u32.f32(i64 %tex, i32 %idx,
float %x, float %y,
float %z, float %level)
获取一个四纹素的双线性插值足迹:
%float4 @llvm.nvvm.tld4.unified.r.2d.v4f32.f32(i64 %tex, float %x, float %y)
%float4 @llvm.nvvm.tld4.unified.g.2d.v4f32.f32(i64 %tex, float %x, float %y)
%float4 @llvm.nvvm.tld4.unified.b.2d.v4f32.f32(i64 %tex, float %x, float %y)
%float4 @llvm.nvvm.tld4.unified.a.2d.v4f32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.r.2d.v4s32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.g.2d.v4s32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.b.2d.v4s32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.a.2d.v4s32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.r.2d.v4u32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.g.2d.v4u32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.b.2d.v4u32.f32(i64 %tex, float %x, float %y)
%int4 @llvm.nvvm.tld4.unified.a.2d.v4u32.f32(i64 %tex, float %x, float %y)
15.5.2. 表面载荷
在以下内置函数中，clamp、trap或zero。
对于处理8位数据通道的表面加载指令，输出操作数的类型为i16。高八位是未定义的。
读取一维表面：
i16 @llvm.nvvm.suld.1d.i8.<clamp>(i64 %tex, i32 %x)
i16 @llvm.nvvm.suld.1d.i16.<clamp>(i64 %tex, i32 %x)
i32 @llvm.nvvm.suld.1d.i32.<clamp>(i64 %tex, i32 %x)
i64 @llvm.nvvm.suld.1d.i64.<clamp>(i64 %tex, i32 %x)
%short2 @llvm.nvvm.suld.1d.v2i8.<clamp>(i64 %tex, i32 %x)
%short2 @llvm.nvvm.suld.1d.v2i16.<clamp>(i64 %tex, i32 %x)
%int2 @llvm.nvvm.suld.1d.v2i32.<clamp>(i64 %tex, i32 %x)
%long2 @llvm.nvvm.suld.1d.v2i64.<clamp>(i64 %tex, i32 %x)
%short4 @llvm.nvvm.suld.1d.v4i8.<clamp>(i64 %tex, i32 %x)
%short4 @llvm.nvvm.suld.1d.v4i16.<clamp>(i64 %tex, i32 %x)
%int4 @llvm.nvvm.suld.1d.v4i32.<clamp>(i64 %tex, i32 %x)
读取一维表面数组：
i16 @llvm.nvvm.suld.1d.array.i8.<clamp>(i64 %tex, i32 %idx, i32 %x)
i16 @llvm.nvvm.suld.1d.array.i16.<clamp>(i64 %tex, i32 %idx, i32 %x)
i32 @llvm.nvvm.suld.1d.array.i32.<clamp>(i64 %tex, i32 %idx, i32 %x)
i64 @llvm.nvvm.suld.1d.array.i64.<clamp>(i64 %tex, i32 %idx, i32 %x)
%short2 @llvm.nvvm.suld.1d.array.v2i8.<clamp>(i64 %tex, i32 %idx, i32 %x)
%short2 @llvm.nvvm.suld.1d.array.v2i16.<clamp>(i64 %tex, i32 %idx, i32 %x)
%int2 @llvm.nvvm.suld.1d.array.v2i32.<clamp>(i64 %tex, i32 %idx, i32 %x)
%long2 @llvm.nvvm.suld.1d.array.v2i64.<clamp>(i64 %tex, i32 %idx, i32 %x)
%short4 @llvm.nvvm.suld.1d.array.v4i8.<clamp>(i64 %tex, i32 %idx, i32 %x)
%short4 @llvm.nvvm.suld.1d.array.v4i16.<clamp>(i64 %tex, i32 %idx, i32 %x)
%int4 @llvm.nvvm.suld.1d.array.v4i32.<clamp>(i64 %tex, i32 %idx, i32 %x)
读取二维表面：
i16 @llvm.nvvm.suld.2d.i8.<clamp>(i64 %tex, i32 %x, i32 %y)
i16 @llvm.nvvm.suld.2d.i16.<clamp>(i64 %tex, i32 %x, i32 %y)
i32 @llvm.nvvm.suld.2d.i32.<clamp>(i64 %tex, i32 %x, i32 %y)
i64 @llvm.nvvm.suld.2d.i64.<clamp>(i64 %tex, i32 %x, i32 %y)
%short2 @llvm.nvvm.suld.2d.v2i8.<clamp>(i64 %tex, i32 %x, i32 %y)
%short2 @llvm.nvvm.suld.2d.v2i16.<clamp>(i64 %tex, i32 %x, i32 %y)
%int2 @llvm.nvvm.suld.2d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %y)
%long2 @llvm.nvvm.suld.2d.v2i64.<clamp>(i64 %tex, i32 %x, i32 %y)
%short4 @llvm.nvvm.suld.2d.v4i8.<clamp>(i64 %tex, i32 %x, i32 %y)
%short4 @llvm.nvvm.suld.2d.v4i16.<clamp>(i64 %tex, i32 %x, i32 %y)
%int4 @llvm.nvvm.suld.2d.v4i32.<clamp>(i64 %tex, i32 %x, i32 %y)
读取二维表面数组：
i16 @llvm.nvvm.suld.2d.array.i8.<clamp>(i64 %tex, i32 %idx, i32 %x, i32 %y)
i16 @llvm.nvvm.suld.2d.array.i16.<clamp>(i64 %tex, i32 %idx, i32 %x, i32 %y)
i32 @llvm.nvvm.suld.2d.array.i32.<clamp>(i64 %tex, i32 %idx, i32 %x, i32 %y)
i64 @llvm.nvvm.suld.2d.array.i64.<clamp>(i64 %tex, i32 %idx, i32 %x, i32 %y)
%short2 @llvm.nvvm.suld.2d.array.v2i8.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%short2 @llvm.nvvm.suld.2d.array.v2i16.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%int2 @llvm.nvvm.suld.2d.array.v2i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%long2 @llvm.nvvm.suld.2d.array.v2i64.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%short4 @llvm.nvvm.suld.2d.array.v4i8.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%short4 @llvm.nvvm.suld.2d.array.v4i16.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y)
%int4 @llvm.nvvm.suld.2d.array.v4i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y)
读取3D表面：
i16 @llvm.nvvm.suld.3d.i8.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
i16 @llvm.nvvm.suld.3d.i16.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
i32 @llvm.nvvm.suld.3d.i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
i64 @llvm.nvvm.suld.3d.i64.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
%short2 @llvm.nvvm.suld.3d.v2i8.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
%short2 @llvm.nvvm.suld.3d.v2i16.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
%int2 @llvm.nvvm.suld.3d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
%long2 @llvm.nvvm.suld.3d.v2i64.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z)
%short4 @llvm.nvvm.suld.3d.v4i8.<clamp>(i64 %tex, i32 %x, i32 %y,
i32 %z)
%short4 @llvm.nvvm.suld.3d.v4i16.<clamp>(i64 %tex, i32 %x, i32 %y,
i32 %z)
%int4 @llvm.nvvm.suld.3d.v4i32.<clamp>(i64 %tex, i32 %x, i32 %y,
i32 %z)
15.5.3. 表面存储
在以下内置函数中，trap，而对于非格式化存储可以是以下之一：clamp、trap或zero。
对于处理8位数据通道的表面存储指令，输入操作数的类型为i16。高阶八位将被忽略。
编写一维表面：
;; Unformatted
void @llvm.nvvm.sust.b.1d.i8.<clamp>(i64 %tex, i32 %x, i16 %r)
void @llvm.nvvm.sust.b.1d.i16.<clamp>(i64 %tex, i32 %x, i16 %r)
void @llvm.nvvm.sust.b.1d.i32.<clamp>(i64 %tex, i32 %x, i32 %r)
void @llvm.nvvm.sust.b.1d.i64.<clamp>(i64 %tex, i32 %x, i64 %r)
void @llvm.nvvm.sust.b.1d.v2i8.<clamp>(i64 %tex, i32 %x, i16 %r, i16 %g)
void @llvm.nvvm.sust.b.1d.v2i16.<clamp>(i64 %tex, i32 %x, i16 %r, i16 %g)
void @llvm.nvvm.sust.b.1d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %r, i32 %g)
void @llvm.nvvm.sust.b.1d.v2i64.<clamp>(i64 %tex, i32 %x, i64 %r, i64 %g)
void @llvm.nvvm.sust.b.1d.v4i8.<clamp>(i64 %tex, i32 %x,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.1d.v4i16.<clamp>(i64 %tex, i32 %x,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.1d.v4i32.<clamp>(i64 %tex, i32 %x,
i32 %r, i32 %g, i32 %b, i32 %a)
;; Formatted
void @llvm.nvvm.sust.p.1d.i32.<clamp>(i64 %tex, i32 %x, i32 %r)
void @llvm.nvvm.sust.p.1d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %r, i32 %g)
void @llvm.nvvm.sust.p.1d.v4i32.<clamp>(i64 %tex, i32 %x,
i32 %r, i32 %g, i32 %b, i32 %a)
编写一维表面数组：
;; Unformatted
void @llvm.nvvm.sust.b.1d.array.i8.<clamp>(i64 %tex, i32 %idx, i32 %x,
i16 %r)
void @llvm.nvvm.sust.b.1d.array.i16.<clamp>(i64 %tex, i32 %idx, i32 %x,
i16 %r)
void @llvm.nvvm.sust.b.1d.array.i32.<clamp>(i64 %tex, i32 %idx, i32 %x,
i32 %r)
void @llvm.nvvm.sust.b.1d.array.i64.<clamp>(i64 %tex, i32 %idx, i32 %x,
i64 %r)
void @llvm.nvvm.sust.b.1d.array.v2i8.<clamp>(i64 %tex, i32 %idx, i32 %x,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.1d.array.v2i16.<clamp>(i64 %tex, i32 %idx, i32 %x,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.1d.array.v2i32.<clamp>(i64 %tex, i32 %idx, i32 %x,
i32 %r, i32 %g)
void @llvm.nvvm.sust.b.1d.array.v2i64.<clamp>(i64 %tex, i32 %idx, i32 %x,
i64 %r, i64 %g)
void @llvm.nvvm.sust.b.1d.array.v4i8.<clamp>(i64 %tex, i32 %idx, i32 %x,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.1d.array.v4i16.<clamp>(i64 %tex, i32 %idx, i32 %x,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.1d.array.v4i32.<clamp>(i64 %tex, i32 %idx, i32 %x,
i32 %r, i32 %g, i32 %b, i32 %a)
;; Formatted
void @llvm.nvvm.sust.p.1d.array.i32.<clamp>(i64 %tex, i32 %idx, i32 %x,
i32 %r)
void @llvm.nvvm.sust.p.1d.array.v2i32.<clamp>(i64 %tex, i32 %idx, i32 %x,
i32 %r, i32 %g)
void @llvm.nvvm.sust.p.1d.array.v4i32.<clamp>(i64 %tex, i32 %idx, i32 %x,
i32 %r, i32 %g, i32 %b, i32 %a)
编写2D曲面：
;; Unformatted
void @llvm.nvvm.sust.b.2d.i8.<clamp>(i64 %tex, i32 %x, i32 %y, i16 %r)
void @llvm.nvvm.sust.b.2d.i16.<clamp>(i64 %tex, i32 %x, i32 %y, i16 %r)
void @llvm.nvvm.sust.b.2d.i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %r)
void @llvm.nvvm.sust.b.2d.i64.<clamp>(i64 %tex, i32 %x, i32 %y, i64 %r)
void @llvm.nvvm.sust.b.2d.v2i8.<clamp>(i64 %tex, i32 %x, i32 %y,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.2d.v2i16.<clamp>(i64 %tex, i32 %x, i32 %y,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.2d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %y,
i32 %r, i32 %g)
void @llvm.nvvm.sust.b.2d.v2i64.<clamp>(i64 %tex, i32 %x, i32 %y,
i64 %r, i64 %g)
void @llvm.nvvm.sust.b.2d.v4i8.<clamp>(i64 %tex, i32 %x, i32 %y,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.2d.v4i16.<clamp>(i64 %tex, i32 %x, i32 %y,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.2d.v4i32.<clamp>(i64 %tex, i32 %x, i32 %y,
i32 %r, i32 %g, i32 %b, i32 %a)
;; Formatted
void @llvm.nvvm.sust.p.2d.i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %r)
void @llvm.nvvm.sust.p.2d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %y,
i32 %r, i32 %g)
void @llvm.nvvm.sust.p.2d.v4i32.<clamp>(i64 %tex, i32 %x, i32 %y,
i32 %r, i32 %g, i32 %b, i32 %a)
写入二维表面数组：
;; Unformatted
void @llvm.nvvm.sust.b.2d.array.i8.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y, i16 %r)
void @llvm.nvvm.sust.b.2d.array.i16.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y, i16 %r)
void @llvm.nvvm.sust.b.2d.array.i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y, i32 %r)
void @llvm.nvvm.sust.b.2d.array.i64.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y, i64 %r)
void @llvm.nvvm.sust.b.2d.array.v2i8.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.2d.array.v2i16.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.2d.array.v2i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i32 %r, i32 %g)
void @llvm.nvvm.sust.b.2d.array.v2i64.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i64 %r, i64 %g)
void @llvm.nvvm.sust.b.2d.array.v4i8.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.2d.array.v4i16.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.2d.array.v4i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i32 %r, i32 %g, i32 %b, i32 %a)
;; Formatted
void @llvm.nvvm.sust.p.2d.array.i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y, i32 %r)
void @llvm.nvvm.sust.p.2d.array.v2i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i32 %r, i32 %g)
void @llvm.nvvm.sust.p.2d.array.v4i32.<clamp>(i64 %tex, i32 %idx,
i32 %x, i32 %y,
i32 %r, i32 %g, i32 %b, i32 %a)
编写3D曲面：
;; Unformatted
void @llvm.nvvm.sust.b.3d.i8.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z, i16 %r)
void @llvm.nvvm.sust.b.3d.i16.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z, i16 %r)
void @llvm.nvvm.sust.b.3d.i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z, i32 %r)
void @llvm.nvvm.sust.b.3d.i64.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z, i64 %r)
void @llvm.nvvm.sust.b.3d.v2i8.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.3d.v2i16.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i16 %r, i16 %g)
void @llvm.nvvm.sust.b.3d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i32 %r, i32 %g)
void @llvm.nvvm.sust.b.3d.v2i64.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i64 %r, i64 %g)
void @llvm.nvvm.sust.b.3d.v4i8.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.3d.v4i16.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i16 %r, i16 %g, i16 %b, i16 %a)
void @llvm.nvvm.sust.b.3d.v4i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i32 %r, i32 %g, i32 %b, i32 %a)
;; Formatted
void @llvm.nvvm.sust.p.3d.i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z, i32 %r)
void @llvm.nvvm.sust.p.3d.v2i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i32 %r, i32 %g)
void @llvm.nvvm.sust.p.3d.v4i32.<clamp>(i64 %tex, i32 %x, i32 %y, i32 %z,
i32 %r, i32 %g, i32 %b, i32 %a)
15.6. Warp级操作
15.6.1. 屏障同步
以下内置函数在warp内的线程子集之间执行屏障同步。
declare void @llvm.nvvm.bar.warp.sync(i32 %membermask)
该内置函数会使执行线程等待，直到与%membermask对应的所有线程都执行了具有相同%membermask值的相同内置函数后，才会恢复执行。
参数 %membership 是一个32位掩码，每个位对应warp中的一个线程通道。值为1表示该线程属于该子集。
如果执行线程不在%membermask中，则该内置函数的行为是未定义的。
对于compute_62或更低版本，%membermask中的所有线程必须在收敛时调用相同的@llvm.nvvm.bar.warp.sync()，并且只有属于%membermask的线程可以在调用该内部函数时处于活动状态。否则，行为将是未定义的。
15.6.2. 数据移动
以下内置函数同步线程束中的一部分线程，然后在这些线程之间执行数据移动。
declare {i32, i1} @llvm.nvvm.shfl.sync.i32(i32 %membermask, i32 %mode, i32 %a, i32 %b, i32 %c)
该内置函数使执行线程等待，直到与%membermask对应的所有线程都使用相同的%membermask值执行了相同的内置函数后，才会从同一warp中的其他线程读取数据。
参数 %membership 是一个32位掩码，每个位对应warp中的一个线程通道。值为1表示该线程属于该子集。
当前执行warp中的每个线程将根据输入参数%b、%c和%mode计算源通道索引j。如果计算的源通道索引j在有效范围内，返回的i32值将是通道j的%a值；否则，将返回当前线程的%a值。如果与通道j对应的线程处于非活动状态，则返回的i32值未定义。如果源通道j在有效范围内，返回的i1值设为1，否则设为0。
参数 %mode 必须是一个常量，其编码在下表中指定。
编码 |
含义 |
|---|---|
0 |
IDX |
1 |
UP |
2 |
下线 |
3 |
BFLY |
参数 %b 根据 %mode 指定源车道或源车道偏移量。
参数 %c 包含两个打包值，分别用于指定将warp逻辑分割为子段的掩码，以及用于钳位源通道索引的上限值。
以下伪代码说明了该内置函数的语义。
wait until all threads in %membermask have arrived;
%lane[4:0] = current_lane_id; // position of thread in warp
%bval[4:0] = %b[4:0]; // source lane or lane offset (0..31)
%cval[4:0] = %c[4:0]; // clamp value
%mask[4:0] = %c[12:8];
%maxLane = (%lane[4:0] & %mask[4:0]) | (%cval[4:0] & ~%mask[4:0]);
%minLane = (%lane[4:0] & %mask[4:0]);
switch (%mode) {
case UP: %j = %lane - %bval; %pval = (%j >= %maxLane); break;
case DOWN: %j = %lane + %bval; %pval = (%j <= %maxLane); break;
case BFLY: %j = %lane ^ %bval; %pval = (%j <= %maxLane); break;
case IDX: %j = %minLane | (%bval[4:0] & ~%mask[4:0]); %pval = (%j <= %maxLane); break;
}
if (!%pval) %j = %lane; // copy from own lane
if (thread at lane %j is active)
%d = %a from lane %j
else
%d = undef
return {%d, %pval}
请注意，如果源通道的线程不在%membermask中，则返回值是未定义的。
如果执行线程不在%membermask中，则该内置函数的行为是未定义的。
对于compute_62或更低版本，%membermask中的所有线程必须在收敛时调用相同的@llvm.nvvm.shfl.sync.i32()，并且只有属于%membermask的线程可以在调用该内部函数时处于活动状态。否则，行为将是未定义的。
15.6.3. 投票
以下内置函数同步warp中的一部分线程，然后对该子集内所有线程的谓词执行归约-广播操作。
declare {i32, i1} @llvm.nvvm.vote.sync(i32 %membermask, i32 %mode, i1 %predicate)
该内置函数使执行线程等待，直到与%membermask对应的所有线程都执行了具有相同%membermask值的相同内置函数后，才会对该子集中所有线程的谓词执行归约-广播操作。
参数 %membermask 是一个32位掩码，每个位对应warp中的一个线程通道。1表示该线程属于子集。
@llvm.nvvm.vote.sync() 在同步后对%membermask中所有线程的源%predicate执行归约操作。返回值在%membermask中的所有线程间保持一致。返回聚合体中保存返回值的元素取决于%mode。
参数 %mode 必须是一个常量，其编码在下表中指定。
编码方式 |
含义 |
返回值 |
|---|---|---|
0 |
全部 |
|
1 |
任意 |
|
2 |
EQ |
|
3 |
投票 |
|
对于BALLOT模式，i32值表示投票数据，其中包含%membermask中每个线程的%predicate值，这些值位于与线程的land id对应的比特位位置。不在%membermask中的线程对应的比特值为0。
请注意，如果源通道中的线程不在%membermask中，则返回值未定义。
如果执行线程不在%membermask中，则该内置函数的行为是未定义的。
对于compute_62或更低版本，%membermask中的所有线程必须在收敛时调用相同的@llvm.nvvm.vote.sync()，并且只有属于%membermask的线程在调用该内部函数时才能处于活动状态。否则，行为将是未定义的。
15.6.4. 匹配
以下内置函数可同步warp中的部分线程，然后在子集线程间广播并比较某个值。
declare i32 @llvm.nvvm.match.any.sync.i32(i32 %membermask, i32 %value)
declare i32 @llvm.nvvm.match.any.sync.i64(i32 %membermask, i64 %value)
declare {i32, i1} @llvm.nvvm.match.all.sync.i32(i32 %membermask, i32 %value)
declare {i32, i1} @llvm.nvvm.match.all.sync.i64(i32 %membermask, i64 %value)
这些内置函数会使执行线程等待，直到与%membermask对应的所有线程都执行了具有相同%membermask值的相同内置函数后，才会在该子集中的所有线程之间对操作数%value进行广播和比较。
参数 %membership 是一个32位掩码，每个位对应warp中的一个线程通道。值为1表示该线程属于该子集。
i32返回值是一个32位掩码，其中掩码中的位位置对应于线程的laneid。
在any版本中，i32返回值被设置为%membermask中活动线程的掩码，这些线程的值与操作数%value相同。
在all版本中，如果%membermask中的所有活动线程与操作数%value具有相同的值，则i32返回值被设置为%membermask，且i1值被设置为1。否则，i32返回值被设置为0，且i1返回值也被设置为0。
如果执行线程不在%membermask中，则该内置函数的行为是未定义的。
这些内置函数仅在compute_70或更高版本中可用。
15.6.5. 矩阵运算
此为预览功能。未来版本可能会移除支持。
NVVM提供用于矩阵乘法运算的warp级内置函数。核心运算形式为矩阵乘加操作：
D = A*B + C, or
C = A*B + C
其中A是一个MxK矩阵，B是一个KxN矩阵，而C和D是MxN矩阵。C和D也被称为累加器。A和B矩阵的元素类型为16位浮点数。累加器的元素类型可以是32位浮点数或16位浮点数。
一个warp中的所有线程将共同持有矩阵A、B、C和D的内容。每个线程仅持有矩阵A的一个片段、矩阵B的一个片段、矩阵C的一个片段以及结果矩阵D的一个片段。矩阵元素如何在片段间分布对用户是不透明的，并且对于矩阵A、B和累加器是不同的。
片段由一系列元素值表示。对于fp32矩阵，元素类型为float。对于fp16矩阵，元素类型为i32（每个i32值包含两个fp16值）。元素数量随矩阵形状而变化。
15.6.5.1. 加载片段
以下内置函数会同步线程束中的所有线程，然后为每个线程加载矩阵的一个片段。
; load fragment A
declare {i32, i32, i32, i32, i32, i32, i32, i32} @llvm.nvvm.hmma.m16n16k16.ld.a.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {i32, i32, i32, i32, i32, i32, i32, i32} @llvm.nvvm.hmma.m32n8k16.ld.a.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {i32, i32, i32, i32, i32, i32, i32, i32} @llvm.nvvm.hmma.m8n32k16.ld.a.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
; load fragment B
declare {i32, i32, i32, i32, i32, i32, i32, i32} @llvm.nvvm.hmma.m16n16k16.ld.b.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {i32, i32, i32, i32, i32, i32, i32, i32} @llvm.nvvm.hmma.m32n8k16.ld.b.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {i32, i32, i32, i32, i32, i32, i32, i32} @llvm.nvvm.hmma.m8n32k16.ld.b.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
; load fragment C
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m16n16k16.ld.c.f32.p<n>f32(float addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m32n8k16.ld.c.f32.p<n>f32(float addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m8n32k16.ld.c.f32.p<n>f32(float addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
; load fragment C
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m16n16k16.ld.c.f16.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m32n8k16.ld.c.f16.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m8n32k16.ld.c.f16.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol);
这些内置函数从内存位置%ptr加载并返回一个矩阵片段。内存中的矩阵必须采用规范矩阵布局，其主维度为%ldm。%rowcol指定内存中的矩阵是行优先(0)还是列优先(1)。%rowcol必须是一个常量值。
返回的值序列表示调用线程持有的片段。矩阵元素如何分布在各个片段中对用户是不透明的，并且对于矩阵A、B和累加器是不同的。因此，提供了三种变体（即ld.a、ld.b和ld.c）。
这些内置函数根据地址空间进行了重载。地址空间编号
如果warp中的任何线程已退出，则此内置函数的行为是未定义的。
这些内置函数仅在compute_70或更高版本中可用。
15.6.5.2. 存储片段
以下内置函数会同步一个warp中的所有线程，然后为每个线程存储矩阵的一个片段。
; The last 8 arguments are the elements of the C fragment
declare void @llvm.nvvm.hmma.m16n16k16.st.c.f32.p<n>float(float addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol, float, float, float, float, float, float, float, float);
declare void @llvm.nvvm.hmma.m32n8k16.st.c.f32.p<n>float(float addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol, float, float, float, float, float, float, float, float);
declare void @llvm.nvvm.hmma.m8n32k16.st.c.f32.p<n>float(float addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol, float, float, float, float, float, float, float, float);
; The last 4 arguments are the elements of the C fragment
declare void @llvm.nvvm.hmma.m16n16k16.st.c.f16.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol, i32, i32, i32, i32);
declare void @llvm.nvvm.hmma.m32n8k16.st.c.f16.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol, i32, i32, i32, i32);
declare void @llvm.nvvm.hmma.m8n32k16.st.c.f16.p<n>i32(i32 addrspace(<n>)* %ptr, i32 %ldm, i32 %rowcol, i32, i32, i32, i32);
这些内置函数将累加器片段存储到内存位置%ptr。内存中的矩阵必须采用规范矩阵布局，其主维度为%ldm。%rowcol指定内存中的矩阵是行优先(0)还是列优先(1)。%rowcol必须是一个常量值。
这些内置函数根据地址空间进行了重载。地址空间编号
如果warp中的任何线程已退出，则此内置函数的行为是未定义的。
这些内置函数仅在compute_70或更高版本中可用。
15.6.5.3. 矩阵乘积累加运算
以下内置函数会同步warp中的所有线程，然后执行矩阵乘加操作。
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m16n16k16.mma.f16.f16(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, i32 %c0, i32 %c1, i32 %c2, i32 %c3);
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m32n8k16.mma.f16.f16(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, i32 %c0, i32 %c1, i32 %c2, i32 %c3);
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m8n32k16.mma.f16.f16(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, i32 %c0, i32 %c1, i32 %c2, i32 %c3);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m16n16k16.mma.f32.f16(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, i32 %c0, i32 %c1, i32 %c2, i32 %c3);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m32n8k16.mma.f32.f16(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, i32 %c0, i32 %c1, i32 %c2, i32 %c3);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m8n32k16.mma.f32.f16(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, i32 %c0, i32 %c1, i32 %c2, i32 %c3);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m16n16k16.mma.f32.f32(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, float %c0, float %c1, float %c2, float %c3, float %c4, float %c5, float %c6, float %c7);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m32n8k16.mma.f32.f32(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, float %c0, float %c1, float %c2, float %c3, float %c4, float %c5, float %c6, float %c7);
declare {float, float, float, float, float, float, float, float} @llvm.nvvm.hmma.m8n32k16.mma.f32.f32(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, float %c0, float %c1, float %c2, float %c3, float %c4, float %c5, float %c6, float %c7);
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m16n16k16.mma.f16.f32(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, float %c0, float %c1, float %c2, float %c3, float %c4, float %c5, float %c6, float %c7);
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m32n8k16.mma.f16.f32(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, float %c0, float %c1, float %c2, float %c3, float %c4, float %c5, float %c6, float %c7);
declare {i32, i32, i32, i32} @llvm.nvvm.hmma.m8n32k16.mma.f16.f32(i32 %rowcol, i32 %satf, i32 %a0, i32 %a1, i32 %a2, i32 %a3, i32 %a4, i32 %a5, i32 %a6, i32 %a7, i32 %b0, i32 %b1, i32 %b2, i32 %b3, i32 %b4, i32 %b5, i32 %b6, i32 %b7, float %c0, float %c1, float %c2, float %c3, float %c4, float %c5, float %c6, float %c7);
这些内部函数执行矩阵乘加运算。%rowcol指定了A和B片段的布局。它必须是一个常量值，可以具有以下值和语义。
编码 |
含义 |
|---|---|
0 |
片段A为行优先，片段B为行优先 |
1 |
A片段为行优先，B片段为列优先 |
2 |
A片段采用列主序，B片段采用行主序 |
3 |
片段A为列优先存储，片段B为列优先存储 |
对%satf的支持已被移除，该操作数必须为常量零。
如果warp中的任何线程已退出，这些内部函数的行为将是未定义的。
这些内置函数仅在compute_70或更高版本中可用。
16. 源代码级调试支持
为了支持IR模块的源代码级调试，NVVM IR提供了调试内置函数和调试信息描述符来表达调试信息。调试信息描述符使用专门的元数据节点表示。当前NVVM IR调试元数据版本为3.1。
NVVM IR调试支持基于LLVM 7.0.1（适用于Blackwell之前的架构目标）和LLVM 18.1.8（适用于Blackwell及后续架构目标）。关于IR的完整语义，本章读者应参考官方LLVM IR文档中的Specialized Metadata Nodes和Source Level Debugging部分。Blackwell及后续架构目标请分别参阅this和this文档。
当请求调试支持时，模块中需要存在以下元数据节点：
命名的元数据节点
!llvm.dbg.cu"Debug Info Version"标志的模块标志元数据：行为标志应为Error。该标志的值应为DEBUG_METADATA_VERSION，即 3。名为
!nvvmir.version的元数据，包含一个元数据节点，其中包含NVVM IR的主版本号和次版本号，后跟NVVM IR调试元数据的主版本号和次版本号。当前的NVVM IR调试元数据版本是3.1。-
调试分辨率（例如，完整、仅行信息）由DICompileUnit的
emissionKind字段控制：FullDebug (value: 1): 生成符号调试和行号信息。这需要在编译时指定libNVVM的-g选项。DebugDirectivesOnly (value: 3): 生成行信息。
源级调试仅支持单个调试编译单元。如果存在多个输入NVVM IR模块，最多只能有一个模块包含单个调试编译单元。
17. PTX的NVVM ABI
17.1. 链接类型
下表提供了与函数和全局变量关联的NVVM IR链接类型到PTX链接器指令的映射关系。
LLVM 链接类型 |
PTX 链接器指令 |
||
|---|---|---|---|
|
这是默认的链接类型，不需要链接器指令。 |
||
|
带定义的函数 |
|
|
带初始化的全局变量 |
|||
未定义的函数 |
|
||
未初始化的全局变量 |
|||
|
|
||
|
|
||
所有其他链接类型 |
不支持。 |
||
17.2. 参数传递与返回
下表展示了NVVM IR中函数参数和返回类型与PTX类型的映射关系。
源类型 |
位大小 |
PTX类型 |
|---|---|---|
整数类型 |
<= 32 |
|
64 |
|
|
指针类型（不包含 |
32 |
|
64 |
|
|
浮点类型 |
32 |
|
64 |
|
|
聚合类型 |
任意尺寸 |
其中align表示整体聚合或向量的字节对齐方式，name是与聚合或向量关联的变量名称，size表示聚合或向量的字节大小。 |
使用 |
32 或 64 |
|
向量类型 |
任意大小 |
18. 修订历史
版本 1.0
初始版本。
版本 1.1
在全局属性注解中新增了对UVM托管变量的支持。详见Supported Properties。
版本 1.2
为CUDA 7.0更新至LLVM 3.4版本。
移除地址空间内置函数，改用
addrspacecast。添加关于源代码级调试支持的信息。
版本 1.3
为CUDA 8.0添加对LLVM 3.8的支持。
版本 1.4
增加对warp级内置函数的支持。
版本 1.5
为CUDA 9.2添加对LLVM 5.0的支持。
版本 1.6
为CUDA 11.2更新至LLVM 7.0.1版本。
版本 1.7
添加对动态大小alloca的支持。
版本 1.8
在数据布局中添加对i128的支持。
版本 1.9
修改了关于忽略共享变量初始化的文本。
版本 1.10
为CUDA 11.7添加了对grid_constant内核参数的支持。
版本 1.11
为CUDA 11.8新增了对Hopper+集群内部函数和max_blocks_per_cluster内核属性的支持。
已弃用对32位编译的支持。
版本 2.0
将NVVM IR更新至2.0版本，该版本与1.x版本的NVVM IR不兼容
移除了地址空间转换内置函数。当这些内置函数存在于2.0版本的中间表示(IR)时，IR验证器会报错。建议libNVVM的用户改用addrspacecast指令。
对支持的数据布局进行更严格的错误检查。
不再支持旧式循环展开pragma元数据应用于循环回边。建议客户端使用LLVM框架定义的新循环pragma元数据。
不再支持使用非未定义值初始化共享变量。在1.x版本中，这些初始化器会被静默忽略。此特性导致2.0版本与1.x版本不兼容。
19. 通知
19.1. 通知
本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“NVIDIA”）对本文件所含信息的准确性或完整性不作任何明示或暗示的陈述或保证，并对其中可能存在的错误不承担任何责任。NVIDIA对于因使用此类信息而产生的后果、或因使用该信息导致的第三方专利或其他权利侵权概不负责。本文件不构成对开发、发布或交付任何材料（定义见下文）、代码或功能的承诺。
NVIDIA保留随时对本文件进行更正、修改、增强、改进以及任何其他变更的权利，恕不另行通知。
客户在下单前应获取最新的相关信息，并确认这些信息是最新且完整的。
除非NVIDIA与客户授权代表签署的单独销售协议中另有约定，否则NVIDIA产品的销售均以订单确认时提供的NVIDIA标准销售条款和条件为准（以下简称"销售条款"）。NVIDIA特此明确反对将任何客户通用条款适用于本文件所述NVIDIA产品的采购。本文件不直接或间接构成任何合同义务。
NVIDIA产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于那些可以合理预期NVIDIA产品故障或失灵会导致人身伤害、死亡、财产或环境损害的应用场景。NVIDIA对于在此类设备或应用中使用和/或包含NVIDIA产品不承担任何责任，因此客户需自行承担相关风险。
NVIDIA不声明或保证基于本文档的产品适用于任何特定用途。NVIDIA未必会对每个产品的所有参数进行测试。客户应全权负责评估和确定本文档所含信息的适用性，确保产品适合并满足客户计划的应用需求，并执行必要的应用测试以避免应用或产品出现故障。客户产品设计中的缺陷可能会影响NVIDIA产品的质量和可靠性，并可能导致超出本文档范围的其他或不同的条件和/或要求。对于任何因以下原因导致的故障、损坏、成本或问题，NVIDIA不承担任何责任：(i) 以违反本文档的任何方式使用NVIDIA产品或(ii) 客户产品设计。
本文档不授予任何NVIDIA专利权、版权或其他NVIDIA知识产权的明示或暗示许可。NVIDIA发布的关于第三方产品或服务的信息，不构成NVIDIA对这些产品或服务的使用许可或担保认可。使用此类信息可能需要获得第三方基于其专利或其他知识产权的许可，或需要获得NVIDIA基于其专利或其他知识产权的许可。
本文件中的信息仅可在获得NVIDIA事先书面批准、未经改动完整复制且完全符合所有适用的出口法律法规，并附带所有相关条件、限制和声明的情况下进行复制。
本文件及所有NVIDIA设计规格、参考板、文件、图纸、诊断工具、清单和其他文档（统称及单独称为"材料"）均以"现状"提供。NVIDIA不对材料作出任何明示或默示的保证，包括但不限于对不侵权、适销性和特定用途适用性的默示保证免责。在法律允许的最大范围内，NVIDIA不就因使用本文件导致的任何损害承担责任，包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论损害成因如何，也无论责任理论为何，即使NVIDIA已被告知发生此类损害的可能性。不论客户因任何原因可能遭受的任何损害，NVIDIA对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。
19.2. OpenCL
OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。
19.3. 商标
NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。