缺少CUDA在線PTX約束爲8個變量信以禁用L1緩存爲8位變量（布爾）

Question

介紹

在this question我們可以學習如何禁用L1緩存爲一個單個可變。 這裏是公認的答案：

As mentioned above you can use inline PTX, here is an example:

__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; 
} 

You can easily vary this by swapping .f64 for .f32 (float) or .s32 (int) etc., the constraint of return_value "=d" for "=d" (float) or "=r" (int) etc. Note that the last constraint before (addr) - "l" - denotes 64 bit addressing, if you are using 32 bit addressing, it should be "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位整數unsinged（從this question，我推導出這一點，因爲它指出，對於> = sm_20，「U8」用於一個布爾值）。 Howevever，我找不到NVIDIA的文檔「Using inline PTX Assembly in CUDA」適當約束字母（第6頁中列出的一些約束字母）。所以我的問題是：

問題

1. 是否有任何類型的任何CUDA在線PTX約束字母：

   • 布爾
   • 無符號的8位整數
   • 或evtl 8位二進制變量

2. 如果不能，我能在我的情況做（引言中所述）？ - 參數「b0」，「b1」等簡要討論here是否有幫助？

非常感謝您提前任何幫助或意見！

UPDATE

我還需要一個存儲功能從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 
} 

然而，編譯器會發出警告「參數‘地址’設置，但從來沒有使用過」和PROGRAMM在運行時失敗，「未指定發射失敗」 。 我也試過.u16而不是.u8，因爲我不知道它到底是什麼。但結果是一樣的。

（附加信息）在PTX 3以下段落。1個文件，似乎對這個問題很重要：

5.2.2 Restricted Use of Sub-Word Sizes The .u8, .s8, and .b8 instruction types are restricted to ld, st, and cvt instructions. The .f16 floating-point type is allowed only in conversions to and from .f32 and .f64 types. All floating-point instructions operate only on .f32 and .f64 types. For convenience, ld, st, and cvt instructions permit source and destination data operands to be wider than the instruction-type size, so that narrow values may be loaded, stored, and converted using regular-width registers. For example, 8-bit or 16-bit values may be held directly in 32-bit or 64-bit registers when being loaded, stored, or converted to other types and sizes.

Answer 1

根據該文件「中使用內聯PTX在CUDA」，有一個字節長度的操作數沒有限制。最好我可以告訴你，最接近你想要的功能是通過一箇中間的'短'來移動數據。這會導致一個額外的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; 
}