Triton Gluon AMD CDNA3 目标内建 API 全解buffer_load/buffer_store/buffer_atomic_* 与 mfma【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton导读本文围绕 Triton 仓库中 docs/gluon/api/amd.cdna3.rst 所定义的 AMD CDNA 3 目标特定 Gluon 内建 API系统讲解buffer_load、buffer_store、七个buffer_atomic_*原子操作以及矩阵乘核心mfma的语义、参数约束与底层实现。阅读本文后你将能够在 Gluon 编程模型下直接使用 CDNA 3gfx942如 MI300 系列的硬件原生 buffer 指令与矩阵核心写出绕过指针张量、以标量基址 偏移张量为访存模型的高性能内核。一、背景Gluon 的 AMD 目标特定 API 体系Gluon 是 Triton 中面向显式布局与硬件内建编程的新一代内核编写模型实验特性。它的语言层按硬件目标划分命名空间AMD 侧的入口模块为triton.experimental.gluon.language.amd由 python/triton/experimental/gluon/language/amd/init.py 导出结构如下通用 APIAMDMFMALayout、AMDWMMALayout、slice、warp_pipeline_stage、get_scaled_upcast_fp4_scale_layout代际子模块cdna3、cdna4、cdna5、rdna3、rdna4、gfx1250。文档 docs/gluon/api/amd.rst 中把cdna3与cdna4、cdna5、rdna3、rdna4并列挂接在 GPU Generations 目录下而本文主角amd.cdna3对应 CDNA 3 代际即AMDMFMALayout.version 3对应的gfx942架构见 python/triton/experimental/gluon/language/amd/_layouts.py。API 参考入口可追溯到 docs/gluon/api/index.rst 中的 AMD Intrinsics 一节。二、CDNA3 模块导出的完整内建函数清单根据 docs/gluon/api/amd.cdna3.rst 的 autosummary 列表triton.experimental.gluon.language.amd.cdna3模块文档化导出的内建函数共 10 个类别内建函数全局内存加载buffer_load全局内存存储buffer_store原子 RMW读改写buffer_atomic_add、buffer_atomic_and、buffer_atomic_max、buffer_atomic_min、buffer_atomic_or、buffer_atomic_xchg、buffer_atomic_xor矩阵核心mfma实际源码模块 python/triton/experimental/gluon/language/amd/cdna3/init.py 的__all__中还包括scaled_upcast软件模拟的 FP4/FP8 升精度缩放路径以及未在 autosummary 中列出的内部辅助函数_buffer_atomic_rmw_impl等。下文逐一展开这 10 个文档化 API 的用法与实现细节。三、buffer_load以标量基址 偏移张量加载全局内存3.1 签名与语义buffer_load(ptr, offsets, maskNone, otherNone, cacheNone)ptr指向标量的全局内存基址指针offsets偏移张量distributed_type元素类型必须为int32或uint32mask可选的谓词掩码张量用于按元素屏蔽加载other可选的填充值张量或标量为被mask屏蔽的元素提供默认值当other非空时必须同时提供maskcache可选的缓存修饰符字符串例如测试中出现的.ca、.cg、.cv。与普通load相比核心差异在于不构造指针张量而是通过一个标量基址 一个偏移张量直接生成 AMD 的amdg.buffer_load指令数据直接从全局内存载入寄存器。测试 python/test/gluon/test_frontend.py 中的buffer_load_store_kernel展示了典型用法并在HIP_TARGET_CDNA3目标下校验生成的 IRa ttgl.amd.cdna3.buffer_load(ptrx, offsetsoffsets, maskmask, otherother, cache.ca)对应的 MLIR 形如amdg.buffer_load %arg0[%2], %cst, %cst_1 {cachePolicy ...}印证了该内建直接映射到 AMDGPU buffer load 指令。3.2 实现要点实现位于 cdna3/init.py处理流程如下调用_verify_buffer_ops校验ptr为标量指针、offsets为分布式张量且元素类型为int32/uint32并检查other与mask的配对关系对mask与other执行广播broadcast_impl_valueother会先被 cast 成ptr指向的元素类型将cache字符串经_semantic._str_to_load_cache_modifier解析为CACHE_MODIFIER枚举缺省为NONE以offsets的形状为骨架、以ptr的元素类型为元素类型构造返回类型调用builder.create_buffer_load生成 IR 并包装为 Gluon 张量。四、buffer_store按偏移张量写回全局内存buffer_store(stored_value, ptr, offsets, maskNone, cacheNone)buffer_store是buffer_load的镜像操作把张量stored_value通过标量基址ptr与偏移张量offsets直接写入全局内存同样支持谓词mask与缓存修饰符cache测试中使用了.cs。实现见 cdna3/init.py其额外约束是广播后offsets的形状必须与调用前一致否则抛出ValueError防止存储时悄悄改变访存范围。五、buffer_atomic_*七个全局内存原子 RMW 操作5.1 公共语义CDNA3 模块提供 7 个原子读改写RMW内建add、and、max、min、or、xchg、xor。它们共享同一行为模型模块 docstring见 cdna3/init.py通过标量基址 偏移张量访问全局内存而非指针张量对ptr offsets处的元素执行op(value)并写回mask为布尔张量mask[i] 0的元素被跳过不执行原子操作返回全局内存中的操作前旧值sem指定内存语义描述符缺省为acq_relscope指定同步范围缺省映射到gpuAMDGPU 术语中称agent更细的 CDNA3gfx942内存模型语义可参考 LLVM AMDGPUUsage 文档中关于 gfx942 的 Memory Model 一节。统一签名7 个函数一致buffer_atomic_xxx(ptr, offsets, value, maskNone, semNone, scopeNone)统一实现路径为_buffer_atomic_rmw_impl(op, ptr, offsets, value, arch, mask, sem, scope, _semantic)见 cdna3/init.py校验参数后根据元素类型把字符串操作名映射为ir.ATOMIC_OP枚举经语义层把maskcast 为i1并广播最后调用builder.create_buffer_atomic_rmw。5.2 元素类型支持矩阵CDNA3 限定_verify_element_type_and_dispatch_opcdna3/init.py给出了硬编码的类型约束这是使用本模块前必须核对的关键事实操作支持元素类型映射的原子指令and/or/xor/xchg仅int32、int64AND/OR/XOR/XCHGmax/minint32、int64→ 有符号uint32、uint64→ 无符号SMAX/SMIN、UMAX/UMINfp16/bf16/fp32 不支持adduint32、uint64→ 整数加法float16、float32、float64→ 浮点加法ADDiadd、FADDfadd注意两点边界全部原子操作的基础类型白名单为float16/float32/bfloat16/float64/int32/int64/uint32/uint64其余类型一律断言失败bfloat16的fadd原子在 CDNA3 上不支持源码中明确断言 Buffer atomic fadd with bf16 is only supported on CDNA4 for now即该能力是 CDNA4 才引入的。六、mfma调用 AMD 原生矩阵核心mfma(a, b, acc)mfma计算a * b acc直接使用 AMD 原生矩阵核心Matrix Fused Multiply-Add单元。三个参数中acc为必填项缺失即断言失败返回张量的类型与acc一致。实现见 cdna3/init.py其本质是转发到语义层的_semantic.dot_semantic.dot(a, b, acc, input_precisionknobs.language.fp32_default, max_num_imprecise_accNone, out_dtypeacc.dtype)input_precision取自全局开关knobs.language.fp32_default即默认的 fp32 输入精度策略由 Triton 的 knob 体系统一控制max_num_imprecise_accNone表示不限制不精确累加的次数上限输出类型由acc.dtype决定因此累加器的类型即结果的类型。6.1 配套布局AMDMFMALayout要让mfma高效执行操作数的分布式布局必须与之匹配。CDNA3 对应的布局类AMDMFMALayoutpython/triton/experimental/gluon/language/amd/_layouts.py关键字段如下version架构代号CDNA3 即gfx942对应version 3version 1/2/4 分别对应 gfx908/gfx90a/gfx950instr_shape指令形状(M, N, K)合法 M×N 组合为[32,32]、[16,16]、[64,4]、[4,64]transposed结果是否转置使线程持有同行的连续元素利于链式 dot 与全局写回warps_per_cta块内 warp 布局element_bitwidth输出元素位宽仅支持32或64缺省32tiles_per_warpwarp 内的 tile 布局缺省为单位 tilecga_layoutCTA 平铺基向量。该布局通过builder.get_amd_mfma_layout(...)落到 IR 层在 python/test/gluon/test_frontend.py 中可以见到其 MLIR 形态例如#ttg.amd_mfma{version 3, warpsPerCTA [4, 1], instrShape [32, 32, 8], isTransposed true}。七、实践示例结合 buffer 访存与 mfma 的 CDNA3 内核骨架综合以上 API一个典型的 CDNA3 内核骨架如下基于测试 python/test/gluon/test_frontend.py 的访存形态与mfma的调用约定接口细节以 cdna3/init.py 为准import triton import triton.language as tl import triton.experimental.gluon.language as ttgl triton.jit def cdna3_buffer_kernel(x_ptr, y_ptr, M, N, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr): offs ttgl.arange(0, BLOCK_M * BLOCK_N).reshape(BLOCK_M, BLOCK_N) # 示意生成偏移张量 # 标量基址 偏移张量 直接访存避免指针张量 a ttgl.amd.cdna3.buffer_load(ptrx_ptr, offsetsoffs, maskNone, other0.0, cache.ca) ttgl.amd.cdna3.buffer_store(stored_valuea, ptry_ptr, offsetsoffs, maskNone, cache.cs) # 原子累加fp32 映射为 FADD默认 acq_rel / gpu(agent) 范围 ttgl.amd.cdna3.buffer_atomic_add(ptry_ptr, offsetsoffs, valuea) # 矩阵核心acc 必填输出类型跟随 acc acc ttgl.zeros((BLOCK_M, BLOCK_N), dtypetl.float32) c ttgl.amd.cdna3.mfma(a, a, acc)使用注意事项汇总offsets必须是元素类型为int32/uint32的分布式张量且必须配合显式布局使用否则_verify_buffer_ops会直接断言失败原子操作前先核对 5.2 节的类型支持矩阵——例如在 CDNA3 上对bf16做原子add会直接报错and/or/xor/xchg只接受 32/64 位整数mfma依赖AMDMFMALayout布局version3 对应 CDNA3且acc不可省略sem/scope缺省语义为acq_rel/gpuagent需要更精细的同步控制时再显式传入。八、结论triton.experimental.gluon.language.amd.cdna3是 Triton Gluon 在 AMD CDNA 3 硬件上的直接硬件抽象层buffer_load/buffer_store以标量基址 偏移张量替代指针张量访存buffer_atomic_*家族提供带类型分派与内存语义控制的全局原子 RMWmfma则把矩阵乘下沉到原生矩阵核心并由AMDMFMALayout显式描述数据分布。无论是手写 CDNA3 内核还是深入理解 Triton 的 AMD 后端代码生成路径本文所列的函数签名、类型约束与源码位置cdna3/init.py、amd/_layouts.py、test_frontend.py都是最直接的参考依据。【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 Triton Gluon AMD CDNA3 目标内建 API 全解：buffer_load/buffer_store/buffer_atomic_* 与 mfma
Triton Gluon AMD CDNA3 目标内建 API 全解buffer_load/buffer_store/buffer_atomic_* 与 mfma【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton导读本文围绕 Triton 仓库中 docs/gluon/api/amd.cdna3.rst 所定义的 AMD CDNA 3 目标特定 Gluon 内建 API系统讲解buffer_load、buffer_store、七个buffer_atomic_*原子操作以及矩阵乘核心mfma的语义、参数约束与底层实现。阅读本文后你将能够在 Gluon 编程模型下直接使用 CDNA 3gfx942如 MI300 系列的硬件原生 buffer 指令与矩阵核心写出绕过指针张量、以标量基址 偏移张量为访存模型的高性能内核。一、背景Gluon 的 AMD 目标特定 API 体系Gluon 是 Triton 中面向显式布局与硬件内建编程的新一代内核编写模型实验特性。它的语言层按硬件目标划分命名空间AMD 侧的入口模块为triton.experimental.gluon.language.amd由 python/triton/experimental/gluon/language/amd/init.py 导出结构如下通用 APIAMDMFMALayout、AMDWMMALayout、slice、warp_pipeline_stage、get_scaled_upcast_fp4_scale_layout代际子模块cdna3、cdna4、cdna5、rdna3、rdna4、gfx1250。文档 docs/gluon/api/amd.rst 中把cdna3与cdna4、cdna5、rdna3、rdna4并列挂接在 GPU Generations 目录下而本文主角amd.cdna3对应 CDNA 3 代际即AMDMFMALayout.version 3对应的gfx942架构见 python/triton/experimental/gluon/language/amd/_layouts.py。API 参考入口可追溯到 docs/gluon/api/index.rst 中的 AMD Intrinsics 一节。二、CDNA3 模块导出的完整内建函数清单根据 docs/gluon/api/amd.cdna3.rst 的 autosummary 列表triton.experimental.gluon.language.amd.cdna3模块文档化导出的内建函数共 10 个类别内建函数全局内存加载buffer_load全局内存存储buffer_store原子 RMW读改写buffer_atomic_add、buffer_atomic_and、buffer_atomic_max、buffer_atomic_min、buffer_atomic_or、buffer_atomic_xchg、buffer_atomic_xor矩阵核心mfma实际源码模块 python/triton/experimental/gluon/language/amd/cdna3/init.py 的__all__中还包括scaled_upcast软件模拟的 FP4/FP8 升精度缩放路径以及未在 autosummary 中列出的内部辅助函数_buffer_atomic_rmw_impl等。下文逐一展开这 10 个文档化 API 的用法与实现细节。三、buffer_load以标量基址 偏移张量加载全局内存3.1 签名与语义buffer_load(ptr, offsets, maskNone, otherNone, cacheNone)ptr指向标量的全局内存基址指针offsets偏移张量distributed_type元素类型必须为int32或uint32mask可选的谓词掩码张量用于按元素屏蔽加载other可选的填充值张量或标量为被mask屏蔽的元素提供默认值当other非空时必须同时提供maskcache可选的缓存修饰符字符串例如测试中出现的.ca、.cg、.cv。与普通load相比核心差异在于不构造指针张量而是通过一个标量基址 一个偏移张量直接生成 AMD 的amdg.buffer_load指令数据直接从全局内存载入寄存器。测试 python/test/gluon/test_frontend.py 中的buffer_load_store_kernel展示了典型用法并在HIP_TARGET_CDNA3目标下校验生成的 IRa ttgl.amd.cdna3.buffer_load(ptrx, offsetsoffsets, maskmask, otherother, cache.ca)对应的 MLIR 形如amdg.buffer_load %arg0[%2], %cst, %cst_1 {cachePolicy ...}印证了该内建直接映射到 AMDGPU buffer load 指令。3.2 实现要点实现位于 cdna3/init.py处理流程如下调用_verify_buffer_ops校验ptr为标量指针、offsets为分布式张量且元素类型为int32/uint32并检查other与mask的配对关系对mask与other执行广播broadcast_impl_valueother会先被 cast 成ptr指向的元素类型将cache字符串经_semantic._str_to_load_cache_modifier解析为CACHE_MODIFIER枚举缺省为NONE以offsets的形状为骨架、以ptr的元素类型为元素类型构造返回类型调用builder.create_buffer_load生成 IR 并包装为 Gluon 张量。四、buffer_store按偏移张量写回全局内存buffer_store(stored_value, ptr, offsets, maskNone, cacheNone)buffer_store是buffer_load的镜像操作把张量stored_value通过标量基址ptr与偏移张量offsets直接写入全局内存同样支持谓词mask与缓存修饰符cache测试中使用了.cs。实现见 cdna3/init.py其额外约束是广播后offsets的形状必须与调用前一致否则抛出ValueError防止存储时悄悄改变访存范围。五、buffer_atomic_*七个全局内存原子 RMW 操作5.1 公共语义CDNA3 模块提供 7 个原子读改写RMW内建add、and、max、min、or、xchg、xor。它们共享同一行为模型模块 docstring见 cdna3/init.py通过标量基址 偏移张量访问全局内存而非指针张量对ptr offsets处的元素执行op(value)并写回mask为布尔张量mask[i] 0的元素被跳过不执行原子操作返回全局内存中的操作前旧值sem指定内存语义描述符缺省为acq_relscope指定同步范围缺省映射到gpuAMDGPU 术语中称agent更细的 CDNA3gfx942内存模型语义可参考 LLVM AMDGPUUsage 文档中关于 gfx942 的 Memory Model 一节。统一签名7 个函数一致buffer_atomic_xxx(ptr, offsets, value, maskNone, semNone, scopeNone)统一实现路径为_buffer_atomic_rmw_impl(op, ptr, offsets, value, arch, mask, sem, scope, _semantic)见 cdna3/init.py校验参数后根据元素类型把字符串操作名映射为ir.ATOMIC_OP枚举经语义层把maskcast 为i1并广播最后调用builder.create_buffer_atomic_rmw。5.2 元素类型支持矩阵CDNA3 限定_verify_element_type_and_dispatch_opcdna3/init.py给出了硬编码的类型约束这是使用本模块前必须核对的关键事实操作支持元素类型映射的原子指令and/or/xor/xchg仅int32、int64AND/OR/XOR/XCHGmax/minint32、int64→ 有符号uint32、uint64→ 无符号SMAX/SMIN、UMAX/UMINfp16/bf16/fp32 不支持adduint32、uint64→ 整数加法float16、float32、float64→ 浮点加法ADDiadd、FADDfadd注意两点边界全部原子操作的基础类型白名单为float16/float32/bfloat16/float64/int32/int64/uint32/uint64其余类型一律断言失败bfloat16的fadd原子在 CDNA3 上不支持源码中明确断言 Buffer atomic fadd with bf16 is only supported on CDNA4 for now即该能力是 CDNA4 才引入的。六、mfma调用 AMD 原生矩阵核心mfma(a, b, acc)mfma计算a * b acc直接使用 AMD 原生矩阵核心Matrix Fused Multiply-Add单元。三个参数中acc为必填项缺失即断言失败返回张量的类型与acc一致。实现见 cdna3/init.py其本质是转发到语义层的_semantic.dot_semantic.dot(a, b, acc, input_precisionknobs.language.fp32_default, max_num_imprecise_accNone, out_dtypeacc.dtype)input_precision取自全局开关knobs.language.fp32_default即默认的 fp32 输入精度策略由 Triton 的 knob 体系统一控制max_num_imprecise_accNone表示不限制不精确累加的次数上限输出类型由acc.dtype决定因此累加器的类型即结果的类型。6.1 配套布局AMDMFMALayout要让mfma高效执行操作数的分布式布局必须与之匹配。CDNA3 对应的布局类AMDMFMALayoutpython/triton/experimental/gluon/language/amd/_layouts.py关键字段如下version架构代号CDNA3 即gfx942对应version 3version 1/2/4 分别对应 gfx908/gfx90a/gfx950instr_shape指令形状(M, N, K)合法 M×N 组合为[32,32]、[16,16]、[64,4]、[4,64]transposed结果是否转置使线程持有同行的连续元素利于链式 dot 与全局写回warps_per_cta块内 warp 布局element_bitwidth输出元素位宽仅支持32或64缺省32tiles_per_warpwarp 内的 tile 布局缺省为单位 tilecga_layoutCTA 平铺基向量。该布局通过builder.get_amd_mfma_layout(...)落到 IR 层在 python/test/gluon/test_frontend.py 中可以见到其 MLIR 形态例如#ttg.amd_mfma{version 3, warpsPerCTA [4, 1], instrShape [32, 32, 8], isTransposed true}。七、实践示例结合 buffer 访存与 mfma 的 CDNA3 内核骨架综合以上 API一个典型的 CDNA3 内核骨架如下基于测试 python/test/gluon/test_frontend.py 的访存形态与mfma的调用约定接口细节以 cdna3/init.py 为准import triton import triton.language as tl import triton.experimental.gluon.language as ttgl triton.jit def cdna3_buffer_kernel(x_ptr, y_ptr, M, N, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr): offs ttgl.arange(0, BLOCK_M * BLOCK_N).reshape(BLOCK_M, BLOCK_N) # 示意生成偏移张量 # 标量基址 偏移张量 直接访存避免指针张量 a ttgl.amd.cdna3.buffer_load(ptrx_ptr, offsetsoffs, maskNone, other0.0, cache.ca) ttgl.amd.cdna3.buffer_store(stored_valuea, ptry_ptr, offsetsoffs, maskNone, cache.cs) # 原子累加fp32 映射为 FADD默认 acq_rel / gpu(agent) 范围 ttgl.amd.cdna3.buffer_atomic_add(ptry_ptr, offsetsoffs, valuea) # 矩阵核心acc 必填输出类型跟随 acc acc ttgl.zeros((BLOCK_M, BLOCK_N), dtypetl.float32) c ttgl.amd.cdna3.mfma(a, a, acc)使用注意事项汇总offsets必须是元素类型为int32/uint32的分布式张量且必须配合显式布局使用否则_verify_buffer_ops会直接断言失败原子操作前先核对 5.2 节的类型支持矩阵——例如在 CDNA3 上对bf16做原子add会直接报错and/or/xor/xchg只接受 32/64 位整数mfma依赖AMDMFMALayout布局version3 对应 CDNA3且acc不可省略sem/scope缺省语义为acq_rel/gpuagent需要更精细的同步控制时再显式传入。八、结论triton.experimental.gluon.language.amd.cdna3是 Triton Gluon 在 AMD CDNA 3 硬件上的直接硬件抽象层buffer_load/buffer_store以标量基址 偏移张量替代指针张量访存buffer_atomic_*家族提供带类型分派与内存语义控制的全局原子 RMWmfma则把矩阵乘下沉到原生矩阵核心并由AMDMFMALayout显式描述数据分布。无论是手写 CDNA3 内核还是深入理解 Triton 的 AMD 后端代码生成路径本文所列的函数签名、类型约束与源码位置cdna3/init.py、amd/_layouts.py、test_frontend.py都是最直接的参考依据。【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考