13

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

私が Web で行った検索に基づいて、考えられる解決策は PTX アセンブリ コードを使用することです。

4

3 に答える 3

15

上記のように、インライン 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" にする必要があります。

于 2012-09-23T20:39:48.510 に答える
5

インライン PTX を使用して、変数を読み込んで保存できます。ld.cg および st.cg 命令は、L2 にデータをキャッシュするだけです。キャッシュ オペレーターについては、 PTX ISA 2.3ドキュメントのセクション 8.7.8.1 キャッシュ オペレーターで説明されています。命令または対象はldおよびstです。インライン PTX については、CUDA でのインライン PTX アセンブリの使用で説明されています。

于 2012-09-23T16:56:22.707 に答える
0

変数を と宣言するとvolatile、Fermi GPU の L2 キャッシュにのみキャッシュされます。繰り返しロードの削除など、一部のコンパイラの最適化は、揮発性変数に対して実行されないことに注意してください。

于 2012-09-23T17:59:02.660 に答える