在 cuda 中预取（通过 C 代码）答案

【问题标题】：Prefetch in cuda (through C code)在 cuda 中预取（通过 C 代码）
【发布时间】：2019-07-28 19:50:59
【问题描述】：

我正在通过 C 代码在 CUDA (Fermi GPU) 中进行数据预取。 Cuda 参考手册讨论了 ptx 级别代码而不是 C 级别代码的预取。

谁能给我一些关于通过 cuda 代码（cu 文件）预取的文档或内容。任何帮助将不胜感激。

【问题讨论】：

【解决方案1】：

根据PTX manual，这是预取在 PTX 中的工作方式：

您可以将 PTX 指令嵌入到 CUDA 内核中。这是来自NVIDIA's documentation 的一个小样本：

__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

【讨论】：

您能否提供更多有关预取工作原理的文档，例如对概念本身的解释。
当然！检查此 GPGPU 预取研究并查看参考资料以了解有关该概念的更多信息：cc.gatech.edu/~hyesoon/lee_taco12.pdf

【解决方案2】：

根据this thread，下面是不同缓存预取技术的代码：

#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));
}

【讨论】：

良好的提升。现在我们只需要一个例子，这些实际上提供了好处。
@tera 我有一个通用的经验法则：如果 Nisght Compute 将长记分牌摊位列为最大摊位贡献者，那么您将主要从预取中受益。这个经验法则对我来说适用于 9/10 的案例。这对于占用率较低的内核非常重要（例如当您被限制为扭曲或块时）。