Skip to content

Commit 6d614a8

Browse files
Konstantin PyzhovKonstantin Pyzhov
Konstantin Pyzhov
authored and
Konstantin Pyzhov
committed
Summary:
This CL adds clang declarations of built-in functions for AMDGPU MFMA intrinsics and instructions. OpenCL tests for new built-ins are included. Differential Revision: https://reviews.llvm.org/D72723
1 parent de2dfc8 commit 6d614a8

File tree

5 files changed

+145
-99
lines changed

5 files changed

+145
-99
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -212,5 +212,30 @@ BUILTIN(__builtin_r600_read_tidig_z, "Ui", "nc")
212212
BUILTIN(__builtin_r600_recipsqrt_ieee, "dd", "nc")
213213
BUILTIN(__builtin_r600_recipsqrt_ieeef, "ff", "nc")
214214

215+
//===----------------------------------------------------------------------===//
216+
// MFMA builtins.
217+
//===----------------------------------------------------------------------===//
218+
219+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x1f32, "V32fffV32fIiIiIi", "nc", "mai-insts")
220+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x1f32, "V16fffV16fIiIiIi", "nc", "mai-insts")
221+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x1f32, "V4fffV4fIiIiIi", "nc", "mai-insts")
222+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x2f32, "V16fffV16fIiIiIi", "nc", "mai-insts")
223+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4f32, "V4fffV4fIiIiIi", "nc", "mai-insts")
224+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4f16, "V32fV4hV4hV32fIiIiIi", "nc", "mai-insts")
225+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4f16, "V16fV4hV4hV16fIiIiIi", "nc", "mai-insts")
226+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x4f16, "V4fV4hV4hV4fIiIiIi", "nc", "mai-insts")
227+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x8f16, "V16fV4hV4hV16fIiIiIi", "nc", "mai-insts")
228+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x16f16, "V4fV4hV4hV4fIiIiIi", "nc", "mai-insts")
229+
TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x4i8, "V32iiiV32iIiIiIi", "nc", "mai-insts")
230+
TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x4i8, "V16iiiV16iIiIiIi", "nc", "mai-insts")
231+
TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_4x4x4i8, "V4iiiV4iIiIiIi", "nc", "mai-insts")
232+
TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x8i8, "V16iiiV16iIiIiIi", "nc", "mai-insts")
233+
TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x16i8, "V4iiiV4iIiIiIi", "nc", "mai-insts")
234+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x2bf16, "V32fV2sV2sV32fIiIiIi", "nc", "mai-insts")
235+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x2bf16, "V16fV2sV2sV16fIiIiIi", "nc", "mai-insts")
236+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x2bf16, "V4fV2sV2sV4fIiIiIi", "nc", "mai-insts")
237+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4bf16, "V16fV2sV2sV16fIiIiIi", "nc", "mai-insts")
238+
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x8bf16, "V4fV2sV2sV4fIiIiIi", "nc", "mai-insts")
239+
215240
#undef BUILTIN
216241
#undef TARGET_BUILTIN

