1

介绍

在这个问题中，我们可以学习如何为单个变量禁用 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 (float) 或 .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 位无符号整数（从这个问题，我推断出这一点，因为注意到对于 >=sm_20，“u8”用于布尔值）。但是，我在 nvidias 文档“ Using inline PTX Assembly in CUDA ”中找不到合适的约束字母（第 6 页列出了一些约束字母）。所以我的问题是：

问题

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

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

  2. 如果没有，我该怎么办（在介绍中解释）？-这里简短讨论的参数“b0”、“b1”等有帮助吗？

非常感谢您提前提供任何帮助或意见！

更新

我还需要一个从 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 位寄存器中。

4

1 回答 1

3

根据文档“在 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;
}
于 2013-01-12T02:31:52.137 回答