问缺少8位变量的内联PTX约束字母，以便禁用8位变量的L1缓存(bool)

EN

INTRODUCTION

在这个问题中，我们可以学习如何禁用单个变量的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”(浮点数)或"=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文档"嵌入式PTX组件在CUDA中的应用“中找不到合适的约束字母(在第6页列出了一些约束字母)。所以我的问题是：

问题

1. 下列任何类型是否有任何CUDA内联PTX约束信函：

- boolean 
- unsigned 8 bit integer
- or evtl 8 bit binary variable

​

1. 如果不是，在我的情况下(在导言中解释)，我能做些什么？-参数"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，因为我不知道它到底指的是什么。但结果是一样的。

(附加信息)PTX3.1文档中的以下段落对这个问题似乎很重要：

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