缺少 8 位变量的 CUDA 内联 PTX 约束字母，以禁用 8 位变量的 L1 缓存（布尔）

Posted 2023-02-16

【中文标题】缺少 8 位变量的 CUDA 内联 PTX 约束字母，以禁用 8 位变量的 L1 缓存（布尔）

【英文标题】：Missing CUDA inline PTX constraint letter for 8 bit variables in order to disable L1 cache for 8 bit variable (bool)

【发布时间】：2013-01-10 15:55:00

【问题描述】：

简介

在this question 中，我们可以学习如何为单个变量禁用 L1 缓存。 这是公认的答案：

如上所述，您可以使用内联 PTX，这是一个示例：

__device__ __inline__ double ld_gbl_cg(const double *addr) 
  double return_value;
  asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr));
  return return_value;

您可以通过将 .f64 替换为 .f32（浮点数）或 .s32 来轻松更改此设置 (int) 等，return_value "=d" 的约束为 "=d" (float) 或 "=r" (int) 等注意 (addr) 之前的最后一个约束 - "l" - 表示 64 位寻址，如果您使用 32 位寻址，它 应该是“r”。

但是，现在我想加载一个布尔值（1 字节）而不是浮点数。所以，我认为我可以做这样的事情（对于架构 >=sm_20）：

__device__ inline bool ld_gbl_cg(const bool* addr)
  bool return_value;
  asm("ld.global.cg.u8 %0, [%1];" : "=???"(return_value) : "l"(addr));
  return return_value;

，其中“？？？”应该是布尔值的适当约束字母，分别是 8 位无符号整数（从this question，我推断出这一点，因为注意到对于 >=sm_20，“u8”用于布尔值）。但是，我在 nvidias 文档“Using inline PTX Assembly in CUDA”中找不到合适的约束字母（第 6 页列出了一些约束字母）。所以我的问题是：

问题

是否有任何类型的任何 CUDA 内联 PTX 约束字母：

布尔值 无符号 8 位整数 或evtl 8位二进制变量

如果没有，我该怎么办（在介绍中解释）？ - 参数“b0”、“b1”等很快讨论过here，有帮助吗？

非常感谢您的任何帮助或 cmets！

更新

我还需要一个从 L2 缓存而不是全局内存读取的存储函数 - 即与上述 ld_gbl_cg 函数互补的 存储函数（只有拥有此函数后，我才能完全验证njuffa 的回答有效）。 根据下面 njuffa 的回答，我的最佳猜测是：

__device__ __forceinline__ void st_gbl_cg (const bool *addr, bool t)

#if defined(__LP64__) || defined(_WIN64)
    asm ("st.global.cg.u8 [%0], %1;" : "=l"(addr) : "h"((short)t));
#else
    asm ("st.global.cg.u8 [%0], %1;" : "=r"(addr) : "h"((short)t));
#endif

但是，编译器给出警告“参数“addr”已设置但从未使用”，并且程序在运行时失败并出现“未指定的启动失败”。 我也尝试使用 .u16 而不是 .u8，因为我不知道它到底指的是什么。然而结果是一样的。

（附加信息） PTX 3.1 文档中的以下段落似乎对这个问题很重要：

5.2.2 子字大小的限制使用 .u8、.s8 和 .b8 指令类型仅限于 ld、st 和 cvt 指令。这 .f16 浮点类型仅允许在往返转换中使用 .f32 和 .f64 类型。所有浮点指令仅在 .f32 和 .f64 类型。为方便起见，使用 ld、st 和 cvt 指令 允许源和目标数据操作数比 指令类型大小，以便可以加载、存储窄值， 并使用常规宽度寄存器进行转换。例如，8 位或 16 位值可以直接保存在 32 位或 64 位寄存器中 正在加载、存储或转换为其他类型和大小。

【问题讨论】：

根据PTX ISA guide 5.2 节，.u8 是无符号的 8 位整数。我认为没有布尔内置类型。

@Robert - 感谢您的评论。是的，我知道 .u8 代表 PTX 中的 8 位无符号整数。但是我找不到的是 CUDA inline PTX 中的约束字母，它对应于 PTX 中的 .u8 寄存器。 IE。我不知道用什么字母来代替“？？？”在上面的代码中（介绍中的第二个代码）。例如，如果将其替换为“r”，则会收到以下错误：“错误：asm 操作数类型大小（1）与约束 'r' 隐含的类型/大小不匹配”。出现此错误是因为 r 代表 4 字节无符号整数，而不是 1 字节无符号整数...

文档“在 CUDA 中使用内联 PTX”列出了可用的约束。字节大小的操作数没有限制。这似乎是有道理的，因为没有可以将字节大小的变量绑定到的字节大小的寄存器。尝试加载到使用 .reg .u32 声明的 32 位临时寄存器中，并使用“=r”约束。

【参考方案1】：

根据文档“在 CUDA 中使用内联 PTX”，对于字节大小的操作数没有限制。我能说的最好的，最接近所需功能的是通过中间“短”移动数据。这会产生一个额外的 SASS 指令，用于从“short”到“bool”的转换。

__device__ __forceinline__ bool ld_gbl_cg (const bool *addr)

    short t;
#if defined(__LP64__) || defined(_WIN64)
    asm ("ld.global.cg.u8 %0, [%1];" : "=h"(t) : "l"(addr));
#else
    asm ("ld.global.cg.u8 %0, [%1];" : "=h"(t) : "r"(addr));
#endif
    return (bool)t;

【讨论】：

非常感谢您的回答！它似乎工作！ （而且额外的转换对我来说应该不重要:)）要知道它是否确实有效，我需要与上述 ld_gbl_cg 函数互补的 store 函数。我将此添加到上面的问题中（请参阅UPDATE）。如果你能帮我解决这个问题，那就太棒了！

加载有缓存模式后缀，我不知道商店也有它们。您的 store 函数无法编译的原因是您绑定不正确。 %0 和 %1 都是读绑定，所以 addr 应该用“l”/“r”绑定，没有“=l”/“=r”。你可能需要一个“记忆”破坏者。在接下来的两天里，我没有时间研究细节并实际尝试。

非常感谢您的评论。我会在此基础上尝试一下（但 PTX 对我来说真的是陌生的土地）。如果我没有误解 PTX 3.1 文档 (docs.nvidia.com/cuda/pdf/ptx_isa_3.1.pdf)，“缓存模式后缀”确实存在：表 84，第 120 页列出了“内存存储指令的缓存操作符”。

如果我在 'st_gbl_cg' 函数中使用以下 asm 语句，它会在没有警告的情况下编译并允许将数据存储到全局位置（对于 64 位地址的情况）：asm ("st.global.cg.u8 [%0], %1;" :: "l"(addr) , "h"((short)t)); 唯一的事情，我还不确定，'ld_gbl_cg' 和 'st_gbl_cg' 是否真的作用于 L2 缓存而不是全局内存。事实上，到目前为止，将缓存运算符 .cg 用于“ld_gbl_cg”和“st_gbl_cg”可以得到与使用 .cv（加载易失性）和 .wt（直写）相同的性能。 IE。两者似乎都从同一个位置读取...

我用一个带有一个循环的内核进行了一些测试，该循环包含对 ld_gbl_cg 和 st_gbl_cg 的调用。无论我使用 .cg 进行读写还是使用 .cv（加载易失性）和 .wt（直写），迭代的时间都是一样的。尽管如此，我接受上面给出的答案，因为它基本上回答了我原来的问题。如果有人报告缓存操作符对加载存储性能影响的一些经验，我当然会非常高兴...