cuda - cuda でのプリフェッチ (C コード経由)

Question

C コードを使用して CUDA (Fermi GPU) でデータのプリフェッチに取り組んでいます。Cuda のリファレンス マニュアルでは、C レベルのコードではなく、ptx レベルのコードでのプリフェッチについて説明しています。

cudaコード（cuファイル）を介したプリフェッチに関するドキュメントまたは何かを誰かに教えてもらえますか。どんな助けでも大歓迎です。

Answer 1

PTX マニュアルによると、 PTXでプリフェッチがどのように機能するかは次のとおりです。

PTX 命令を CUDA カーネルに埋め込むことができます。以下は、 NVIDIA のドキュメントからの小さなサンプルです。

__device__ int cube (int x)
{
  int y;
  asm("{\n\t"                       // use braces for local scope
      " .reg .u32 t1;\n\t"           // temp reg t1,
      " mul.lo.u32 t1, %1, %1;\n\t" // t1 = x * x
      " mul.lo.u32 %0, t1, %1;\n\t" // y = t1 * x
      "}"
      : "=r"(y) : "r" (x));
  return y;
}

C の次のプリフェッチ関数で結論付けることができます。

__device__ void prefetch_l1 (unsigned int addr)
{

  asm(" prefetch.global.L1 [ %1 ];": "=r"(addr) : "r"(addr));
}

注意: プリフェッチには Compute Capability 2.0 以上の GPU が必要です。それに応じて適切なコンパイルフラグを渡します-arch=sm_20

Answer 2

このスレッドによると、以下はさまざまなキャッシュプリフェッチ手法のコードです。

#define DEVICE_STATIC_INTRINSIC_QUALIFIERS  static __device__ __forceinline__

#if (defined(_MSC_VER) && defined(_WIN64)) || defined(__LP64__)
#define PXL_GLOBAL_PTR   "l"
#else
#define PXL_GLOBAL_PTR   "r"
#endif

DEVICE_STATIC_INTRINSIC_QUALIFIERS void __prefetch_global_l1(const void* const ptr)
{
  asm("prefetch.global.L1 [%0];" : : PXL_GLOBAL_PTR(ptr));
}

DEVICE_STATIC_INTRINSIC_QUALIFIERS void __prefetch_global_uniform(const void* const ptr)
{
  asm("prefetchu.L1 [%0];" : : PXL_GLOBAL_PTR(ptr));
}

DEVICE_STATIC_INTRINSIC_QUALIFIERS void __prefetch_global_l2(const void* const ptr)
{
  asm("prefetch.global.L2 [%0];" : : PXL_GLOBAL_PTR(ptr));
}