1

前書き

この質問では、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) などに交換することで、これを簡単に変更できます。 (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" が設定されていましたが、使用されていません」という警告を出し、プログラムは実行時に「不特定の起動エラー」で失敗します。また、.u8 の代わりに .u16 を試してみました。正確に何を参照しているのかわからないからです。それでも結果は同じです。

(追加情報) 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の使用」によると、バイトサイズのオペランドに制約はありません。私が知る限り、目的の機能に最も近いのは、中間の「short」を介してデータを移動することです。これにより、「short」から「bool」に変換するためのSASS命令が1つ追加されます。

__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 に答える