Author: Shilei Tian Date: 2026-03-04T08:30:42-05:00 New Revision: 47766d7f8c397da857d9f78db36c71e8a2ca1fbf
URL: https://github.com/llvm/llvm-project/commit/47766d7f8c397da857d9f78db36c71e8a2ca1fbf DIFF: https://github.com/llvm/llvm-project/commit/47766d7f8c397da857d9f78db36c71e8a2ca1fbf.diff LOG: [AMDGPU][Clang][Doc] Add documentation for WMMA builtins (#183939) Added: Modified: clang/include/clang/Basic/BuiltinsAMDGPU.td clang/include/clang/Basic/BuiltinsAMDGPUDocs.td Removed: ################################################################################ diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 38e35bd7d3b71..acd0a34a79253 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -431,23 +431,71 @@ def __builtin_amdgcn_s_wait_event : AMDGPUBuiltin<"void(_Constant short)", [], " // Postfix w32 indicates the builtin requires wavefront size of 32. // Postfix w64 indicates the builtin requires wavefront size of 64. //===----------------------------------------------------------------------===// -def __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, float>)", [Const], "wmma-256b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, float>)", [Const], "wmma-256b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x16_f16_w32 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32 : AMDGPUBuiltin<"_ExtVector<16, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<16, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32 : AMDGPUBuiltin<"_ExtVector<16, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<16, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<4, int>, _Constant bool, _ExtVector<4, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32">; - -def __builtin_amdgcn_wmma_f32_16x16x16_f16_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<4, float>)", [Const], "wmma-256b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<4, float>)", [Const], "wmma-256b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f16_16x16x16_f16_w64 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, _ExtVector<4, int>, _Constant bool, _ExtVector<4, int>, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64">; +def __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, float>)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAF32_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, float>)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAF32_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f16_16x16x16_f16_w32 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAOpsel_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32 : AMDGPUBuiltin<"_ExtVector<16, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<16, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAOpsel_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32 : AMDGPUBuiltin<"_ExtVector<16, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAOpselTied_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32 : AMDGPUBuiltin<"_ExtVector<16, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<16, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAOpselTied_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<4, int>, _Constant bool, _ExtVector<4, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAIU_16x16x16_GFX11]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAIU_16x16x16_GFX11]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} + +def __builtin_amdgcn_wmma_f32_16x16x16_f16_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<4, float>)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAF32_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<4, float>)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAF32_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f16_16x16x16_f16_w64 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAOpsel_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAOpsel_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, _Float16>, _ExtVector<16, _Float16>, _ExtVector<8, _Float16>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAOpselTied_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<16, short>, _ExtVector<16, short>, _ExtVector<8, short>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAOpselTied_16x16x16_GFX11]; + let ArgNames = ["a", "b", "c", "opsel"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, _ExtVector<4, int>, _Constant bool, _ExtVector<4, int>, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAIU_16x16x16_GFX11]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-256b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAIU_16x16x16_GFX11]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} def __builtin_amdgcn_s_sendmsg_rtn : AMDGPUBuiltin<"unsigned int(_Constant unsigned int)", [], "gfx11-insts">; def __builtin_amdgcn_s_sendmsg_rtnl : AMDGPUBuiltin<"uint64_t(_Constant unsigned int)", [], "gfx11-insts">; @@ -686,33 +734,99 @@ def __builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn : AMDGPUBuiltin<"_ExtVector<2, // elements. Therefore, we add an "_gfx12" suffix to distinguish them from the // existing builtins. //===----------------------------------------------------------------------===// -def __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, _Float16>, _ExtVector<8, _Float16>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, short>, _ExtVector<8, short>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, _Float16>, _ExtVector<8, _Float16>, _ExtVector<8, _Float16>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<8, short>, _ExtVector<8, short>, _ExtVector<8, short>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, int, _Constant bool, int, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize32">; +def __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, _Float16>, _ExtVector<8, _Float16>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAF32_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, short>, _ExtVector<8, short>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAF32_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, _Float16>, _ExtVector<8, _Float16>, _ExtVector<8, _Float16>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAHalf_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, short>(_ExtVector<8, short>, _ExtVector<8, short>, _ExtVector<8, short>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAHalf_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAIU_16x16x16_GFX12]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, int, _Constant bool, int, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAIU_16x16x16_GFX12]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} // These are gfx1170 and gfx12 only, but for consistency with the other WMMA // variants we're keeping the "_gfx12" suffix. -def __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize32">; - -def __builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, short>, _ExtVector<4, short>, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(_ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, short>(_ExtVector<4, short>, _ExtVector<4, short>, _ExtVector<4, short>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, int, _Constant bool, int, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, int, _Constant bool, int, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize64">; +def __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<2, int>, _ExtVector<2, int>, _ExtVector<8, float>)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<2, int>, _Constant bool, _ExtVector<2, int>, _ExtVector<8, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize32"> { + let Documentation = [DocWMMAIU4_16x16x32_GFX12]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} + +def __builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAF32_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, short>, _ExtVector<4, short>, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAF32_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(_ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAHalf_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, short>(_ExtVector<4, short>, _ExtVector<4, short>, _ExtVector<4, short>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAHalf_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, int, _Constant bool, int, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAIU_16x16x16_GFX12]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} +def __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, int, _Constant bool, int, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAIU_16x16x16_GFX12]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} // These are gfx1170 and gfx12 only, but for consistency with the other WMMA // variants we're keeping the "_gfx12" suffix. -def __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64">; -def __builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, int, _Constant bool, int, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize64">; +def __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, _ExtVector<4, float>)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAFP8_16x16x16_GFX12]; + let ArgNames = ["a", "b", "c"]; +} +def __builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12 : AMDGPUBuiltin<"_ExtVector<4, int>(_Constant bool, int, _Constant bool, int, _ExtVector<4, int>, _Constant bool)", [Const], "wmma-128b-insts,wavefrontsize64"> { + let Documentation = [DocWMMAIU4_16x16x32_GFX12]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "clamp"]; +} def __builtin_amdgcn_swmmac_f32_16x16x32_f16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, __fp16>, _ExtVector<16, __fp16>, _ExtVector<8, float>, int)", [Const], "wmma-128b-insts,wavefrontsize32">; def __builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, short>, _ExtVector<16, short>, _ExtVector<8, float>, int)", [Const], "wmma-128b-insts,wavefrontsize32">; @@ -940,35 +1054,122 @@ def __builtin_amdgcn_pk_add_min_i16 : AMDGPUBuiltin<"_ExtVector<2, short>(_ExtVe def __builtin_amdgcn_pk_add_min_u16 : AMDGPUBuiltin<"_ExtVector<2, unsigned short>(_ExtVector<2, unsigned short>, _ExtVector<2, unsigned short>, _ExtVector<2, unsigned short>, _Constant bool)", [Const], "pk-add-min-max-insts">; // GFX1250 WMMA builtins -def __builtin_amdgcn_wmma_f32_16x16x4_f32 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<2, float>, _Constant bool, _ExtVector<2, float>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<16, __bf16>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_bf16_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<16, __bf16>, _Constant short, _ExtVector<8, __bf16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_bf16f32_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<16, __bf16>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_i32_16x16x64_iu8 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<8, int>, _Constant bool, _ExtVector<8, int>, _ExtVector<8, int>, _Constant bool, _Constant bool, ...)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant int, _Constant int, int, _Constant int, _Constant int, int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_16x16x32_f16 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<16, _Float16>, _Constant bool, _ExtVector<16, _Float16>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f16_16x16x32_f16 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_Constant bool, _ExtVector<16, _Float16>, _Constant bool, _ExtVector<16, _Float16>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_scale_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, int, _Constant int, _Constant int, int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; -def __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; +def __builtin_amdgcn_wmma_f32_16x16x4_f32 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<2, float>, _Constant bool, _ExtVector<2, float>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_f32_16x16x4_f32_GFX1250]; + let ArgNames = ["a_neg", "a", "b_neg", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<16, __bf16>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_f32_16x16x32_GFX1250]; + let ArgNames = ["a_neg", "a", "b_neg", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_bf16_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<16, __bf16>, _Constant short, _ExtVector<8, __bf16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_half_16x16x32_GFX1250]; + let ArgNames = ["a_neg", "a", "b_neg", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_bf16f32_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<16, __bf16>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_bf16f32_16x16x32_GFX1250]; + let ArgNames = ["a_neg", "a", "b_neg", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<8, int>, _ExtVector<8, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x64_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_i32_16x16x64_iu8 : AMDGPUBuiltin<"_ExtVector<8, int>(_Constant bool, _ExtVector<8, int>, _Constant bool, _ExtVector<8, int>, _ExtVector<8, int>, _Constant bool, _Constant bool, ...)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_i32_16x16x64_iu8_GFX1250]; + let ArgNames = ["a_sign", "a", "b_sign", "b", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_f8f6f4_GFX1250]; + let ArgNames = ["matrix_a_fmt", "a", "matrix_b_fmt", "b", "c_mod", "c"]; +} +def __builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<8, float>(_ExtVector<16, int>, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_fp8_16x16x128_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant int, _Constant int, int, _Constant int, _Constant int, int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_scale_GFX1250]; + let ArgNames = ["matrix_a_fmt", "a", "matrix_b_fmt", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant int, _ExtVector<16, int>, _Constant int, _ExtVector<16, int>, _Constant short, _ExtVector<8, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_scale16_GFX1250]; + let ArgNames = ["matrix_a_fmt", "a", "matrix_b_fmt", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_16x16x32_f16 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<16, _Float16>, _Constant bool, _ExtVector<16, _Float16>, _Constant short, _ExtVector<8, float>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_f32_16x16x32_GFX1250]; + let ArgNames = ["a_neg", "a", "b_neg", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f16_16x16x32_f16 : AMDGPUBuiltin<"_ExtVector<8, _Float16>(_Constant bool, _ExtVector<16, _Float16>, _Constant bool, _ExtVector<16, _Float16>, _Constant short, _ExtVector<8, _Float16>, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_half_16x16x32_GFX1250]; + let ArgNames = ["a_neg", "a", "b_neg", "b", "c_mod", "c", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_f4_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c"]; +} +def __builtin_amdgcn_wmma_scale_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, int, _Constant int, _Constant int, int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_scale_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; +} +def __builtin_amdgcn_wmma_scale16_f32_32x16x128_f4 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, int>, _ExtVector<8, int>, _Constant short, _ExtVector<16, float>, _Constant int, _Constant int, long int, _Constant int, _Constant int, long int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32"> { + let Documentation = [DocWMMA_scale16_GFX1250]; + let ArgNames = ["a", "b", "c_mod", "c", "matrix_a_scale", "matrix_a_scale_fmt", "matrix_a_scale_exp", "matrix_b_scale", "matrix_b_scale_fmt", "matrix_b_scale_exp", "matrix_a_reuse", "matrix_b_reuse"]; +} def __builtin_amdgcn_swmmac_f32_16x16x64_bf16 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<32, __bf16>, _ExtVector<8, float>, int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; def __builtin_amdgcn_swmmac_bf16_16x16x64_bf16 : AMDGPUBuiltin<"_ExtVector<8, __bf16>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<32, __bf16>, _ExtVector<8, __bf16>, int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; def __builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16 : AMDGPUBuiltin<"_ExtVector<8, float>(_Constant bool, _ExtVector<16, __bf16>, _Constant bool, _ExtVector<32, __bf16>, _ExtVector<8, float>, int, _Constant bool, _Constant bool)", [Const], "gfx1250-insts,wavefrontsize32">; diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td index 0ce6b105eef7a..683c41a414c4e 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td @@ -266,3 +266,329 @@ def DocABIGridSizeZ : Documentation { Returns the total grid size in the Z dimension. }]; } + +//===----------------------------------------------------------------------===// +// WMMA (Wave Matrix Multiply-Accumulate) Builtins +//===----------------------------------------------------------------------===// + +def DocCatWMMA : DocumentationCategory<"WMMA (Wave Matrix Multiply-Accumulate) Builtins"> { + let Content = [{ +WMMA builtins perform wave-cooperative matrix multiply-accumulate of the form +``D = A * B + C``. The work is distributed across all lanes in the wavefront. +The ``_w32`` suffix indicates a wave32 variant, ``_w64`` for wave64. +}]; +} + +//===----------------------------------------------------------------------===// +// GFX11 WMMA — Documentation records +//===----------------------------------------------------------------------===// + +def DocWMMAF32_16x16x16_GFX11 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Multiplies a 16x16 matrix ``A`` by a 16x16 matrix ``B`` (both f16 or bf16) +and adds the 16x16 f32 accumulator ``C`` using fused multiply-add. Returns +the f32 result matrix. + +On GFX11, the A and B input matrices require replication (full 256-bit +operands): 2 copies for wave32, 4 copies for wave64. The total VGPR +footprint is the same for both wave sizes. +}]; +} + +def DocWMMAOpsel_16x16x16_GFX11 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Multiplies a 16x16 matrix ``A`` by a 16x16 matrix ``B`` and adds the +accumulator ``C``, where all operands and the result are f16 or bf16. + +The ``opsel`` argument selects which 16-bit half of the accumulator +registers to read from and write to. The content of the other 16-bit +half is undefined. + +On GFX11, the A and B input matrices require replication (full 256-bit +operands). +}]; +} + +def DocWMMAOpselTied_16x16x16_GFX11 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Same as the non-tied f16/bf16 WMMA variant, but the destination register +is tied to the input accumulator: the unselected 16-bit half is preserved +from the input rather than being undefined. + +On GFX11, the A and B input matrices require replication (full 256-bit +operands). +}]; +} + +def DocWMMAIU_16x16x16_GFX11 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Integer wave matrix multiply-accumulate on packed 8-bit or 4-bit integer +inputs. Multiplies a 16x16 matrix ``A`` by a 16x16 matrix ``B`` and adds +the i32 accumulator ``C``. + +- ``a_sign`` / ``b_sign``: if true, the corresponding operand is treated + as signed; if false, as unsigned. +- ``clamp``: if true, the result saturates instead of wrapping on overflow. + +On GFX11, the A and B input matrices require replication: 2 copies for +wave32 and 4 copies for wave64. In the builtin API, IU8 A/B operands are +128-bit and IU4 A/B operands are 64-bit. +}]; +} + +//===----------------------------------------------------------------------===// +// GFX12 WMMA — Documentation records +//===----------------------------------------------------------------------===// + +def DocWMMAF32_16x16x16_GFX12 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Multiplies a 16x16 matrix ``A`` by a 16x16 matrix ``B`` (both f16 or bf16) +and adds the 16x16 f32 accumulator ``C`` using fused multiply-add. Returns +the f32 result matrix. + +Unlike the GFX11 variants, these GFX12 builtins do not require A/B matrix +replication in their input arguments, so A/B vectors are smaller (128-bit +for wave32 and 64-bit for wave64). +}]; +} + +def DocWMMAHalf_16x16x16_GFX12 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Multiplies a 16x16 matrix ``A`` by a 16x16 matrix ``B`` and adds the +accumulator ``C``, where all operands and the result are f16 or bf16. +Returns the result matrix. + +On GFX12, the output always uses the low 16-bit half of the accumulator +registers (the ``opsel`` bit is implicitly false). In these builtins, +no explicit A/B matrix replication is required in the arguments. +}]; +} + +def DocWMMAIU_16x16x16_GFX12 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Integer wave matrix multiply-accumulate on packed 8-bit or 4-bit integer +inputs. Multiplies a 16x16 matrix ``A`` by a 16x16 matrix ``B`` and adds +the i32 accumulator ``C``. + +- ``a_sign`` / ``b_sign``: if true, the corresponding operand is treated + as signed; if false, as unsigned. +- ``clamp``: if true, the result saturates instead of wrapping on overflow. + +In these builtins, no explicit A/B matrix replication is required in the +arguments. +}]; +} + +def DocWMMAFP8_16x16x16_GFX12 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +WMMA with 8-bit floating-point inputs (FP8 or BF8) and f32 result. +Multiplies a 16x16 matrix ``A`` by a 16x16 matrix ``B`` and adds the f32 +accumulator ``C``. + +A and B contain packed 8-bit floats (4 values per 32-bit word). Since FP8 +and BF8 are not first-class LLVM types, the inputs are passed as integer +vectors. +}]; +} + +def DocWMMAIU4_16x16x32_GFX12 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Integer wave matrix multiply-accumulate on packed 4-bit integer inputs with +a 32-deep K dimension (16x16x32). Multiplies a 16x32 matrix ``A`` by a +32x16 matrix ``B`` and adds the i32 accumulator ``C``. + +- ``a_sign`` / ``b_sign``: if true, the corresponding operand is treated + as signed; if false, as unsigned. +- ``clamp``: if true, the result saturates instead of wrapping on overflow. +}]; +} + +//===----------------------------------------------------------------------===// +// GFX1250 WMMA — Documentation records +//===----------------------------------------------------------------------===// + +def DocWMMA_f32_16x16x4_f32_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Wave matrix multiply-accumulate with f32 inputs and f32 output, using a +4-deep K dimension. Multiplies a 16x4 matrix ``A`` by a 4x16 matrix ``B`` +and adds the 16x16 f32 accumulator ``C``. + +- ``a_neg`` / ``b_neg``: if true, negate the corresponding input matrix. +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} + +def DocWMMA_f32_16x16x32_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Wave matrix multiply-accumulate with f16 or bf16 inputs and f32 output, +using a 32-deep K dimension. Multiplies a 16x32 matrix ``A`` by a 32x16 +matrix ``B`` and adds the 16x16 f32 accumulator ``C``. + +- ``a_neg`` / ``b_neg``: if true, negate the corresponding input matrix. +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} + +def DocWMMA_bf16f32_16x16x32_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Wave matrix multiply-accumulate with bf16 inputs, f32 accumulator, and bf16 +output. Uses a 32-deep K dimension. Multiplies a 16x32 matrix ``A`` by a +32x16 matrix ``B``, adds the 16x16 f32 accumulator ``C``, and converts the +result to bf16. + +- ``a_neg`` / ``b_neg``: if true, negate the corresponding input matrix. +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} + +def DocWMMA_half_16x16x32_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Wave matrix multiply-accumulate with f16 or bf16 inputs and matching +f16 or bf16 output, using a 32-deep K dimension. Multiplies a 16x32 +matrix ``A`` by a 32x16 matrix ``B`` and adds the 16x16 accumulator ``C``. + +- ``a_neg`` / ``b_neg``: if true, negate the corresponding input matrix. +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} + +def DocWMMA_fp8_16x16x64_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Wave matrix multiply-accumulate with FP8 or BF8 inputs and f32 or f16 +output, using a 64-deep K dimension. Multiplies a 16x64 matrix ``A`` by a +64x16 matrix ``B`` and adds the 16x16 accumulator ``C``. + +A and B contain packed 8-bit floats passed as integer vectors. + +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} + +def DocWMMA_fp8_16x16x128_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Wave matrix multiply-accumulate with FP8 or BF8 inputs and f32 or f16 +output, using a 128-deep K dimension. Multiplies a 16x128 matrix ``A`` +by a 128x16 matrix ``B`` and adds the 16x16 accumulator ``C``. + +A and B contain packed 8-bit floats passed as integer vectors. + +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} + +def DocWMMA_i32_16x16x64_iu8_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Integer wave matrix multiply-accumulate on packed 8-bit integer inputs +with a 64-deep K dimension. Multiplies a 16x64 matrix ``A`` by a 64x16 +matrix ``B`` and adds the i32 accumulator ``C``. + +- ``a_sign`` / ``b_sign``: if true, the corresponding operand is treated + as signed; if false, as unsigned. +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +- ``clamp`` (optional): an optional 8th argument. When present, it must be + a compile-time constant integer expression convertible to bool. If true, + the result saturates instead of wrapping on overflow. If omitted, defaults + to false. +}]; +} + +def DocWMMA_f8f6f4_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Mixed-format wave matrix multiply-accumulate with a 128-deep K dimension. +Multiplies a 16x128 matrix ``A`` by a 128x16 matrix ``B`` and adds the +16x16 f32 accumulator ``C``. The input element types are selected +independently per operand. + +- ``matrix_a_fmt`` / ``matrix_b_fmt``: format selectors for A and B + (FP8, BF8, FP6, BF6, or FP4). Must be compile-time constants. +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +}]; +} + +def DocWMMA_f4_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +FP4 wave matrix multiply-accumulate with a 128-deep K dimension and a +32x16 output. Multiplies a 32x128 matrix ``A`` by a 128x16 matrix ``B`` +and adds the 32x16 f32 accumulator ``C``. + +- ``c_mod``: accumulator modifier (0 = none, 1 = neg, 2 = abs, + 3 = neg(abs)). +}]; +} + +def DocWMMA_scale_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Scaled wave matrix multiply-accumulate. Extends the ``f8f6f4`` or ``f4`` +WMMA with per-operand scale factors applied during the computation. + +- For ``..._f8f6f4`` variants, ``matrix_a_fmt`` / ``matrix_b_fmt`` are format + selectors for A and B. They must be compile-time constants. +- ``c_mod``: accumulator modifier. +- ``matrix_a_scale`` / ``matrix_b_scale``: scale factor selectors. + Must be compile-time constants. +- ``matrix_a_scale_fmt`` / ``matrix_b_scale_fmt``: scale format selectors. + Must be compile-time constants. +- ``matrix_a_scale_exp`` / ``matrix_b_scale_exp``: 32-bit scale exponents. +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} + +def DocWMMA_scale16_GFX1250 : Documentation { + let Category = DocCatWMMA; + let Content = [{ +Scaled wave matrix multiply-accumulate with 64-bit (16-bit per element) +scale exponents, as opposed to the 32-bit exponents used by the non-16 +scale variant. + +- For ``..._f8f6f4`` variants, ``matrix_a_fmt`` / ``matrix_b_fmt`` are format + selectors for A and B. They must be compile-time constants. +- ``c_mod``: accumulator modifier. +- ``matrix_a_scale`` / ``matrix_b_scale``: scale factor selectors. + Must be compile-time constants. +- ``matrix_a_scale_fmt`` / ``matrix_b_scale_fmt``: scale format selectors. + Must be compile-time constants. +- ``matrix_a_scale_exp`` / ``matrix_b_scale_exp``: 64-bit scale exponents. +- ``matrix_a_reuse`` / ``matrix_b_reuse``: hints to the hardware that + matrix A or B data can be reused from a previous WMMA instruction. +}]; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
