1. 项目概述一场被误读的“技术碾压”以及我们真正该学的底层优化思维你点开这篇文章大概率是被标题里那个极具冲击力的短语击中了——“How Deepseek Destroyed OpenAI”。这确实是个抓眼球的好标题但作为在GPU高性能计算领域摸爬滚打十多年、亲手调过上万行CUDA内核、在A100和H100集群上熬过无数个通宵的从业者我得先给你泼一盆冷静的水DeepSeek没有、也不可能“摧毁”OpenAI。这两家机构根本不在同一个竞争维度上。OpenAI是通用大模型能力与产品生态的集大成者而DeepSeek是专注在模型架构、训练效率与推理优化上打出差异化优势的技术型选手。把它们放在“摧毁”这个语境里就像拿一个顶级赛车引擎工程师去对比一家F1车队的总冠军——前者造出了史上热效率最高的V6涡轮增压引擎后者赢下了蒙扎赛道。他们都在推动边界但路径完全不同。那么标题里真正值得我们掏心掏肺去拆解的其实是后半句“and How You Can Do it Too!”。这才是这篇博文存在的全部意义。它指向的不是一场虚构的商业战争而是一套真实存在、可被复现、能立竿见影提升你代码性能的硬核技术路径CUDA PTX与内联汇编Inline ASM的深度协同优化。这不是什么玄学而是NVIDIA GPU计算栈里最锋利、也最常被忽视的一把手术刀。当你用PyTorch写完一个模型torch.compile帮你做了图优化当你用triton写了一个kernel它自动帮你做了shared memory调度但当你需要把一块已经高度优化的kernel再榨出5%、10%甚至20%的峰值算力时你就必须亲手握住PTX这把刀去雕琢每一个warp的指令发射、每一条寄存器的生命周期、每一次global memory的访存模式。我试过在一个关键的attention softmax归一化kernel里把一段循环展开手动向量化寄存器重用的逻辑从高级C降级到PTX手写最终在A100上将延迟从8.7ms压到了6.9ms吞吐量提升了26%。这个数字背后没有魔法只有一行行对着nvdisasm反汇编出来的SASS指令逐条比对、逐条修改的枯燥工作。所以这篇文章的核心关键词绝不是“DeepSeek”或“OpenAI”而是PTX、Inline ASM、CUDA Kernel Optimization、GPU Microarchitecture、Performance Bottleneck Analysis。它适合三类人第一类是正在用CUDA写高性能计算、科学仿真或AI推理服务的工程师你们正卡在“为什么我的kernel跑不满GPU的TFLOPS”的瓶颈上第二类是算法研究员你们的模型结构创新已经触顶下一步想从系统侧打开新空间第三类是刚入门的GPU开发者你们需要知道除了__global__函数和cudaMemcpyGPU编程的深水区究竟长什么样。接下来的内容不会教你如何“毁灭”谁但会手把手带你把你的CUDA kernel从“能跑”变成“飞起来”。2. 内容整体设计与思路拆解为什么是PTX/ASM而不是别的2.1 高层抽象的甜蜜陷阱与性能悬崖在开始讲PTX之前我们必须先直面一个残酷的现实现代GPU编程框架无论是CUDA C、HIP还是更上层的Triton、CuBLAS都建立在一个巨大的、精妙的抽象金字塔之上。这个金字塔的顶端是让你写几行Python就能启动一个分布式训练的torch.distributed中间层是让你用cuda.jit就能定义一个kernel的Numba而最底层是nvcc编译器为你生成的、面向特定GPU架构如sm_80, sm_90的二进制机器码SASS。这个抽象带来了无与伦比的开发效率但也埋下了一个巨大的隐患抽象泄漏Abstraction Leakage。举个最典型的例子。你在CUDA C里写了一个简单的向量加法__global__ void vec_add(float* a, float* b, float* c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) c[idx] a[idx] b[idx]; }这段代码nvcc会把它编译成什么它会生成一个包含大量ld.global.f32加载全局内存、add.f32浮点加法、st.global.f32存储全局内存指令的SASS序列。但问题在于nvcc的优化器是一个基于统计模型和启发式规则的“通用”优化器。它不知道你的a和b数组在物理内存上是否连续、是否对齐、是否会被其他kernel频繁访问。它只能按照一套预设的、适用于“大多数情况”的规则来调度指令。当你的场景足够特殊——比如你的数据是16字节对齐的你的kernel是warp-level同步的你的访存模式是完全可预测的——这套通用规则就会成为你性能的天花板。我亲眼见过一个客户他们的kernel在V100上跑得飞快但迁移到A100后性能反而掉了30%。原因就是nvcc为V100生成的指令序列在A100的全新Tensor Core和L2 cache架构下触发了更严重的bank conflict。而这个问题nvcc自己无法感知也无法修复。这就是PTX存在的根本逻辑它是一个可控的、可审计的、位于高级语言与硬件机器码之间的“可信中间层”。PTX不是汇编它是一种虚拟ISAInstruction Set Architecture由NVIDIA定义独立于具体的GPU硬件。你可以把它想象成Java的字节码——javac把Java源码编译成.class文件JVM再把字节码解释或JIT编译成x86或ARM指令。同理nvcc把CUDA C编译成.ptx文件而GPU驱动里的JIT编译器再把PTX编译成针对当前GPU的SASS。这个设计带来了两个不可替代的价值可移植性和可控性。可移植性意味着一份为compute_70Volta编写的PTX代码可以在compute_80Ampere甚至compute_90Hopper的GPU上运行驱动会负责做必要的适配。而可控性则意味着你终于可以绕过nvcc那个“黑盒”优化器亲手去编写、调试、验证每一行指令的行为。2.2 Inline ASM在PTX的土壤上种下最锋利的刺如果说PTX是那片可供耕作的土地那么Inline ASM就是你手里那把最锋利的锄头。CUDA C标准允许你在kernel内部用asm volatile语法直接嵌入PTX指令。这听起来很危险也确实如此——它绕过了C编译器的所有类型检查、寄存器分配和依赖分析。但正是这种“危险”赋予了它无与伦比的力量。我们来看一个教科书级的例子Warp-level reduction线程块内规约。这是几乎所有并行算法的核心操作比如求和、求最大值、softmax归一化。一个标准的、教科书式的实现是用共享内存shared memory做两级规约// 标准shared memory reduction __shared__ float sdata[256]; int tid threadIdx.x; sdata[tid] input[tid]; __syncthreads(); for (int s blockDim.x / 2; s 0; s 1) { if (tid s) sdata[tid] sdata[tid s]; __syncthreads(); } if (tid 0) output[blockIdx.x] sdata[0];这段代码清晰、安全、易于理解。但它有一个致命的弱点同步开销。每一次__syncthreads()都会让整个warp停下来等待而GPU的warp scheduler最怕的就是这种停顿。在A100上一次__syncthreads()的开销可能相当于执行几十条ALU指令。而Inline ASM的解决方案是彻底抛弃共享内存转而利用warp内部的shuffle指令shfl_down.sync。这个指令允许一个线程直接从同一warp内另一个线程的寄存器里“偷”数据全程无需任何内存访问也无需同步。nvcc当然也支持__shfl_down_sync这个内置函数但它的生成逻辑依然是“保守”的。而Inline ASM让你可以精确控制shuffle的mask、predicate、甚至指令的发射顺序。下面是一段真实的、用于float32求和的Inline ASM核心片段float sum val; #pragma unroll for (int offset 16; offset 0; offset / 2) { asm volatile ( shfl.down.b32 %0, %0, %1, 0x1f; : r(sum) : r(offset) : cc ); }这里shfl.down.b32是PTX指令%0代表输出操作数sum%1代表偏移量offset0x1f是warp mask31即所有32个线程。r表示这是一个输入输出寄存器约束cc表示这条指令会修改条件码寄存器。这段代码编译后会生成一条极简的SASS指令没有任何分支、没有内存访问、没有同步。实测下来在32线程的warp内做一次float32求和它比共享内存版本快了整整40%且功耗更低。这就是Inline ASM的价值它不是为了炫技而是为了在那些nvcc无法触及的、最细微的微架构缝隙里塞进最高效的指令流。2.3 方案选型的终极考量何时该拿起这把刀看到这里你可能会问既然这么强大那我是不是该把所有kernel都重写成PTX答案是绝对否定的。这把刀锋利但也沉重。我的经验是只有当同时满足以下三个条件时才值得投入PTX/ASM的开发性能是生死线你的kernel是端到端pipeline的瓶颈其执行时间占总耗时的30%以上且经过nvprof或Nsight Compute分析确认瓶颈在指令级如IPC低、ALU Utilization低、L1/Tensor Cache Miss Rate高而非内存带宽或PCIe瓶颈。模式高度固定你的数据规模、访存模式、计算逻辑是已知且稳定的。PTX代码不具备C的泛型能力一个为n1024优化的PTX kernel很可能在n1025时就因边界检查失败而崩溃。团队具备微架构知识你和你的团队必须熟悉目标GPU的微架构文档如NVIDIA的Ampere Architecture Whitepaper知道什么是warp divergence、什么是memory coalescing、什么是tensor core的指令格式。没有这些知识PTX对你而言不是工具而是地雷。我见过太多团队为了追求“极致性能”一头扎进PTX结果花了三个月只把一个kernel的性能提升了2%却拖垮了整个项目的迭代速度。这完全违背了工程的本质。PTX/ASM永远是你的“最后一招”是当你已经榨干了所有高级优化手段算法重构、内存布局调整、__restrict__关键字、#pragma unroll之后才亮出的底牌。它解决的不是“能不能跑”的问题而是“能不能跑得更快、更省、更稳”的问题。3. 核心细节解析与实操要点从理论到落地的完整链路3.1 PTX语言全景不只是“GPU汇编”而是一套精巧的虚拟机要驾驭PTX首先得理解它到底是什么。很多初学者把它简单等同于“GPU的汇编语言”这是一个危险的误解。真正的GPU机器码SASS是高度硬件相关的sm_80的指令集和sm_90的指令集差异巨大直接写SASS等于给自己判了死刑。而PTX是NVIDIA精心设计的一套虚拟指令集架构Virtual ISA它的设计哲学是“一次编写处处运行”。PTX的指令集被严格划分为几个功能域每个域都有其明确的语义和硬件映射ALU指令add.f32,mul.f32,fma.rn.f32融合乘加等负责基本的算术运算。注意fma指令它在现代GPU上是单周期完成的比muladd两步快得多且精度更高。Memory指令ld.global.f32,st.global.f32,ld.shared.f32,st.shared.f32等负责不同地址空间的加载/存储。这里的global、shared、local对应着GPU的物理内存层级。Control Flow指令bra,call,ret,exit等负责程序跳转。PTX的分支预测非常弱因此bra指令的代价很高应尽量避免。Warp Shuffle指令shfl.down.b32,shfl.up.b32,shfl.xor.b32等这是PTX区别于传统CPU汇编的最大特色专为warp级通信设计。Predicate指令p1 add.f32其中p1是一个谓词predicate只有当谓词为真时后面的指令才会执行。这是实现warp内条件分支warp divergence的底层机制。PTX的寄存器模型也与CPU截然不同。它没有rax,rbx这样的物理寄存器概念而是使用%r1,%r2,%f1,%f2这样的虚拟寄存器。nvcc的后端编译器ptxas会负责把这些虚拟寄存器映射到GPU真实的32位通用寄存器GPR上。这意味着你在PTX里写的%r1在SASS里可能是R4也可能是R128完全取决于ptxas的寄存器分配策略。这也是为什么PTX代码本身是可移植的——它不关心底层物理寄存器的编号。一个完整的PTX文件其结构非常清晰。以下是一个简化版的、用于向量加法的PTX代码vec_add.ptx.version 7.8 .target sm_80 .address_size 64 .visible .entry _Z8vec_addPfS_S_i( .param .u64 _Z8vec_addPfS_S_i_param_0, .param .u64 _Z8vec_addPfS_S_i_param_1, .param .u64 _Z8vec_addPfS_S_i_param_2, .param .u32 _Z8vec_addPfS_S_i_param_3 ) { .reg .f32 %f10; .reg .u32 %r10; .reg .u64 %rd10; ld.param.u64 %rd1, [_Z8vec_addPfS_S_i_param_0]; ld.param.u64 %rd2, [_Z8vec_addPfS_S_i_param_1]; ld.param.u64 %rd3, [_Z8vec_addPfS_S_i_param_2]; ld.param.u32 %r1, [_Z8vec_addPfS_S_i_param_3]; cvt.s32.u32 %r2, %r1; mov.s32 %r3, %tid.x; mul.wide.s32 %rd4, %r3, 4; add.s64 %rd5, %rd1, %rd4; add.s64 %rd6, %rd2, %rd4; add.s64 %rd7, %rd3, %rd4; ld.global.f32 %f1, [%rd5]; ld.global.f32 %f2, [%rd6]; add.f32 %f3, %f1, %f2; st.global.f32 [%rd7], %f3; }这段代码展示了PTX的典型结构.version声明PTX版本.target声明目标架构.entry定义入口函数.reg声明寄存器然后是具体的指令序列。注意%tid.x这个特殊寄存器它代表当前线程在block内的x维索引是PTX提供的硬件内置变量无需你手动计算。这个例子虽然简单但它揭示了PTX的全部力量你不再需要blockIdx和threadIdx的复杂计算%tid.x直接给你最原始的线程ID你也不再需要cudaMemcpy的抽象ld.global.f32和st.global.f32直接操作内存地址。这是一种回归本质的编程体验。3.2 Inline ASM的语法精要在C的躯壳里注入PTX的灵魂在CUDA C中嵌入PTX是通过asm volatile内联汇编语法实现的。它的语法结构比GCC的内联汇编更为简洁但也更易出错。一个标准的asm volatile语句由四部分组成asm volatile ( PTX指令模板字符串 : 输出操作数列表 : 输入操作数列表 : 破坏列表clobber list );PTX指令模板字符串这是核心里面可以包含占位符%0,%1,%2...分别对应输入/输出操作数列表中的第0、1、2...个操作数。字符串里还可以包含直接的PTX指令比如add.f32 %0, %1, %2;。输出操作数列表用括号()包裹格式为r(variable)。表示这是一个输出r表示使用任意通用寄存器variable是C变量名。r则表示这是一个输入输出操作数即该变量既提供初始值又接收计算结果。输入操作数列表格式为r(variable)r表示输入variable是C变量名。你也可以用f表示浮点寄存器l表示64位寄存器。破坏列表告诉编译器这条内联汇编会“破坏”即修改哪些寄存器或状态。最常见的是cccondition code条件码和r1明确指定破坏某个寄存器。如果你不写编译器会假设你破坏了所有寄存器导致它生成大量冗余的保存/恢复指令严重拖慢性能。下面我们用一个更复杂的、实际生产环境中的例子来展示这四部分是如何协同工作的。这是一个用于计算exp(x)的快速近似函数它利用了GPU的ex2.approx.f32指令一个硬件级的指数近似比软件库快10倍__device__ __forceinline__ float fast_expf(float x) { // 将x分解为整数部分i和小数部分f: x i f, 其中f in [0,1) int i; float f modff(x, i); // 计算2^i * 2^f因为exp(x) e^x 2^(x * log2(e)) ≈ 2^(x * 1.442695) float scale 1.442695f * x; int i_scale; float f_scale modff(scale, i_scale); // 使用硬件指令计算2^f_scale float exp_f; asm volatile ( ex2.approx.f32 %0, %1; : f(exp_f) : f(f_scale) : cc ); // 将2^i_scale转换为浮点数这需要查表或位操作此处略 float pow2_i __int_as_float(i_scale 127) * 0.5f; // 简化版实际需更精确 return exp_f * pow2_i; }在这个例子里ex2.approx.f32是PTX指令%0和%1分别对应输出变量exp_f和输入变量f_scale。f表示输出到一个浮点寄存器f表示输入来自一个浮点寄存器。cc是必须的因为ex2.approx.f32会修改条件码。这个函数在我们的一个实时语音识别pipeline中将logits的softmax前计算耗时从1.2ms降到了0.3ms效果立竿见影。提示Inline ASM的调试是最大的痛点。nvcc不会像C编译器那样给你报出“undefined symbol”或“type mismatch”的错误。它只会静默地编译失败或者在运行时产生illegal instruction的CUDA error。因此我强烈建议任何Inline ASM代码都必须配合nvdisasm工具进行反汇编验证。编译你的CUDA程序后用nvdisasm -c your_program.cubin | grep -A 10 your_kernel_name查看生成的SASS指令确保你的PTX真的被编译成了你想要的那条硬件指令。3.3 实战避坑指南那些只有踩过才知道的“天坑”PTX/ASM的世界充满了优雅的数学和精妙的逻辑但也遍布着只有亲身趟过才能发现的“天坑”。以下是我在过去十年里用无数个崩溃的kernel和core dump换来的血泪教训坑一寄存器溢出Register Spilling的无声杀手这是最隐蔽、也最致命的坑。当你在Inline ASM里声明了过多的%r或%f寄存器而GPU的物理寄存器池通常是255个32位GPR不够用时ptxas编译器会自动启用“寄存器溢出”机制它会把一部分寄存器的值临时存放到local memory即GPU的stack里。Local memory是DRAM访问延迟是寄存器的上百倍。结果就是你精心优化的PTX kernel性能不升反降而且下降得毫无征兆。如何发现用Nsight Computeprofiling看Stall Memory Throttle和Stall Instruction Fetch这两个指标。如果它们异常高十有八九就是寄存器溢出了。解决方案在nvcc编译时强制指定寄存器数量例如nvcc -Xptxas -v -maxrregcount64 your_code.cu-v参数会打印出详细的寄存器使用报告-maxrregcount64则强制每个线程最多使用64个寄存器逼迫你写出更紧凑的代码。坑二warp divergence的“幽灵”PTX指令本身是warp级执行的但Inline ASM嵌入在C kernel里而C的if/else、for循环在PTX层面依然会产生warp divergence。一个常见的错误是在Inline ASM外面写了一个if (threadIdx.x n)然后在里面写了一堆PTX指令。你以为只有满足条件的线程在执行但实际上warp内的所有32个线程都在执行相同的PTX指令流只是不满足条件的线程其结果被谓词predicate屏蔽了。这会导致大量的“空转”指令。正确的做法是把整个逻辑包括边界检查都用PTX的谓词指令来实现。例如用p1 ld.global.f32 %f1, [%rd5];其中p1是一个根据threadIdx.x n计算出来的谓词。坑三PTX版本与GPU架构的“代沟”PTX是一个向前兼容的虚拟ISA但不是向后兼容的。一个为ptx 7.8对应CUDA 11.8编写的PTX文件可以在sm_80A100和sm_90H100上运行但如果你试图在sm_70V100上运行就会失败因为sm_70的驱动不支持ptx 7.8的新指令。因此你的PTX代码必须声明一个足够老、但又能支持你所需指令的.target。我的经验是除非你明确需要sm_90的全新指令如mma.sync.aligned.m16n8k16.row.col.f32否则sm_80是一个完美的平衡点它覆盖了A100和RTX 3090等主流卡且指令集足够丰富。4. 实操过程与核心环节实现一个完整的、可运行的PTX优化案例4.1 案例背景一个被低估的性能瓶颈——Batched GEMM的尾部处理在AI推理服务中batched GEMM批量矩阵乘是一个无处不在的核心算子。PyTorch的torch.bmm、TensorRT的IConstantLayer背后都是它。然而几乎所有框架的batched GEMM实现都默认假设batch size是32的整数倍。当你的实际batch size是31、33、65时框架会自动向上取整到32或64然后用padding填充零来凑满。这看似无害但padding带来的后果是灾难性的它引入了大量无效的、但依然要执行的计算浪费了宝贵的GPU算力。在一个为边缘设备优化的语音唤醒模型中我们发现batch_size1的请求其batched GEMM的耗时竟然是batch_size32的1.8倍原因就是框架为了兼容把1x128x128的矩阵强行扩展成了32x128x128然后执行了32次计算最后只取第一个结果。这就是我们选择的优化战场为任意batch size尤其是小batch定制一个零padding的、纯PTX实现的batched GEMMkernel。我们的目标很明确对于batch_size1的请求性能必须优于任何padding方案对于batch_size32性能不能低于cuBLAS。4.2 设计蓝图从算法到PTX指令的逐层映射一个标准的GEMM计算C A * B其中A是MxK矩阵B是KxN矩阵C是MxN矩阵。batched GEMM则是对batch_size个这样的三元组同时进行计算。我们的PTX kernel将采用经典的分块tiling warp-level shared memory策略但所有分块大小、循环展开、寄存器分配都将由PTX指令直接控制。核心设计思想如下Warp Mapping每个warp32线程负责计算C矩阵的一个16x16的tile。这意味着一个128x128的C矩阵需要8x864个warp来完成。Shared Memory Tiling我们将A和B的子块预先加载到shared memory中以减少global memory的访问次数。每个warp会加载一个16x8的A_tile和一个8x16的B_tile到shared memory。Register Tiling在计算C_tile时我们将C_tile的16x16256个元素全部存放在warp的32个线程的寄存器中。每个线程负责4个元素256/328但考虑到寄存器压力我们让每个线程负责4个共128个其余由相邻线程分担。这要求我们精确计算每个线程的寄存器需求并用-maxrregcount进行约束。PTX指令选择核心计算将使用fma.rn.f32融合乘加指令因为它在一个周期内完成a*bc且精度符合IEEE 754。访存将使用ld.shared.f32和st.shared.f32以获得最佳的shared memory带宽。4.3 完整代码实现与逐行注释下面是这个batched GEMMkernel的完整PTX实现为简洁起见省略了错误检查和边界处理实际生产环境必须加上// batched_gemm_ptx.cuh #include cuda_runtime.h #include cuda.h // 声明一个外部的PTX kernel其入口点名为batched_gemm_kernel extern C const char* batched_gemm_ptx; // 这是一个包装函数用于从host端调用PTX kernel __host__ cudaError_t launch_batched_gemm_ptx( float* d_A, float* d_B, float* d_C, int batch_size, int M, int N, int K, int lda, int ldb, int ldc ) { // 计算grid和block尺寸 dim3 block(256); // 一个block包含256个线程即8个warp dim3 grid((M 15) / 16, (N 15) / 16, batch_size); // 每个C矩阵的tile对应一个block // 获取PTX模块 CUmodule module; CUfunction function; cuModuleLoadData(module, batched_gemm_ptx); cuModuleGetFunction(function, module, batched_gemm_kernel); // 设置kernel参数 void* args[] {d_A, d_B, d_C, batch_size, M, N, K, lda, ldb, ldc}; cuLaunchKernel(function, grid.x, grid.y, grid.z, block.x, block.y, block.z, 0, 0, args, 0); return cudaSuccess; }而真正的灵魂在于batched_gemm_ptx这个字符串常量里。它是一个巨大的、由nvcc编译生成的PTX代码字符串。由于篇幅限制我无法在此处粘贴全部数千行代码但我将为你展示其中最关键、最体现PTX威力的核心计算循环部分并附上详尽的注释// 这是PTX代码中计算一个16x16 C_tile的核心循环 // 它被展开unrolled了8次以隐藏指令延迟 // 每次迭代计算C_tile中的一列16个元素 // 初始化C_tile的16个寄存器每个寄存器对应C_tile的一行的一个元素 .reg .f32 %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7; .reg .f32 %c8, %c9, %c10, %c11, %c12, %c13, %c14, %c15; mov.f32 %c0, 0.0; mov.f32 %c1, 0.0; // ... 初始化所有%c0-%c15 ... // 外层循环遍历K维度每次处理8行因为A_tile是16x8 .set K_ITER, 0; .repeat K_LOOP, 8; // 假设K128所以需要16次迭代这里简化为8次 // 加载A_tile的一列16个元素到寄存器 .reg .f32 %a0, %a1, %a2, %a3, %a4, %a5, %a6, %a7; .reg .f32 %a8, %a9, %a10, %a11, %a12, %a13, %a14, %a15; ld.shared.f32 %a0, [%rs_a 0]; ld.shared.f32 %a1, [%rs_a 4]; // ... 加载所有16个元素 ... // 加载B_tile的一行8个元素到寄存器 .reg .f32 %b0, %b1, %b2, %b3, %b4, %b5, %b6, %b7; ld.shared.f32 %b0, [%rs_b 0]; ld.shared.f32 %b1, [%rs_b 4]; // ... 加载所有8个元素 ... // 执行16x8的FMA计算C_row[i] A_col[j] * B_row[k] // 这里是8次展开的FMA每次处理B_row的一个元素 fma.rn.f32 %c0, %a0, %b0, %c0; fma.rn.f32 %c1, %a1, %b0, %c1; fma.rn.f32 %c2, %a2, %b0, %c2; // ... 对%c0-%c15全部执行一遍 ... fma.rn.f32 %c0, %a0, %b1, %c0; fma.rn.f32 %c1, %a1, %b1, %c1; // ... 对b1, b2, ..., b7全部执行一遍 ... // 更新shared memory指针为下一次迭代做准备 add.s32 %rs_a, %rs_a, 64