为什么在 ROCm 7.x 时代要关注 Triton
如果你最近开始在 AMD Instinct GPU 上折腾大模型，大概率会听到两个词：一个是 ROCm 7.x，另一个就是 Triton。以前大家聊 AMD 加速，总绕不开“手写 HIP C++"这道高门槛——不仅要懂 GPU 架构，还得跟各种指针、内存布局死磕，稍有不慎就是 Segfault。但现在情况变了，随着 ROCm 7.x 的成熟，Triton 编译器在 AMD 平台上的支持已经从“实验性”迈向了“生产可用”。
对于关注底层优化的开发者来说，这绝对是个好消息。Triton 不再只是 NVIDIA 生态的专属玩具，它正在成为连接 PyTorch 高层逻辑与 AMD 硬件底层算力的关键桥梁。今天我就结合最近的实战体验，聊聊怎么用 Triton 在 ROCm 7.x 环境下开发自定义 Kernel，顺便给一段能跑通的矩阵乘法代码，帮你省去那些重复造轮子的时间。
Triton 如何替代手写 HIP 代码
在传统的 AMD GPU 开发流程里，想要优化一个特定算子（比如某种特殊的 Attention 变体），通常得走这条路：写 HIP C++ 代码 -> 手动管理 Shared Memory -> 处理 Warp 级别的同步 -> 编译链接 -> 调试。这个过程不仅耗时，而且极易出错，尤其是当硬件架构从 gfx90a 升级到 gfx942（MI300 系列）时，很多底层的调优参数都得重新摸索。
Triton 的出现把这个问题简化了。它允许你用类似 Python 的语法描述并行计算逻辑，编译器会自动帮你处理分块（Blocking）、预取（Prefetching）以及寄存器分配。在 ROCm 7.x 版本中，Triton 的后端已经能够正确识别 AMD 的架构特性，生成高效的机器码。这意味着你不需要再去纠结hipLaunchKernel的具体参数，也不用担心 Shared Memory 的大小限制，只需专注于算法逻辑本身。
更重要的是，Triton 生成的 Kernel 可以直接被 PyTorch 调用。你在前端用 PyTorch 写模型结构，遇到性能瓶颈的算子直接用 Triton 重写，两者无缝衔接。这种“高层灵活 + 底层高效”的模式，特别适合那些需要快速迭代算法的研究团队，或者想要在不修改主框架的前提下提升推理速度的工程团队。
实战：用 Triton 编写矩阵乘法 Kernel
光说不练假把式。下面这段代码展示了一个基础的矩阵乘法（MatMul）Kernel，专门针对 AMD GPU 进行了适配。这段代码可以在安装了 ROCm 7.x 和对应版本 Triton 的环境中直接运行。
importtorchimporttritonimporttriton.languageastl@triton.jitdefmatmul_kernel(a_ptr,b_ptr,c_ptr,M,N,K,stride_am,stride_ak,stride_bk,stride_bn,stride_cm,stride_cn,BLOCK_SIZE_M:tl.constexpr,BLOCK_SIZE_N:tl.constexpr,BLOCK_SIZE_K:tl.constexpr,GROUP_SIZE_M:tl.constexpr,):pid=tl.program_id(axis=0)num_pid_m=tl.cdiv(M,BLOCK_SIZE_M)num_pid_n=tl.cdiv(N,BLOCK_SIZE_N)num_pid_in_group=GROUP_SIZE_M*num_pid_n group_id=pid//num_pid_in_group first_pid_m=group_id*GROUP_SIZE_M group_size_m=min(num_pid_m-first_pid_m,GROUP_SIZE_M)pid_m=first_pid_m+(pid%group_size_m)pid_n=(pid%num_pid_in_group)//group_size_m offs_am=(pid_m*BLOCK_SIZE_M+tl.arange(0,BLOCK_SIZE_M))%M offs_bn=(pid_n*BLOCK_SIZE_N+tl.arange(0,BLOCK_SIZE_N))%N offs_k=tl.arange(0,BLOCK_SIZE_K)a_ptrs=a_ptr+(offs_am[:,None]*stride_am+offs_k[None,:]*stride_ak)b_ptrs=b_ptr+(offs_k[:,None]*stride_bk+offs_bn[None,:]*stride_bn)accumulator=tl.zeros((BLOCK_SIZE_M,BLOCK_SIZE_N),dtype=tl.float32)forkinrange(0,tl.cdiv(K,BLOCK_SIZE_K)):a=tl.load(a_ptrs,mask=offs_k[None,:]<K-k*BLOCK_SIZE_K,other=0.0)b=tl.load(b_ptrs,mask=offs_k[:,None]<K-k*BLOCK_SIZE_K,other=0.0)accumulator+=tl.dot(a,b)a_ptrs+=BLOCK_SIZE_K*stride_ak b_ptrs+=BLOCK_SIZE_K*stride_bk c_ptrs=c_ptr+stride_cm*offs_am[:,None]+stride_cn*offs_bn[None,:]c_mask=(offs_am[:,None]<M)&(offs_bn[None,:]<N)tl.store(c_ptrs,accumulator,mask=c_mask)defmatmul(a,b):asserta.shape[1]==b.shape[0],"Incompatible dimensions"asserta.is_contiguous(),"Matrix A must be contiguous"assertb.is_contiguous(),"Matrix B must be contiguous"M,K=a.shape K,N=b.shape c=torch.empty((M,N),device=a.device,dtype=torch.float32)# 配置 Grid 和 Block 大小，针对 MI300 系列可适当调大 BLOCK_SIZEBLOCK_SIZE_M,BLOCK_SIZE_N,BLOCK_SIZE_K=64,64,32grid=(triton.cdiv(M,BLOCK_SIZE_M)*triton.cdiv(N,BLOCK_SIZE_N),)matmul_kernel[grid](a,b,c,M,N,K,a.stride(0),a.stride(1),b.stride(0),b.stride(1),c.stride(0),c.stride(1),BLOCK_SIZE_M,BLOCK_SIZE_N,BLOCK_SIZE_K,GROUP_SIZE_M=8)returnc# 测试运行if__name__=="__main__":# 确保在 AMD GPU 环境下运行device=torch.device("cuda"iftorch.cuda.is_available()else"cpu")# 注意：在 ROCm 中 torch.cuda 通常兼容，具体视版本而定a=torch.randn(1024,1024,device=device,dtype=torch.float16)b=torch.randn(1024,1024,device=device,dtype=torch.float16)triton_output=matmul(a,b)torch_output=torch.matmul(a,b).to(torch.float32)print(f"Triton 输出最大误差：{torch.max(torch.abs(triton_output-torch_output))}")这段代码的核心在于matmul_kernel函数。你可以看到，我们没有显式地分配 Shared Memory，也没有写复杂的线程索引计算，Triton 编译器会自动将这些逻辑映射到 AMD GPU 的硬件资源上。在 ROCm 7.x 环境下，只要设置好PYTORCH_ROCM_ARCH环境变量（例如gfx942），这段代码就能编译并通过验证。实测在 MI300X 上，对于中等规模的矩阵，其性能已经非常接近手写 HIP 的水平，但开发效率却提升了数倍。
优化潜力与落地建议
当然，Triton 在 ROCm 上的应用不仅仅是跑通一个 MatMul。对于大模型推理中的关键算子，如 FlashAttention 的变体、自定义的量化反量化逻辑，Triton 都提供了极大的优化空间。特别是在处理非标准形状或非标准精度的运算时，手写 CUDA/HIP 往往成本过高，而 Triton 能让你在几天内就完成原型的验证和部署。
不过，目前仍有一些细节需要注意。首先是版本匹配，Triton 的 ROCm 分支更新较快，务必确保其与你的 PyTorch 及 ROCm 驱动版本兼容。其次，虽然编译器自动化程度很高，但在极端性能要求下，手动调整BLOCK_SIZE等参数依然能带来显著收益。建议大家在开发时，多参考 Github 上活跃的 Triton ROCm 相关 Issue，社区里有很多关于特定架构调优的实战讨论。
总的来说，Triton 正在让 AMD GPU 的底层开发变得前所未有的友好。如果你之前因为 HIP 的学习曲线而犹豫是否深入 AMD 生态，现在或许是个重新尝试的好时机。不用再去啃晦涩的底层文档，用熟悉的 Python 思维去挖掘硬件潜力，这才是技术演进该有的样子。
200小时GPU算力已就位，快来领取：https://marketing.csdn.net/questions/Q2604140858304426315?utm_source=AIpaper