‎clang/lib/Basic/Targets/AMDGPU.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -163,6 +163,7 @@ bool AMDGPUTargetInfo::initFeatureMap(
163163
Features["dot4-insts"] = true;
164164
Features["dot5-insts"] = true;
165165
Features["dot6-insts"] = true;
166+
Features["mai-insts"] = true;
166167
LLVM_FALLTHROUGH;
167168
case GK_GFX906:
168169
Features["dl-insts"] = true;

‎clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl

Whitespace-only changes.

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx908-param.cl

Whitespace-only changes.

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 119 additions & 99 deletions
Original file line numberDiff line numberDiff line change
@@ -1725,105 +1725,125 @@ def int_amdgcn_buffer_atomic_fadd : AMDGPUBufferAtomicNoRtn;
17251725
def int_amdgcn_global_atomic_fadd : AMDGPUGlobalAtomicNoRtn;
17261726

17271727
// llvm.amdgcn.mfma.f32.* vdst, srcA, srcB, srcC, cbsz, abid, blgp
1728-
def int_amdgcn_mfma_f32_32x32x1f32 : Intrinsic<[llvm_v32f32_ty],
1729-
[llvm_float_ty, llvm_float_ty, llvm_v32f32_ty,
1730-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1731-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1732-
1733-
def int_amdgcn_mfma_f32_16x16x1f32 : Intrinsic<[llvm_v16f32_ty],
1734-
[llvm_float_ty, llvm_float_ty, llvm_v16f32_ty,
1735-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1736-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1737-
1738-
def int_amdgcn_mfma_f32_4x4x1f32 : Intrinsic<[llvm_v4f32_ty],
1739-
[llvm_float_ty, llvm_float_ty, llvm_v4f32_ty,
1740-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1741-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1742-
1743-
def int_amdgcn_mfma_f32_32x32x2f32 : Intrinsic<[llvm_v16f32_ty],
1744-
[llvm_float_ty, llvm_float_ty, llvm_v16f32_ty,
1745-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1746-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1747-
1748-
def int_amdgcn_mfma_f32_16x16x4f32 : Intrinsic<[llvm_v4f32_ty],
1749-
[llvm_float_ty, llvm_float_ty, llvm_v4f32_ty,
1750-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1751-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1752-
1753-
def int_amdgcn_mfma_f32_32x32x4f16 : Intrinsic<[llvm_v32f32_ty],
1754-
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v32f32_ty,
1755-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1756-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1757-
1758-
def int_amdgcn_mfma_f32_16x16x4f16 : Intrinsic<[llvm_v16f32_ty],
1759-
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
1760-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1761-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1762-
1763-
def int_amdgcn_mfma_f32_4x4x4f16 : Intrinsic<[llvm_v4f32_ty],
1764-
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
1765-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1766-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1767-
1768-
def int_amdgcn_mfma_f32_32x32x8f16 : Intrinsic<[llvm_v16f32_ty],
1769-
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
1770-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1771-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1772-
1773-
def int_amdgcn_mfma_f32_16x16x16f16 : Intrinsic<[llvm_v4f32_ty],
1774-
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
1775-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1776-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1777-
1778-
def int_amdgcn_mfma_i32_32x32x4i8 : Intrinsic<[llvm_v32i32_ty],
1779-
[llvm_i32_ty, llvm_i32_ty, llvm_v32i32_ty,
1780-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1781-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1782-
1783-
def int_amdgcn_mfma_i32_16x16x4i8 : Intrinsic<[llvm_v16i32_ty],
1784-
[llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
1785-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1786-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1787-
1788-
def int_amdgcn_mfma_i32_4x4x4i8 : Intrinsic<[llvm_v4i32_ty],
1789-
[llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
1790-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1791-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1792-
1793-
def int_amdgcn_mfma_i32_32x32x8i8 : Intrinsic<[llvm_v16i32_ty],
1794-
[llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
1795-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1796-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1797-
1798-
def int_amdgcn_mfma_i32_16x16x16i8 : Intrinsic<[llvm_v4i32_ty],
1799-
[llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
1800-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1801-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1802-
1803-
def int_amdgcn_mfma_f32_32x32x2bf16 : Intrinsic<[llvm_v32f32_ty],
1804-
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v32f32_ty,
1805-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1806-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1807-
1808-
def int_amdgcn_mfma_f32_16x16x2bf16 : Intrinsic<[llvm_v16f32_ty],
1809-
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
1810-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1811-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1812-
1813-
def int_amdgcn_mfma_f32_4x4x2bf16 : Intrinsic<[llvm_v4f32_ty],
1814-
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
1815-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1816-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1817-
1818-
def int_amdgcn_mfma_f32_32x32x4bf16 : Intrinsic<[llvm_v16f32_ty],
1819-
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
1820-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1821-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1822-
1823-
def int_amdgcn_mfma_f32_16x16x8bf16 : Intrinsic<[llvm_v4f32_ty],
1824-
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
1825-
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1826-
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1728+
def int_amdgcn_mfma_f32_32x32x1f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x1f32">,
1729+
Intrinsic<[llvm_v32f32_ty],
1730+
[llvm_float_ty, llvm_float_ty, llvm_v32f32_ty,
1731+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1732+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1733+
1734+
def int_amdgcn_mfma_f32_16x16x1f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x1f32">,
1735+
Intrinsic<[llvm_v16f32_ty],
1736+
[llvm_float_ty, llvm_float_ty, llvm_v16f32_ty,
1737+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1738+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1739+
1740+
def int_amdgcn_mfma_f32_4x4x1f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x1f32">,
1741+
Intrinsic<[llvm_v4f32_ty],
1742+
[llvm_float_ty, llvm_float_ty, llvm_v4f32_ty,
1743+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1744+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1745+
1746+
def int_amdgcn_mfma_f32_32x32x2f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x2f32">,
1747+
Intrinsic<[llvm_v16f32_ty],
1748+
[llvm_float_ty, llvm_float_ty, llvm_v16f32_ty,
1749+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1750+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1751+
1752+
def int_amdgcn_mfma_f32_16x16x4f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x4f32">,
1753+
Intrinsic<[llvm_v4f32_ty],
1754+
[llvm_float_ty, llvm_float_ty, llvm_v4f32_ty,
1755+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1756+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1757+
1758+
def int_amdgcn_mfma_f32_32x32x4f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x4f16">,
1759+
Intrinsic<[llvm_v32f32_ty],
1760+
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v32f32_ty,
1761+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1762+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1763+
1764+
def int_amdgcn_mfma_f32_16x16x4f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x4f16">,
1765+
Intrinsic<[llvm_v16f32_ty],
1766+
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
1767+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1768+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1769+
1770+
def int_amdgcn_mfma_f32_4x4x4f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x4f16">,
1771+
Intrinsic<[llvm_v4f32_ty],
1772+
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
1773+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1774+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1775+
1776+
def int_amdgcn_mfma_f32_32x32x8f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x8f16">,
1777+
Intrinsic<[llvm_v16f32_ty],
1778+
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
1779+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1780+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1781+
1782+
def int_amdgcn_mfma_f32_16x16x16f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x16f16">,
1783+
Intrinsic<[llvm_v4f32_ty],
1784+
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
1785+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1786+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1787+
1788+
def int_amdgcn_mfma_i32_32x32x4i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_32x32x4i8">,
1789+
Intrinsic<[llvm_v32i32_ty],
1790+
[llvm_i32_ty, llvm_i32_ty, llvm_v32i32_ty,
1791+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1792+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1793+
1794+
def int_amdgcn_mfma_i32_16x16x4i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_16x16x4i8">,
1795+
Intrinsic<[llvm_v16i32_ty],
1796+
[llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
1797+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1798+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1799+
1800+
def int_amdgcn_mfma_i32_4x4x4i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_4x4x4i8">,
1801+
Intrinsic<[llvm_v4i32_ty],
1802+
[llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
1803+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1804+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1805+
1806+
def int_amdgcn_mfma_i32_32x32x8i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_32x32x8i8">,
1807+
Intrinsic<[llvm_v16i32_ty],
1808+
[llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
1809+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1810+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1811+
1812+
def int_amdgcn_mfma_i32_16x16x16i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_16x16x16i8">,
1813+
Intrinsic<[llvm_v4i32_ty],
1814+
[llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
1815+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1816+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1817+
1818+
def int_amdgcn_mfma_f32_32x32x2bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x2bf16">,
1819+
Intrinsic<[llvm_v32f32_ty],
1820+
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v32f32_ty,
1821+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1822+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1823+
1824+
def int_amdgcn_mfma_f32_16x16x2bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x2bf16">,
1825+
Intrinsic<[llvm_v16f32_ty],
1826+
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
1827+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1828+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1829+
1830+
def int_amdgcn_mfma_f32_4x4x2bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x2bf16">,
1831+
Intrinsic<[llvm_v4f32_ty],
1832+
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
1833+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1834+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1835+
1836+
def int_amdgcn_mfma_f32_32x32x4bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x4bf16">,
1837+
Intrinsic<[llvm_v16f32_ty],
1838+
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
1839+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1840+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1841+
1842+
def int_amdgcn_mfma_f32_16x16x8bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x8bf16">,
1843+
Intrinsic<[llvm_v4f32_ty],
1844+
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
1845+
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1846+
[IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
18271847

18281848
//===----------------------------------------------------------------------===//
18291849
// Special Intrinsics for backend internal use only. No frontend

0 commit comments

Comments
 (0)