CUDA 1つの変数に対してL1キャッシュを無効にする

CUDA 2.0デバイスでは、特定の1つの変数に対してのみL1キャッシュを無効にする方法はありますか？コンパイル時にL1キャッシュを無効にできることがわかっていれば、すべてのメモリ操作でフラグ-Xptxas -dlcm=cgをnvccに追加します。しかし、私は特定のグローバル変数上のメモリ読み取りのためにのみキャッシュを無効にして、残りのすべてのメモリがL1キャッシュを読み取るようにします。CUDA 1つの変数に対してL1キャッシュを無効にする

私がウェブで行った検索に基づいて、可能な解決策はPTXアセンブリコードによるものです。

出典

2012-09-23 zeus2

、ここでの例です：

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

あなたは、.f32（フロート）または.s32（int型）などに.f64を交換することで、これを簡単に変えることができ"= f"（float）または "= r"（int）etcの場合、return_valueの制約は "d"です。 before（addr） - "l"の最後の制約は64ビットアドレッシングを示します.32ビットアドレッシングを使用している場合は "r"でなければなりません。

出典

2012-09-23 20:39:48 Reguj

ありがとう！それは素晴らしい仕事です！ – zeus2

@Reguj、これはどこのNVIDIAのヘッダーでも提供されていませんか？ – einpoklum

[this]（https://nvlabs.github.io/cub/classcub_1_1_cache_modified_input_iterator.html#details）に興味があります –

インラインPTXを使用して、変数をロードして保存することができます。 ld.cgおよびst.cg命令は、L2内のデータのみをキャッシュします。キャッシュ演算子は、8.7.8.1ドキュメントのPTX ISA 2.3のキャッシュ演算子で説明されています。指示または興味はldおよびstです。インラインPTXについては、Using Inline PTX Assembly in CUDAに記載されています。

出典

2012-09-23 16:56:22

変数をvolatileと宣言すると、その変数はFermi GPUのL2キャッシュにのみキャッシュされます。コンパイラが別のスレッドによって書き込まれる可能性があると仮定しているため、反復負荷の除去などの一部のコンパイラ最適化は、揮発性変数に対して実行されません。あなたはインラインPTXを使用することができ、上記で述べたように

出典

2012-09-23 17:59:02 Heatsink

私は、プログラミングモデルが揮発性変数のキャッシャビリティについてどのような表現もしていないと思います。 – ArchaeaSoftware

@Archaeaフェルミアーキテクチャは、キャッシュの一貫性がないため、揮発性データのキャッシュを実行不可能にします。過去にCUDAのドキュメントでエラーが発生したため、CUDAのメモリモデルのドキュメントは信頼できません。 – Heatsink

私は揮発性の変数逸散法で解決策を試しましたが、うまくいかなかったのです。変数が再びキャッシュされているようです。 – zeus2

CUDA 1つの変数に対してL1キャッシュを無効にする

答えて

関連する問題