https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/183939
>From f8f8d5a43d17a4dd77f991cbf9e25426da78b2f6 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Sat, 28 Feb 2026 14:16:08 -0500 Subject: [PATCH 1/2] [Clang][TableGen] Sort undocumented builtins after documented ones in generated docs The builtin documentation emitter previously sorted all categories purely alphabetically, which placed the "Undocumented" section before categories like "WMMA" in the generated RST. This made the output confusing since stub entries appeared before real documentation. Push the "Undocumented" category to the end of the output so that all documented categories appear first, regardless of their names. --- clang/test/TableGen/builtin-docs.td | 30 ++++++++++++------- clang/utils/TableGen/ClangBuiltinsEmitter.cpp | 13 ++++++-- 2 files changed, 30 insertions(+), 13 deletions(-) diff --git a/clang/test/TableGen/builtin-docs.td b/clang/test/TableGen/builtin-docs.td index 99de9767575d1..468268de4e6c6 100644 --- a/clang/test/TableGen/builtin-docs.td +++ b/clang/test/TableGen/builtin-docs.td @@ -1,6 +1,16 @@ // RUN: clang-tblgen -gen-builtin-docs -I%p/../../include %s -o - 2>&1 | \ // RUN: FileCheck %s +// Test that the Undocumented category always sorts after all documented +// categories, even when its name would come earlier alphabetically (e.g. +// "Undocumented" < "Work-Item Builtins"). +// RUN: clang-tblgen -gen-builtin-docs -I%p/../../include %s -o - 2>&1 | \ +// RUN: FileCheck --check-prefix=CAT-ORDER %s +// CAT-ORDER: ABI Builtins +// CAT-ORDER: Instruction Builtins +// CAT-ORDER: Work-Item Builtins +// CAT-ORDER: Undocumented + // Test that mismatched ArgNames count produces an error. // RUN: not clang-tblgen -gen-builtin-docs -I%p/../../include %s -o - \ // RUN: -DERROR_ARGNAMES_MISMATCH 2>&1 | \ @@ -226,7 +236,16 @@ def __builtin_test_no_doc : TestBuiltin<"int(int, int)">; // CHECK: void __builtin_test_store(int val, int address_space<1> * ptr) // CHECK: Stores a value to the given global memory pointer. -// --- Undocumented (alphabetically between Instruction and Work-Item) --- +// --- Work-Item Builtins (inline doc, no ArgNames) --- +// CHECK: Work-Item Builtins +// CHECK-NEXT: ================== +// CHECK: These builtins return work-item identification. +// CHECK: ``__builtin_test_workitem_id_x`` +// CHECK: unsigned int __builtin_test_workitem_id_x() +// CHECK: **Target Features:** some-feature +// CHECK: Returns the work-item id in the x dimension. + +// --- Undocumented (sorted last, after all documented categories) --- // CHECK: Undocumented // CHECK-NEXT: ============ @@ -239,15 +258,6 @@ def __builtin_test_no_doc : TestBuiltin<"int(int, int)">; // CHECK: void __builtin_test_undocumented() // CHECK: No documentation. -// --- Work-Item Builtins (inline doc, no ArgNames) --- -// CHECK: Work-Item Builtins -// CHECK-NEXT: ================== -// CHECK: These builtins return work-item identification. -// CHECK: ``__builtin_test_workitem_id_x`` -// CHECK: unsigned int __builtin_test_workitem_id_x() -// CHECK: **Target Features:** some-feature -// CHECK: Returns the work-item id in the x dimension. - //===----------------------------------------------------------------------===// // Error test: ArgNames count mismatch //===----------------------------------------------------------------------===// diff --git a/clang/utils/TableGen/ClangBuiltinsEmitter.cpp b/clang/utils/TableGen/ClangBuiltinsEmitter.cpp index f381d6dce6261..f628a993a23cc 100644 --- a/clang/utils/TableGen/ClangBuiltinsEmitter.cpp +++ b/clang/utils/TableGen/ClangBuiltinsEmitter.cpp @@ -769,10 +769,17 @@ void clang::EmitClangBuiltinDocs(const RecordKeeper &Records, raw_ostream &OS) { } } - // Sort categories alphabetically by name for deterministic output. + // Sort categories alphabetically by name for deterministic output, but + // push the "Undocumented" category to the end so that documented sections + // always appear first. llvm::sort(SplitDocs, [](const auto &A, const auto &B) { - return A.first->getValueAsString("Name") < - B.first->getValueAsString("Name"); + StringRef NameA = A.first->getValueAsString("Name"); + StringRef NameB = B.first->getValueAsString("Name"); + bool UndocA = (NameA == "Undocumented"); + bool UndocB = (NameB == "Undocumented"); + if (UndocA != UndocB) + return UndocB; + return NameA < NameB; }); // Write out each category and its builtins. >From 88ffa4650b9536eddcbe9fcd98d55591cd52ffdc Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Sat, 28 Feb 2026 14:09:24 -0500 Subject: [PATCH 2/2] [AMDGPU][Clang][Doc] Add documentation for WMMA builtins --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 339 ++++++++++++++---- .../include/clang/Basic/BuiltinsAMDGPUDocs.td | 326 +++++++++++++++++ 2 files changed, 596 insertions(+), 69 deletions(-) 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
