自己的原文哦~ https://blog.51cto.com/whaosoft/14102097
一、NVIDIA GPU Prefetch行为分析
虽然NVIDIA PTX 的 level::prefetch_size 被文档称为“性能提示”，但在大多数 GPU 上几乎不起作用；仅在 H20 上配合特定线程数与访存宽度时才真正触发 L2 预取，并可带来约 12% 的 Kernel 提速。
近期，工作内容上的一些变动促使我更加关注一些算子层面的性能分析与优化工作。在优化一个MoE算子的性能时，我和同事们注意到，Nvidia GPU在使用LDG指令从Global Memory中加载数据到寄存器时，可以通过内联PTX的方式为ld.global指令设置level::prefetch_size这个qualifier(https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-ld)，在生成SASS汇编指令时，可以生成具有L2 Prefetch Modifier的LDG指令（关于SASS汇编指令的细节，可以参考reed大佬的文章(https://zhuanlan.zhihu.com/p/692445145)）。因此，我们尝试在Kernel代码中使用这个特性，但是在效果上，我们却并未观察到预期的性能提升。因此，在空闲的时候，我做了一些实验，基于ncu简要的分析了一下近几代Nvidia GPU架构的Prefetch行为。我将实验的内容和结果记录于本文当中，以供相关的同学进行参考，若文中有不严谨的地方，也欢迎大家在评论区批评指正。
注：本文中所有的实验代码均已上传至Github——https://github.com/HydraQYH/cuda_prefetch_experiment
实验环境
本文中的实验对比了三种架构的GPU：
- A100 -- SM80 Ampere架构
- L4 -- SM89 Ada架构
- H20 -- SM90 Hopper架构
实验内容
CUDA/PTX编程模型是一种SIMT编程模型，编写Kernel代码时，我们是以单个Thread的视角进行编程的。当我们对加载数据的操作有一些特殊的要求时（比如设置Cache Operator(https://docs.nvidia.com/cuda/parallel-thread-execution/#cache-operators)，指定缓存逐出策略(https://docs.nvidia.com/cuda/parallel-thread-execution/#cache-eviction-priority-hints)），我们可以通过内联PTX的方法以满足我们的需求。使用内联PTX时，可以通过设置ld.global指令的type qualifier指定单个线程要加载的数据类型和大小。其中，数据的大小以bit为单位，可以指定为8bits，16bits，32bits，64bits，128bits等。由PTX指令生成SASS汇编指令时，即可生成具有不同Access Size的LDG指令。
虽然从SIMT编程的视角来看，每个Thread都是独立的，但是GPU硬件在运行Kernel时，是以warp为最基本的执行单元进行调度的，执行访存指令时，实际上也是以warp为一个整体，将warp内所有激活线程的访存请求汇聚成一个单一的request，向相关硬件单元发送。请求的数据在Global Memory中是以sector为基本单元存储的，一个sector是一个对齐的32Byte的数据Chunk，一个单一的request通常会请求多个sector，这些sector会经由L1/L2 Cache传输给GPU的SM，在L1/L2 Cache中，4个sector组成一个128Byte的Cache Line。关于这部分细节，大家可以参考ncu的文档中关于访存模型(https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#id28)的部分。
在真实的CUDA Kernel中，加载数据时通常会尽可能的使用较大的Access Size，并使一个warp内的激活线程访问连续的地址空间（Coalesced Access to Global Memory(https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#coalesced-access-to-global-memory)），以提升有效显存带宽（Effective Bandwidth(https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#effective-bandwidth-calculation)）。常用的Access Size有32bits，64bits，128bits等。如果一个warp内的激活线程访问连续的32bits，64bits，128bits的数据，那么从warp的整体角度来说，一个warp将会访问连续的128Byte，256Byte，512Byte的数据。
为了更容易的观察Prefetch的行为，我们希望执行具有不同Access Size的LDG指令时，都只加载一个cache line大小（128Byte）的数据，因此，对于Access Size为32bits的LDG指令，我们正常激活一个warp内所有的线程即可，但是对于Access Size为64bits和128bits的LDG指令，我们就需要分别的只激活一个warp内的16个线程（half warp）和8个线程（quarter warp）。
最终，我们设计的实验程序如下：仅launch一个warp/half warp/quarter warp，首先执行一条带有Prefetch功能的LDG指令加载128Byte的数据，然后sleep一小段时间等待相关的加载操作彻底完成，然后执行两条普通的LDG指令加载接下来的两段连续的128Byte的数据，基于ncu观察L2 Cache中sector的命中率，以分析第一条带有Prefetch功能的LDG指令是否真的预取了接下来的数据。示例代码如下：
__device__ __always_inline float4 ldg_f32v4_prefetch_256B(const float4 *ptr) {
float4 ret;
asm volatile ("ld.global.L2::256B.v4.f32 {%0,%1,%2,%3}, [%4];" : "=f"(ret.x), "=f"(ret.y), "=f"(ret.z), "=f"(ret.w) : "l"(ptr));
return ret;
}
__global__
void prefetch_kernel_load_128_prefetch_256(const float* in, float* out) {
const float4* value_0_ptr = reinterpret_cast<const float4*>(in);
float4 value_0 = ldg_f32v4_prefetch_256B(value_0_ptr + threadIdx.x);
__nanosleep(1000000); // Sleep 1us
float4 value_1 = value_0_ptr[threadIdx.x + blockDim.x * 1];
float4 value_2 = value_0_ptr[threadIdx.x + blockDim.x * 2];
// ...
}我们在前文提到的三种GPU架构上进行了实验，实验结果如下表所示：
基于以上表格可以发现，绝大多数情况下，level::prefetch_size这个qualifier并没有起作用。只有在H20上使用half warp/quarter warp执行Access Size为64bits/128bits的LDG指令时，设置level::prefetch_size为L2::256B，才会预取接下来的128Byte数据。 如果我们观察ncu中的Memory Tables，可以发现在预取生效时，发送到L2 Cache的requests会增多，如下图所示：
图1:未设置Prefetch时的L2 Cache统计数据
图2:设置Prefetch为256Byte时的L2 Cache统计数据
一个例子
既然我们已经知道了GPU硬件什么时候会发生预取行为，那么接下来我们构造一个简单的例子来验证预取对性能的影响。
我们设计了一个程序，在H20上进行实验，对一个Tensor执行一个Pointwise的tanh激活操作，然后进行一个Sum操作，基于Pytorch的代码如下：
torch.sum(torch.tanh(x))现在，我们基于手写CUDA Kernel来实现上述操作，不过，由于需要Sum的数据量较大，我们并不在Kernel内部一次性完成全局的Sum操作，而是在每一个ThreadBlock内部做一个局部的Sum操作，并将局部Sum操作的结果写入Global Memory，后续的全局Sum操作则由另一个Kernel完成，这也是处理大规模数据规约的常用办法。
执行Pointwise的tanh操作时，我们将所有的线程以8个线程（quarter warp）为一组进行划分，每一组线程通过循环处理连续的4 x 8 x LOOP_COUNT个float元素，每一次循环当中，每个线程处理一个float4类型的元素，8个线程一共处理连续的32个float类型的元素，循环LOOP_COUNT次后，即可完成4 x 8 x LOOP_COUNT个float元素的处理，在我们的代码中，模板参数LOOP_COUNT设定为64，也就是说，每一组线程处理连续的2048个float元素。
每一次循环中，每组线程通过执行Access Size为128bits的LDG指令加载128Byte的数据（32个float元素），我们可以通过内联PTX的方式指定level::prefetch_size为L2::256B ，以要求生成的LDG指令预取接下来的128Byte数据到L2 Cache中，同时，设置一组对照实验，使用不带Prefetch的普通LDG指令，对比两个Kernel的性能，其示例代码如下：
__device__ __always_inline float4 ldg_f32v4_prefetch_256B(const float4 *ptr) {
float4 ret;
asm volatile ("ld.global.L2::256B.v4.f32 {%0,%1,%2,%3}, [%4];" : "=f"(ret.x), "=f"(ret.y), "=f"(ret.z), "=f"(ret.w) : "l"(ptr));
return ret;
}
template<int CG_SIZE, int LOOP_COUNT>
__global__ void prefetch_kernel(const float* in, float* out) {
int thread_offset = blockIdx.x * blockDim.x + threadIdx.x;
int group_id = thread_offset / CG_SIZE;
int thread_id = thread_offset % CG_SIZE;
constexpr int elements_per_group = 4 * CG_SIZE * LOOP_COUNT; // unit: float
const float4* ptr = reinterpret_cast<const float4*>(in + elements_per_group * group_id) + thread_id;
float sum = 0.0;
for (int i = 0; i < LOOP_COUNT; i++) {
#ifdef PREFETCH_256
float4 vals = ldg_f32v4_prefetch_256B(ptr);
#else
float4 vals = *ptr;
#endif
// tanh
sum += tanh(vals.x);
sum += tanh(vals.y);
sum += tanh(vals.z);
sum += tanh(vals.w);
ptr += CG_SIZE;
}
sum = blockReduceSum(sum);
if(threadIdx.x == 0) {
out[blockIdx.x] = sum;
}
}我们指定这个Kernel处理256M个float元素，实验结果显示，不使用预取时，该Kernel运行时间为355.90us，而使用预取时，该Kernel运行时间为313.95us，如下图所示：
图3:不使用预取的Kernel
图4:使用预取的Kernel
观察ncu的Memory Chart，可以发现使用预取时，L2 Cache的命中率有显著的提高：
图5:不使用预取时的L2 Cache命中率
图6:使用预取时的L2 Cache命中率
Memory Table也显示，做了预取的Kernel对L2 Cache具有更多的requests & sectors：
图7:不使用预取时的L2 Cache Requests & Sectors
图8:使用预取时的L2 Cache Requests & Sectors
在Warp State统计数据中可以观察到，使用预取时，Stall Long Scoreboard从20个周期降低至17个周期，这是因为使用预取时，在L2 Cache上，针对于某些数据的requests被提前发出了，因此降低了等待这部分数据所带来的固定延迟开销，从而提升了Kernel的性能。
图9:Warp State统计数据
总结
正如PTX文档中所描述的那样，level::prefetch_size只是一个performance hint，我们不能假设使用这个qualifier后，Kernel具有预期的预取行为。
本文通过实验的方式，对近几代Nvidia GPU的预取行为进行了分析，事实上，level::prefetch_size这个qualifier还可以用于异步拷贝指令cp.async(https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async)，这里先留个坑，如果后续有时间和精力的话，笔者会进一步通过实验验证异步拷贝指令的预取行为。