NVVM IR 规范
基于 LLVM IR 的 NVVM 编译器（中间表示）参考指南。
1. 简介
NVVM IR 是一种基于 LLVM IR 的编译器 IR（中间表示）。NVVM IR 旨在表示 GPU 计算内核（例如，CUDA 内核）。高级语言前端，如 CUDA C 编译器前端，可以生成 NVVM IR。NVVM 编译器（基于 LLVM）从 NVVM IR 生成 PTX 代码。
NVVM IR 和 NVVM 编译器在很大程度上与所使用的源语言无关。NVVM 编译器的 PTX 代码生成部分需要知道源语言，因为 DCI（驱动程序/编译器接口）存在差异。
NVVM IR 是一种二进制格式，基于 LLVM IR bitcode 格式的子集。本文档仅使用人类可读的形式来描述 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 版本（请参阅 源代码级调试支持）。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, alignment> 注解对齐方式，其中 alignment 是一个 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. 可见性样式
所有样式（default、hidden 和 protected）均被接受并忽略。
3.4. DLL 存储类
不支持。
3.5. 线程局部存储模型
不支持。
3.6. 运行时抢占说明符
不支持。
3.7. 结构类型
支持。
3.8. 非整数指针类型
不支持。
3.9. Comdats
不支持。
3.10. source_filename
接受并忽略。
3.11. 全局变量
全局变量（不是内建全局变量）可以选择声明为驻留在以下地址空间之一中
globalsharedconstant
如果未显式指定地址空间，则假定全局变量驻留在具有通用地址值的 global 地址空间中。有关详细信息，请参阅 地址空间。
不支持 thread_local 变量。
不允许显式节（元数据节除外）。
不支持 shared 变量的初始化。使用 undef 初始化。
3.12. 函数
以下函数不受支持
对齐
显式节
垃圾回收器名称
前缀数据
序言
人格化
3.13. 别名
仅支持作为非内核函数的别名。
3.14. Ifuncs
不支持。
3.15. 命名元数据
接受并忽略，以下情况除外
!nvvm.annotations：请参阅 全局属性注解!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(<num>) 说明符。
10.6.2. load 指令
不支持 load atomic。
10.6.3. store 指令
不支持 store atomic。
10.6.4. fence 指令
不支持。请改用 NVVM 内建函数。
10.6.5. cmpxchg 指令
支持 i32、i64 和 i128 类型，但有以下限制
指针必须是全局指针、共享指针或指向全局地址空间或共享地址空间的通用指针。
接受并忽略
weak标记和failure ordering。仅在
compute_90及更高版本上支持i128类型。
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的特殊用例，请参阅 转换。
10.8. 其他操作
支持
icmpfcmpphiselectva_argcall（有关其他规则和限制，请参阅 调用约定。）
不支持所有其他操作。
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 和 float/double 向量。映射到 PTX
sqrt.rn.f32和sqrt.rn.f64。 -
llvm.fma支持 float/double 和 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 中定义的语义。 未在下面列出的任何地址空间均不受支持。
名称 |
地址空间编号 |
语义/示例 |
|---|---|---|
code |
0 |
函数、代码
|
generic |
0 |
只能用于限定指针的被指向类型
|
global |
1 |
|
shared |
3 |
|
constant |
4 |
|
local |
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. 通用指针 vs. 非通用指针
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 相同的集群中。 此内联函数仅在 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 使用命名元数据来注解 IR 对象，这些对象的属性在 IR 中无法表示。 NVVM IR 生成器可以使用命名元数据来注解 IR 的属性，NVVM 编译器可以处理这些属性。
13.2. 属性的表示形式
对于每个转换单元（即，每个位代码文件），都有一个名为 nvvm.annotations 的命名元数据。
此命名元数据包含 MDNode 列表。
每个 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}
如果正在链接两个位代码文件，并且两者都具有命名元数据 nvvm.annotations，则链接的文件将具有单个合并的命名元数据。 如果两个文件都为同一实体 foo 定义了属性，则链接的文件将具有两个 MDNode，用于为 foo 定义属性。 文件为同一实体具有冲突的属性是非法的。
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 |
Arrive |
到达集群屏障 |
1 |
Wait |
在集群屏障处等待 |
2-15 |
保留 |
保留 |
集群屏障内存排序编码
编码 |
模式 |
描述 |
|---|---|---|
0 |
Default |
在到达之前，执行条目请求的所有同步内存访问都已执行，并且在等待之后对集群中的所有条目都可见。 |
1 |
Relaxed |
在到达之前，执行条目请求的所有先前已栅栏化的内存访问都已执行，并且在等待之后对集群中的所有条目都可见。 仅当操作模式为“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 级别执行。 内存屏障仅强制执行垂直排序。 它不保证与其他线程的执行同步。 对于水平同步，应使用屏障，或与 membar 结合使用。
%flags 根据下表进行编码
%flags 位 |
含义 |
|---|---|
31-4 |
保留 |
3-0 |
Membar 模式（请参阅 Membar 模式编码。） |
Membar 模式编码
编码 |
模式 |
描述 |
|---|---|---|
0 |
GLOBAL |
全局级别的 Membar |
1 |
CTA |
CTA 级别的 Membar |
2 |
SYSTEM |
系统级别的 Membar |
3 |
保留 |
保留 |
4 |
CLUSTER |
集群级别的 Membar，仅在 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. 纹理读取
采样 1D 纹理
%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)
采样 1D 纹理数组
%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> 表示表面钳位模式，可以是以下之一：clamp、 trap 或 zero。
对于对 8 位数据通道进行操作的表面加载指令，输出操作数的类型为 i16。 高阶八位是未定义的。
读取 1D 表面
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)
读取 1D 表面数组
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)
读取 2D 表面
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)
读取 2D 表面数组
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. 表面存储
在以下内联函数中，<clamp> 表示表面钳位模式。 对于格式化存储，它是 trap，对于非格式化存储，可以是以下之一：clamp、 trap 或 zero。
对于对 8 位数据通道进行操作的表面存储指令，输入操作数的类型为 i16。 高阶八位将被忽略。
写入 1D 表面
;; 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)
写入 1D 表面数组
;; 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)
写入 2D 表面数组
;; 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 中的一个 lane。 1 表示线程位于子集中。
如果执行线程不在 %membermask 中，则此内联函数的行为未定义。
对于 compute_62 或更低版本，%membermask 中的所有线程都必须在收敛中调用相同的 @llvm.nvvm.bar.warp.sync()，并且仅当调用内联函数时，属于 %membermask 的线程才能处于活动状态。 否则，行为未定义。
15.6.2. 数据移动
以下内联函数同步 warp 中的线程子集，然后在这些线程之间执行数据移动。
declare {i32, i1} @llvm.nvvm.shfl.sync.i32(i32 %membermask, i32 %mode, i32 %a, i32 %b, i32 %c)
此内联函数导致执行线程等待，直到与 %membermask 对应的所有线程都使用相同的 %membermask 值执行相同的内联函数，然后才从同一 warp 中的其他线程读取数据。
参数 %membership 是一个 32 位掩码，其中每个位对应于 warp 中的一个 lane。 1 表示线程位于子集中。
当前执行 warp 中的每个线程都将基于输入参数 %b、 %c 和 %mode 计算源 lane 索引 j。 如果计算的源 lane 索引 j 在范围内，则返回的 i32 值将是 lane j 中 %a 的值；否则，它将是当前线程中 %a 的值。 如果与 lane j 对应的线程处于非活动状态，则返回的 i32 值是未定义的。 如果源 lane j 在范围内，则返回的 i1 值设置为 1，否则设置为 0。
参数 %mode 必须是常量，其编码在下表中指定。
编码 |
含义 |
|---|---|
0 |
IDX |
1 |
UP |
2 |
DOWN |
3 |
BFLY |
参数 %b 指定源 lane 或源 lane 偏移量，具体取决于 %mode。
参数 %c 包含两个打包的值，分别指定用于将 warp 逻辑拆分为子段的掩码和用于钳位源 lane 索引的上限。
以下伪代码说明了此内联函数的语义。
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}
请注意，如果源 lane 处的线程不在 %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 中的一个 lane。 1 表示线程位于子集中。
@llvm.nvvm.vote.sync() 在同步后对 %membermask 中所有线程的源 %predicate 执行归约。 返回值在 %membermask 中的所有线程中都相同。 返回聚合中保存返回值的元素取决于 %mode。
参数 %mode 必须是常量，其编码在下表中指定。
编码 |
含义 |
返回值 |
|---|---|---|
0 |
ALL |
|
1 |
ANY |
|
2 |
EQ |
|
3 |
BALLOT |
|
对于 BALLOT 模式，i32 值表示 ballot 数据，其中包含来自 %membermask 中每个线程的 %predicate 值，该值位于与线程通道 ID 对应的位位置。与不在 %membermask 中的线程对应的位值为 0。
请注意，如果源 lane 处的线程不在 %membermask 中，则返回值未定义。
如果执行线程不在 %membermask 中，则此内联函数的行为未定义。
对于 compute_62 或更低版本，%membermask 中的所有线程都必须在收敛时调用相同的 @llvm.nvvm.vote.sync()，并且当调用 intrinsic 时，只有属于 %membermask 的线程才能处于活动状态。否则，行为是未定义的。
15.6.4. 匹配
以下 intrinsic 同步 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)
这些 intrinsic 导致执行线程等待，直到与 %membermask 对应的所有线程都使用相同的 %membermask 值执行了相同的 intrinsic，然后才在子集中的所有线程之间执行操作数 %value 的广播和比较。
参数 %membership 是一个 32 位掩码，其中每个位对应于 warp 中的一个 lane。 1 表示线程位于子集中。
i32 返回值是一个 32 位掩码，其中掩码中的位位置对应于线程的通道 ID。
在 any 版本中，i32 返回值设置为 %membermask 中具有与操作数 %value 相同值的活动线程的掩码。
在 all 版本中，如果 %membermask 中的所有活动线程都具有与操作数 %value 相同的值，则 i32 返回值设置为 %membermask，并且 i1 值设置为 1。否则，i32 返回值设置为 0，并且 i1 返回值也设置为 0。
如果执行线程不在 %membermask 中，则此内联函数的行为未定义。
这些 intrinsic 仅在 compute_70 或更高版本上可用。
15.6.5. 矩阵运算
这是预览功能。未来版本中可能会删除此支持。
NVVM 提供了用于矩阵乘法运算的 warp 级 intrinsic。核心运算是以下形式的矩阵乘法和累加：
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. 加载片段
以下 intrinsic 同步 warp 中的所有线程，然后为每个线程加载矩阵的片段。
; 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);
这些 intrinsic 从内存位置 %ptr 加载并返回矩阵片段。内存中的矩阵必须采用规范矩阵布局，前导维度为 %ldm。%rowcol 指定内存中的矩阵是行优先 (0) 还是列优先 (1)。%rowcol 必须是常量值。
返回的值序列表示调用线程持有的片段。矩阵元素如何在片段之间分布对用户是不透明的，并且对于矩阵 A、B 和累加器是不同的。因此，提供了三种变体（即 ld.a、ld.b 和 ld.c）。
这些 intrinsic 基于地址空间重载。地址空间编号 <n> 必须是 0（通用）、1（全局）或 3（共享）。
如果 warp 中的任何线程已退出，则此 intrinsic 的行为是未定义的。
这些 intrinsic 仅在 compute_70 或更高版本上可用。
15.6.5.2. 存储片段
以下 intrinsic 同步 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);
这些 intrinsic 将累加器片段存储到内存位置 %ptr。内存中的矩阵必须采用规范矩阵布局，前导维度为 %ldm。%rowcol 指定内存中的矩阵是行优先 (0) 还是列优先 (1)。%rowcol 必须是常量值。
这些 intrinsic 基于地址空间重载。地址空间编号 <n> 必须是 0（通用）、1（全局）或 3（共享）。
如果 warp 中的任何线程已退出，则此 intrinsic 的行为是未定义的。
这些 intrinsic 仅在 compute_70 或更高版本上可用。
15.6.5.3. 矩阵乘法和累加
以下 intrinsic 同步 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);
这些 intrinsic 执行矩阵乘法和累加运算。%rowcol 指定 A 和 B 片段的布局。它必须是常量值，可以具有以下值和语义。
编码 |
含义 |
|---|---|
0 |
A 片段是行优先，B 片段是行优先 |
1 |
A 片段是行优先，B 片段是列优先 |
2 |
A 片段是列优先，B 片段是行优先 |
3 |
A 片段是列优先，B 片段是列优先 |
已删除对 %satf 的支持，此操作数必须为常量零。
如果 warp 中的任何线程已退出，则这些 intrinsic 的行为是未定义的。
这些 intrinsic 仅在 compute_70 或更高版本上可用。
16. 源代码级调试支持
为了启用 IR 模块的源代码级调试，NVVM IR 支持调试 intrinsic 和调试信息描述符来表达调试信息。调试信息描述符使用专门的元数据节点表示。当前的 NVVM IR 调试元数据版本为 3.1。
NVVM IR 调试支持基于 LLVM 7.0.1（用于 Blackwell 之前的目标）和 LLVM 18.1.8（用于 Blackwell 及更高版本的目标）。有关 IR 的完整语义，本章的读者应参考官方 LLVM IR 专用元数据节点和 源代码级调试 文档。Blackwell 及更高版本的目标应分别参考 此处 和 此处 文档。
请求调试支持时，模块中需要存在以下元数据节点
命名元数据节点
!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 托管变量的支持。请参阅 支持的属性。
1.2 版本
针对 CUDA 7.0 更新到 LLVM 3.4。
删除地址空间 intrinsic，改用
addrspacecast。添加了有关源代码级调试支持的信息。
1.3 版本
为 CUDA 8.0 添加了对 LLVM 3.8 的支持。
1.4 版本
添加了对 warp 级 intrinsic 的支持。
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+ 集群 intrinsic 和 max_blocks_per_cluster 内核属性的支持。
弃用了对 32 位编译的支持。
2.0 版本
将 NVVM IR 更新到 2.0 版本，该版本与 NVVM IR 1.x 版本不兼容
删除了地址空间转换 intrinsic。当存在这些 intrinsic 时，2.0 IR 上的 IR 验证器将给出错误。建议 libNVVM 的客户端改用 addrspacecast 指令。
对支持的数据布局进行更严格的错误检查。
不再支持循环后沿上的旧式循环展开 pragma 元数据。建议客户端使用 LLVM 框架定义的新循环 pragma 元数据。
不再支持使用非 undef 值初始化共享变量。在 1.x 版本中，这些初始化程序被静默忽略。此功能使 2.0 版本与 1.x 版本不兼容。
19. 声明
19.1. 声明
本文档仅供参考，不应被视为对产品的特定功能、状况或质量的保证。NVIDIA Corporation（“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 已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，但 NVIDIA 对本文所述产品的客户的累计总责任应根据产品的销售条款进行限制。
19.2. OpenCL
OpenCL 是 Apple Inc. 的商标，已获得 Khronos Group Inc. 的许可使用。
19.3. 商标
NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其关联的各自公司的商标。