1. 实验目标实现并比较不同的 Softmax kernel理解共享内存规约、Warp Shuffle、寄存器缓存和 Online Softmax 的性能取舍。2. 实验环境2.1 硬件与软件操作系统Ubuntu LinuxGPUNVIDIA Tesla T4GPU 显存15360 MiBGPU Compute Capability7.5NVIDIA 驱动535.216.01CUDA Toolkit12.4nvcc 版本V12.4.131Nsight Compute2024.1.1.02.2 编译配置编译优化-O2目标架构-archsm_75资源报告-Xptxas-v输入与输出FP32CPU 参考计算FP642.3 计时配置CUDA Event 计时预热 100 次正式运行 1000 次计算单次平均耗时内存分配、数据拷贝、结果验证不计入计时区间3. 正确性验证3.1 CPU 参考实现使用 double 精度计算稳定 Softmax作为 GPU 输出的参考。3.2 检查方法检查输出是否存在 NaN 或 Inf。统计 GPU 与 CPU 结果的最大绝对误差。当前通过阈值1e-6。3.3 测试输入全零输入。每行末尾元素设置为特殊值。对整行输入增加相同常数。固定随机种子的随机输入。不同长度与尾部边界。3.4 输出覆盖与非整除行数验证新增问题各版本共用 d_output后一版若漏写某些位置前一版留下的正确结果可能使检查误通过。在 V8、V9 各自开始计时之前先将输出填为 NaNCHECK_CUDA(cudaMemset(d_output, 0xff, bytes)); CHECK_CUDA(cudaEventRecord(start)); // 再启动该版本的计时循环本实验 FP32 下0xff 填满所有字节对应 NaN未覆盖的位置会被 std::isfinite 检查发现。填充和开始事件在同一 stream 中顺序执行填充操作不计入 kernel 耗时。边界配置M1025、N257、THREADS256、grid_v9256。V8 最后一个 Block 仅一个 Warp 有有效行V9 的 Block 0 处理五行其余 Block 各处理四行。加入 NaN 初始化后两版均通过当前用例的数值及覆盖检查。已有 memcheck 结果按对应版本和输入归档数值与覆盖检查通过不自动代表所有后续版本都通过内存检查。4. Kernel 优化过程版本关系与对照原则版本号表示实验顺序不代表性能排名后一版也不一定保留前一版全部优化。版本主要变化逻辑输入扫描V2分别规约 Max / Sum三遍V3保存输入随后保存指数用于输出一遍V4 / V4_OPTOnline 递推提取共同指数两遍V5两级 Warp 合并状态 (m,d)两遍V6 / DOUBLE / FETCH向量化、双遍向量化、输出预取两遍V7缓存完整线程输入一遍V8从 V5 改起一个 Warp 一行两遍V9整个 Block 通过 Grid Stride 顺序遍历多行两遍逻辑扫描次数不等于实际 DRAM 读取次数。V8 没有叠加 V7 缓存V9 也不是将 V8 简单加一个循环。V9 保留原始两指数 Online 更新与 V6_DOUBLE 的局部更新写法也不完全一致。4.1 V1Shared Memory 规约实现思路引入 Shared Memory 实现数据复用存放各线程规约得到的局部结果求最大值阶段存放局部最大值最终某行最大值由 0 号下标元素负责求和阶段存放各线程负责的局部和最终行和覆盖先前的 0 号下标元素内容。遇到的问题与修正Q1最初只有tid 0更新max_val和sum其他线程也需要用它们计算输出而max_val和sum不属于 Shared Memory是各线程的私有变量因此其他线程无法获取全局最大值/和最终计算出错。A1删除tid 0条件让每一个线程都能读取到max_val和sum。Q2线程同步问题。注意线程同步应该放在什么位置不同轮之间有什么数据依赖同一轮写入位置不重叠为什么仍要考虑同步A2区分两种依赖。第一种是第一轮规约前依赖各线程刚写入的初始值。如果缺少同步某些线程开始规约时其他线程可能还没写完。第二种是下一轮规约依赖上一轮规约结果。例如 s 128时线程64更新 smem[64] s 64时线程0需要读取smem[64]此时就需要同步。__global__ void softmax_v1(float* input, float* output, int M, int N) { extern __shared__ float smem[]; int row blockIdx.x; int tid threadIdx.x; float* x input row * N; float* y output row * N; float max_val -INFINITY; for(int i tid;i lt; N;i blockDim.x){ max_val fmax(max_val,x[i]); } smem[tid] max_val; __syncthreads(); for(int s blockDim.x / 2;s gt; 0;s gt;gt; 1){ if(tid lt; s) smem[tid] fmax(smem[tid],smem[tid s]); __syncthreads(); } max_val smem[0]; __syncthreads(); float sum 0.0f; for(int i tid;i lt; N;i blockDim.x){ sum expf(x[i] - max_val); } smem[tid] sum; __syncthreads(); for(int s blockDim.x / 2;s gt; 0;s gt;gt; 1){ if(tid lt; s) smem[tid] smem[tid s]; __syncthreads(); } sum smem[0]; float inv_sum 1.0f / sum; for(int i tid;i lt; N;i blockDim.x){ y[i] inv_sum * expf(x[i] - max_val); } }4.2 V2Warp Shuffle 规约相比 V1 的变化V1 先由每个线程计算自己的局部结果再将结果写入 Shared Memory通过整个 block 的树形规约得到最终结果。相邻规约轮之间存在数据依赖需要进行 block 同步。Warp Shuffle 在各 Warp 内先规约得到局部结果再指定某一 Warp 负责最终整体规约得到最终结果。好处在于同一 Warp 内的各个线程可以直接交换寄存器值减少规约过程中的 Shared Memory 读写和 block 同步但最后的跨 Warp 结果交换仍需 Shared Memory 和同步。4.3 V3寄存器缓存相比 V2 的变化V2 分三遍访问输入求最大值、求指数和、计算输出。其中指数在求和与输出阶段各计算一次。V3 使用线程私有数组缓存自己负责的输入。得到整行最大值后将数组中的原始值更新为指数同时计算局部指数和最后直接使用保存的指数完成输出。引入固定长度的数组数组长度由模板参数 ITEMS 决定表示每个线程最多处理的元素个数元素位置为 threadIdx.x i * blockDim.x 步长取 blockDim.x 的解释设一行有1024个元素共256个线程则每个线程负责四个元素每轮每个线程各处理一个元素共处理256个元素合计四轮。Local Memory 排查过程最初使用私有数组后V3 在部分测试中反而比 V2 慢。检查早期每线程缓存 4 个 float 的版本编译报告显示16 bytes stack frame0 bytes spill stores0 bytes spill loads进一步检查 PTX发现存在st.local.f32ld.local.f32说明数组引入了 Local Memory 访问没有达到预期的寄存器缓存效果。Local Memory 虽然是线程私有的底层仍使用设备内存。当时数组长度虽然固定但循环次数由运行时参数决定数组通过动态下标访问编译器不容易将各槽位转换为独立寄存器。随后将循环改为编译期固定次数并使用模板参数 ITEMS 控制数组长度配合循环展开。重新编译后局部栈消失寄存器用量增加性能改善。寄存器用量与 occupancy 的取舍occupancy 指一个SM上活跃Warp数与该SM所能容纳的最大Warp数比值。由于SM上寄存器数量固定线程所需寄存器数量越多则该SM所能容纳的线程就越少活跃Warp自然减少occupancy降低但带来的性能收益可能抵消occupancy下降带来的弊端。遇到的问题与修正Q线程私有数组不一定都放在寄存器里没有 spill 不代表没有Local Memory数组可能直接被分配到局部内存并非先进入寄存器再溢出。Afloat* private_elm new float[4];这是设备端动态分配申请的数据位于设备堆中不能作为预期的寄存器缓存。float private_elm[4];以上是对线程私有数组的声明但并不保证数组保存的物理位置是寄存器。受运行时计算、寄存器压力等因素影响都可能使它进入Local Memory。固定长度、小容量、编译器可以确定的访问下标有利于将私有数组拆成寄存器值。但数组是否真正位于寄存器中需要通过编译报告及生成的指令验证不能仅凭声明方式判断。输入最初仍需从 Global Memory 加载之后才复用寄存器中的副本。另外Global Memory 加载可能命中硬件缓存不能将加载次数直接等同于 DRAM 访问次数。4.4 V4Online Softmax实现思路同时维护局部最大值与局部和与朴素Softmax的三遍扫描相比由于将求最大值与求和阶段合并便只需两遍扫描——求最大值与求和、归一化。每次 Online 更新都计算两个指数相比 V3 节省了什么、增加了什么相对 V2输入扫描由三遍减为两遍相对 V3输入读取由一遍变为两遍但不再需要缓存每个线程负责的全部输入。Online 递推增加指数计算与状态依赖但减少保存整份输入所需的寄存器实际寄存器数量以对应配置的编译报告为准。4.5 V4_OPT共享指数计算优化思路每次Online无论 m 是否需要更新都会对 d 进行重复的指数计算因而引入分支减少指数运算推测有可能带来性能提升。遇到的问题与修正Q线程少算指数Warp 不一定同比减少工作各线程执行路径只有一条但由于 if 分支的存在一个Warp内有多个执行路径。SASS 显示两条路径分别保留了指数指令随机输入的分支一致性也从 100% 降到 85.54%。这些证据支持“新增了分歧执行”但不能单独量化其性能损失。A原来的分支化简在数学上成立其中cur 是当前元素值。观察指数可以得到cur - m 与 m - cur 互为相反数可以提取出共同指数 e用同一个 e 更新 d float e expf(-fabsf(cur - m)); if(cur m){ d d e; } else{ d d * e 1; m cur; }SASS 和分支指标验证新 SASS 确认- 局部更新共享了一条指数路径。- 条件更新使用选择、谓词指令。- 不再通过两条数据相关跳转路径分别计算指数。分支一致性恢复到 100%不代表线程判断结果都相同而是编译后的条件处理方式改变了。4.6 V5Block 级并行 Warp Shuffle 合并规约优化思路每个线程先扫描分配给自己的元素维护局部状态m max(x_i)d Σ exp(x_i - m)两组状态合并时将分母统一到相同的最大值基准m max(m1, m2)d d1 * exp(m1 - m) d2 * exp(m2 - m)第一级由每个 Warp 在寄存器间交换 m、d 并合并lane 0 把该 Warp 结果写入共享内存。同步后由第一个 Warp 读取所有 Warp 的状态执行第二级合并。最后通过共享内存及同步将行结果提供给整个 Block。遇到的问题与取舍共享内存保存 float 数值不能误声明为 float 指针数组数组容量使用编译期常量实际有效项数为 Block 的 Warp 数。第二级合并时第一个 Warp 的每个 lane 将共享内存的一组状态读入自己的 m、d再复用 Warp 规约函数。空线程的状态是 (-∞, 0)。两个空状态直接代入公式会出现 -∞-(-∞)产生 NaN乘以 0 也不能消除 NaN。当前 helper 在取出对端状态后对 d2 0 跳过本次合并。输出的倒数必须保存为浮点数误用 int 会截断小于 1 的值。减少分阶段规约并不保证更快。Online 状态合并在每轮包含指数计算而原先独立 Max/Sum 规约的合并操作更简单。V5 保留为实验对照不写成必然优于 V4_OPT。4.7 V6float4 向量化、双遍向量化与预取优化目标在 V5 的 Online 实现上改变访存组织分别验证第一遍向量化、第二遍也向量化以及提前加载下一组数据的效果。版本第一遍扫描第二遍归一化V6float4 加载标量加载、标量存储V6_DOUBLEfloat4 加载float4 加载、float4 存储V6_FETCHfloat4 加载float4 读写并提前加载下一组DOUBLE 指两遍均向量化不是 double 精度。这三版仍需读取输入两遍当前采用一个 Block 负责一行。线程分工与下标单位访问 x4[k] 对应原数组 x[4k] 至 x[4k3]。N4096、THREADS256 时共有 1024 个完整向量每线程最多处理 4 个向量即 16 个标量。constexpr int VEC_ITEMS (N / 4 THREADS - 1) / THREADS; int vec_idx tid i * blockDim.x; // 四个标量位置为 4 * vec_idx jj0,1,2,3当时的错误曾混用标量与向量下标将 tid 再乘 4也曾用 VEC_ITEMS 次标量输出导致覆盖不足。循环次数和下标必须与该阶段的数据单位对应。行首对齐与尾部处理cudaMalloc 基地址对齐但连续 FP32 行的起点还取决于 row×N。N257 时行偏移依次为 0、1028、2056、3084、4112 字节每四行中只有一行满足 float4 的 16 字节对齐。实际修正对齐行采用 float4 加载与标量尾部未对齐行整行回退标量。两条路径共用后续规约。后续 V9 也曾因遗漏回退触发 misaligned address。测量结果与未获得收益的尝试在 M1024、N4096、THREADS256 的较早同批测量中V5 与 V6 的耗时中位数分别为 186.100 μs、180.570 μsV6 耗时减少约2.97%。另一批包含预取版的测量单独记录V6 为 180.505 / 180.552 μsV6_DOUBLE 为 191.923 / 191.920 μsV6_FETCH 为 188.351 / 188.351 μs。双遍向量化减少了读写指令却更慢预取相对双遍版改善约 1.86%仍未超过该批 V6。SASS 确认第一遍出现 128 位加载预取版中 next 的加载在 current 计算前发起使用不同寄存器。这支持“提前加载被保留”还不能说明延迟完全隐藏。N257 的特殊情况N/464小于 256 个线程对齐路径中每线程最多处理一个向量没有实际的下一组预取。思考题为什么访存指令更少不一定更快N257 时即使 FETCH 版更快为什么也不能归因于下一组加载与当前计算重叠解释双遍向量化虽然减少了访存指令但也可能改变了寄存器需求、指令调度和数据依赖而整体耗时是由这些因素共同决定的。N257时向量数为257/464因此每个线程最多处理一个向量未对齐数据走标量回退因此不存在“下一组加载与当前计算重叠”。4.8 V7Online 完整线程输入缓存优化目标第一遍保存线程负责的完整输入跨过规约阶段后仍可复用消除归一化阶段的第二遍全局输入加载。完整四元素组使用固定容量 float4 私有数组最多三个尾部元素使用独立 float 缓存。N4 时 VEC_ITEMS 可能为 0数组声明至少保留一个占位元素避免零长度数组。对齐行使用 float4未对齐行以四次标量访问组成一组同样保存到私有缓存输出再拆成标量存储。未对齐不等于不能缓存。思考题V6 的 float4 data 也可能位于寄存器为什么仍需要第二遍读取 x而 V7 可以避免请从数据的保存范围与生命周期解释。解释V6 的 float4 data 属于在某轮循环中声明的局部变量其生命周期仅在一次循环的作用域中下一轮循环会将其值更新因此后续仍需读取第二遍x。而V7使用私有数组保存线程负责的完整输入使这些值跨过规约阶段后仍然可用。输出直接读取缓存因而消除了第二遍全局输入加载。编译与指令证据M1024、N4096、THREADS256、VEC_ITEMS4 时V7 使用41 个寄存器/线程stack frame、spill stores、spill loads 均为 0。当前配置没有观察到栈分配或 spill不能把任意私有数组都默认视为寄存器数组。当时采集后确认V7 的全局加载、存储 Warp 指令均为 32,768 次对应 V6_DOUBLE 为加载 65,536 次、存储 32,768 次。加载指令减半与消除第二遍输入读取相对应。同一批普通运行中V7 为 140.987 / 140.965 μsV6_DOUBLE 为 191.945 / 191.939 μs按两次平均值耗时减少约26.55%。相对该批 V3 仅改善约 1.76%这一小幅差异保留测量波动的限制。思考题V7 保存的是原始输入V3 在求和时保存的是指数结果两者在指数重算和寄存器需求上有什么取舍解释V3 将指数结果保存在寄存器中那么后续阶段指数重算只需从寄存器中读取数据而V7则需重新计算V7同时维护m、d寄存器用量比V3更多。4.9 V8一个 Warp 负责一行多行并行版本关系V8 从 V5 的标量 Online 方案改起没有叠加 V7 的完整输入缓存。int lane threadIdx.x % 32; int wid threadIdx.x / 32; int num_warps blockDim.x / 32; int row blockIdx.x * num_warps wid; // 本行元素位置lane、lane32、lane64……THREADS256 时一个 Block 的八个 Warp 分别负责八行。Warp 内完成 Online 规约后用两次 Shuffle 将 lane 0 的 m、d 广播给本 Warp不使用共享内存和 Block 同步。原先若继续合并不同 Warp 的状态会把不同行混在一起。行号越界时需要同一 Warp 的线程一起退出再由其他有效 Warp 继续处理。短行收益与长行代价M1024、THREADS256 的两类输入分别测量N257 时V8 为 7.636–7.973 μsV7 为 15.014–15.494 μsN4096 时V8 为 249.590–249.652 μsV7 为 143.282–143.346 μs。当前实现短行更快、长行更慢尚未确定通用切换阈值。同为 N4096V5 每线程递推 16 次V8 每线程递推 128 次。后续一轮使用前一轮产生的 m、d这属于同一线程内的指令依赖不是轮与轮之间需要线程同步。思考题 · 待本人填写为什么一条更长的 m、d 递推依赖链可能抵消减少跨 Warp 规约与同步的收益解释更长的m、d递推依赖链带来的是一个Warp处理行数据的更多循环、更多递推次数例如N为257时V8中每线程需递推257/328次而N4096时则需递推128次行的加长带来的是单线程更长的串行递推其带来的时间开销很有可能抵消减少跨Warp操作的收益。边界验证M1025 时V8 的 Grid129最后一个 Block 只有一个 Warp 对应有效行。运行前将输出填 NaN 后当前用例的数值与覆盖检查通过。最新调度与总指令证据见第 7.6 节。4.10 V9Grid Stride 行遍历与 V8 的区别for (int row blockIdx.x; row M; row gridDim.x) { // 整个 Block 合作处理当前行 }V8 是同一 Block 的不同 Warp 同时处理不同行V9 是整个 Block 完成一行之后再顺序处理下一行。Grid M 时每个 Block 只处理一行没有实际复用 Block 遍历后续行。Grid M 才会走到后续行。每轮 kernel 启动次数仍为一次缩小 Grid 不等于减少 kernel launch 次数。复用共享内存时末尾同步是阶段边界不会清空或重置共享内存。Grid 大小实验M 1024、N 257、THREADS 256Grid每 Block 顺序处理行数运行 1 / μs运行 2 / μs1024117.79517.914256420.31720.628128822.30523.083256 的配置已在日志打印128 的记录按当时预热及计时两处均修改为 128 的实验操作归档旧输出没有打印该参数。上述是分次运行结果不能视为同一进程内交替测量。在已测试的短行输入上缩小 Grid 没有获得收益。更少的 Block 并没有消除每一行所需的规约、共享内存通信和同步也引入了顺序处理多行的执行组织。Grid Stride 的价值是控制 Grid 和安排行遍历不应写成“采用 Grid Stride 就能加速”。V9 保留为当前条件下未获收益的实验。思考题V8 的多行并行与 V9 的多行遍历在行分工、每行参与线程数、同步和执行先后上分别有什么区别为什么缩小 Grid 不保证加速解释行分工上V8一行由一个Warp负责V9一行由一个Block负责V8参与线程数为一个Warp的线程数通常为32V9参与线程数则为blockDim.xV8对某一行的处理由一个Warp循环完成每轮循环间同步但各行各Warp可并行工作V9的多行遍历由一个block循环实现完成一行后才处理下一行行结果间存在两级规约、共享内存与block级同步。缩小Grid会减少block数量这也导致单个block所要处理的行数增加每行的规约和同步、以及kernel启动次数都没有消除或减少。5. 性能测试方法所有版本使用相同输入。预热次数100。正式计时次数1000。使用 CUDA Event 计时。分配内存、数据拷贝和结果验证均放在计时区间之外。普通运行耗时与 Nsight Compute 采集数据分开记录。6. 性能结果6.1 当前配置M 1024N 4096THREADS 256ITEMS 16随机种子42输入分布均匀分布[-10, 10)6.2 一次普通运行结果版本平均耗时 / μs正确性V1195.854通过V2185.475通过V3147.249通过V4187.140通过V4_OPT184.406通过以上为同一次程序运行结果不代表所有尺寸下的性能。6.3 V6 系列与 V7同一编译批次的普通运行配置M1024、N4096、THREADS256、VEC_ITEMS4。本表是加入 V7 后的一批测量与第 4.7 节预取实验的批次分别记录。版本运行 1 / μs运行 2 / μsV3143.467143.542V6180.556180.597V6_DOUBLE191.945191.939V6_FETCH188.378188.335V7140.987140.965按两次平均值V7 相对本批 V6_DOUBLE 耗时减少约26.55%相对 V3 减少约 1.76%。不同批次不拼接计算收益小幅变化保留测量波动的限制。6.4 V8、V9边界验证后的同次普通运行配置M1025、N257、THREADS256、grid_v8129、grid_v9256预热 100 次、正式启动 1000 次。V8、V9 运行前加入 NaN 初始化。版本平均耗时 / μs最大绝对误差V516.22421.96156e-08V715.10251.90672e-08V87.723522.00473e-08V920.48461.96156e-08上述记录来自同一次程序执行不是多次运行中位数。当前用例均通过有限性与误差检查V8 相对该次 V7 的加速比约为1.96 倍。6.5 V8短行与长行分开记录M1024、THREADS256各输入形状分别普通运行两次NV7 耗时范围 / μsV8 耗时范围 / μs25715.014–15.4947.636–7.9734096143.282–143.346249.590–249.652V8 的优势依赖输入形状。这里只验证了两类行长度尚不能据此确定通用的 kernel 切换阈值。第 6.1–6.2 节保留的是早期 V1–V4_OPT 批次不作为当前全部版本的性能排名。7. Nsight Compute 分析不同配置的报告分别记录不能混用。7.1 全局加载指令与 DRAM 流量配置M 128N 4096THREADS 256ITEMS 16。V2 的全局加载指令执行数为 49,152V3 为 16,384减少到原来的三分之一两者的全局存储指令执行数均为 16,384。指标 smsp__sass_inst_executed_op_global_ld.sum 统计 Warp 级指令执行数不是线程读取的元素个数也不是 DRAM transaction 数。每个 Block 有 8 个 Warp每个线程处理 16 个元素一次扫描对应 128 × 8 × 16 16,384 次 Warp 级加载指令。V2 扫描三遍V3 扫描一遍之后复用寄存器中的输入与指数结果。同尺寸的 Memory Workload Analysis 中V2 的 DRAM 读取量为 2.61 MBV3 为 2.41 MB写入量分别为 1.54 MB、1.29 MB。加载指令减少三分之二并不意味着 DRAM 读取量也减少三分之二V2 的重复加载可能命中 L1/L2。V2 的 L1/TEX 命中率为 38.40%V3 为 0%。V3 消除了重复加载也减少了产生缓存命中的机会不能据此认为它没有访存等待。L2 命中率是聚合统计不能认定所有命中都来自输入 x。DRAM 写入量的差异也不能直接解释为输出元素数量不同。采集使用 --cache-control all重放前清理缓存但 Kernel 内部仍可能复用缓存。它与连续预热后的 CUDA Event 计时不是同一缓存条件应分别解读。7.2 寄存器与 occupancy配置M 128N 16384。V3 使用 256 个线程、ITEMS 64每线程需要 161 个寄存器Block Limit Registers 1每个 SM 只能驻留一个 Block即 8 个 Warp理论 occupancy 为 25%实际为 24.91%。这次限制来自寄存器。将线程数改为 512、ITEMS 改为 32 后每线程寄存器降到 89 个但每个 SM 仍只能驻留一个 Block。由于每个 Block 变为 16 个 Warp理论 occupancy 提高到 50%实际为 49.47%。寄存器数量与 occupancy 的关系取决于资源分配阈值不能认为寄存器稍有增加occupancy 就一定下降。两种配置的 Issued Warp Per Scheduler 均约为 0.25Active Warps Per Scheduler 从 1.99 增到 3.94Eligible Warps 从 0.36 增到 0.69。更多 Warp 驻留没有带来同比增加的实际发射。相应 NCU Duration 为 104.74 μs 和 107.07 μs未显示这次 occupancy 提升使 Kernel 更快。同为 N 16384、THREADS 256V4 只需要 20 个寄存器而 V3 需要 161 个。Online 版本不保存整份线程私有输入降低了长行缓存的寄存器压力但仍有额外指数运算和递推依赖。occupancy 衡量驻留的活跃 Warp 数包括正在等待的 Warp它不同于可发射 Warp 数也不同于计算单元利用率不能单独作为性能更好的依据。7.3 Long Scoreboard配置M 128N 4096THREADS 256。Long Scoreboard 从 V2 的 55.81% 降到 V3 的 35.11%下降 20.70 个百分点。指标 smsp__warp_issue_stalled_long_scoreboard_per_warp_active.pct 描述活跃 Warp 因 L1TEX 相关未完成加载的结果依赖而无法发射的等待不等同于 DRAM 缺失率也不是整个 Kernel 耗时中可以直接减去的比例。V3 减少重复加载结果与访存依赖等待减少的判断一致。但最初仍需加载输入后续使用寄存器并不能消除全部 Long Scoreboard。L1 命中率为 0也不能推出 Long Scoreboard 接近 0请求可能需要从 L2 或 DRAM 返回。V3 同时减少了指数重算并改变了循环与指令组织不能仅凭这一指标将全部加速归因于访存等待下降。7.4 分支分歧分支指标采集配置M 128N 16384THREADS 256ITEMS 64。原始 V4 的分支一致性为 100%最初通过 if/else 分别计算指数的优化版为 85.54%提取共同指数后的版本恢复到 100%。第一次改写中每个线程只走一条路径但同一 Warp 的线程判断可能不同。SASS 显示两条路径分别保留指数指令支持存在分歧执行。随后将共同指数提取到条件更新之前SASS 中局部更新共享指数路径条件处理改为选择与谓词指令。因此 100% 并不表示所有线程对 cur 与 m 的比较都相同而是编译后不再通过这两条数据相关跳转分别计算指数。MUFU.EX2 是以 2 为底的指数近似指令expf 的实现还可能包含其他转换与范围处理指令不能把一行 SASS 直接等同于一次完整 expf。另一次固定配置 M 1024、N 4096、THREADS 256 的三次普通运行中V4 中位耗时为 187.198 μs最终 V4_OPT 为 184.339 μs降低约 1.53%。这是当前配置下的小幅改善不代表所有形状都能获益也不能仅凭分支指标量化这部分收益。7.5 V6–V7向量化与完整输入缓存的指令证据配置M1024、N4096、THREADS256、VEC_ITEMS4。以下统计单位为Warp 级指令版本全局加载指令全局存储指令V5262,144131,072V6163,840131,072V6_DOUBLE65,53632,768V732,76832,768V6 第一遍使用 128 位加载V6_DOUBLE 的归一化读写也向量化V7 进一步去掉第二遍输入加载。指令数量、逻辑字节数、缓存请求与实际 DRAM 字节数应分开讨论。较早的 V6_DOUBLE 资源报告中寄存器允许每 SM 驻留 6 个 Block而 Warp 上限只允许 4 个理论 occupancy 仍为 100%实际约 93.86%。不能直接把双遍版退化归因为寄存器限制降低了理论驻留量。SASS 和源码采样定位到输出阶段加载后使用的依赖热点next 提前加载在预取版得到保留。不过采样次数不是精确延迟cycles/issued instruction 的分母随版本改变不能直接当成总等待时间的增减。双遍版退化的全部原因尚未隔离。7.6 V5 与 V8更低的 occupancy更少的总指令配置M1025、N257、THREADS256V5 的 Grid1025V8 的 Grid129。两版理论 occupancy 均为 100%。指标V5V8寄存器 / 线程2034静态共享内存 / Block272 B0 BAchieved Occupancy87.20%63.40%Active Warps / Scheduler7.055.65Eligible Warps / Scheduler1.651.24Issued Warp / Scheduler0.700.54Warp Cycles / Issued Instruction10.1110.54总 Warp 指令数2,442,575634,580全局加载指令数18,45018,450全局存储指令数9,2259,225实测事实V8 的 occupancy、Eligible Warps 和发射率均较低但总 Warp 指令数减少约74.02%全局加载、存储指令数则完全相同。一行 257 个元素每遍需要 ceil(257/32)9 组 Warp 访问。两遍加载对应 1025×9×218,450 次一遍存储对应 1025×99,225 次。思考题为什么发射率更低整体仍然更快请联系“每周期执行多少工作”和“完成任务总共需要执行多少工作”解释。解释发射率反映单位周期内发射多少指令虽然V8的发射率相对较低但由Warp指令的减少可知每周期执行的工作也大幅减少因而完成任务的总耗时也可能减少。从源码数规约工作V5 每行由八个 Warp 各执行一次 Online 规约再由第一个 Warp 合并各 Warp 的结果V8 由一个 Warp 独立完成整行最后广播 m、d。每组规约有五轮每轮分别 Shuffle m、d。V5 每行规约组数9组V8 每行规约组数1组8. 实验结论与局限8.1 已有证据支持的结论已测试的版本与配置通过 CPU FP64 参考检查。基础版本覆盖全零、末尾特殊值、整行平移与随机输入后续版本按各自记录的尺寸验证。V8、V9 另在 M1025、N257 下加入 NaN 输出初始化检查数值与行覆盖。V2 使用 Warp Shuffle 减少规约中的共享内存读写与 Block 同步V3 复用输入与指数结果减少全局加载及指数重算。第 6.2 节 V3 最快的结论只针对早期 V1–V4_OPT 那一批测量不能当成当前全部版本的排名。Online V4 将 V2 的三遍输入扫描降为两遍但比 V3 多一次读取并增加递推与指数计算V4_OPT 共享指数后取得小幅改善。Online 与完整输入缓存体现了计算、访存和寄存器压力的取舍。后续实测中V6_DOUBLE 进一步减少访存指令却没有加速V7 缓存完整输入消除了第二遍全局加载在已测 N4096 配置下取得明确收益。V8 在已测短行上更快、长行上更慢V9 缩小 Grid 的短行实验未获收益。当前最直接的新证据M1025、N257 时V8 的总 Warp 指令减少约 74.02%全局读写指令数保持相同同时其 occupancy 与发射率均低于 V5。原因分析与个人推导保留在第 7.6 节。8.2 结论边界与尚未隔离的因素各版本改变的因素并非完全隔离V3 同时减少加载和指数重算V4_OPT 改变指数及条件处理V8 改变每行线程数、跨 Warp 通信和启动规模。因此不能把全部加速归到单个指标也不能将 NCU 的预计加速当成可实现收益。指令数不等于 DRAM 流量Warp 等待百分比不等于可直接减去的总耗时。本次未单独量化 V8 的 Barrier 等待贡献更长 Online 递推链对长行退化的贡献也尚未独立测定。测试覆盖有限尺寸、连续 FP32、有限数值的输入尚未完整覆盖 NaN/Inf、空行、非连续张量、FP16/BF16 和其他 GPU。通过已有数值检查或某版 memcheck不等于所有版本、所有输入均正确。普通 Event 与 NCU 的重放、缓存和频率条件不同采集期间程序打印的 Average Time 不用于正式性能比较。不同云主机或配置不完整的历史记录分别归档。计时仍按固定版本顺序执行。约 1%–2% 的小幅收益需要交换顺序、多次重复与多组输入继续验证。尚未与框架或成熟 Softmax 系统比较不宣称最优。8.3 后续工作V5–V9 已完成本文记录用例的实现、数值验证与性能实验其中 V8 有调度和总指令证据V9 保留未获收益的 Grid 对照。后续继续按“实现 → 正确性验证 → Event 计时 → NCU/SASS 验证”逐项推进。探索短行 Warp 版与寄存器缓存结合单独保留对照目前待实验。系统扫描 M、N、线程数和 Grid寻找按形状选择方案的依据目前没有通用切换阈值。统一记录版本、编译参数、寄存器、stack/spill、输入分布和原始报告并加入成熟实现作为基线。8.4 实验复盘思考题用自己的话总结本轮最重要的两项性能取舍是什么解释第一是缓存输入与寄存器资源之间的取舍。V7 保留完整输入消除第二遍全局加载在已测 N4096 配置下相对同批 V6_DOUBLE 耗时减少约 26.55%但缓存容量随每线程数据量增加需要检查寄存器、Local Memory 和 spill不能无限增加。第二是每行并行度与规约通信之间的取舍。V8 减少每行参与的线程和规约工作在短行上表现更好长行却让每线程递推链明显增长。