https://github.com/fzou1 updated https://github.com/llvm/llvm-project/pull/115829
>From 9fd6e9e598423b6cc58a25fe70cc12a846483be5 Mon Sep 17 00:00:00 2001 From: Feng Zou <feng....@intel.com> Date: Thu, 7 Nov 2024 11:56:17 +0800 Subject: [PATCH 1/2] [X86][AMX] Add AMX FP8 new APIs This is a follow-up to #113850. Ref.: https://cdrdv2.intel.com/v1/dl/getContent/671368 --- clang/include/clang/Basic/BuiltinsX86_64.def | 4 + clang/lib/Headers/amxfp8intrin.h | 175 ++++++++++++++++--- clang/test/CodeGen/X86/amx_fp8_api.c | 36 ++++ llvm/include/llvm/IR/IntrinsicsX86.td | 25 +++ llvm/lib/Target/X86/X86ExpandPseudo.cpp | 18 +- llvm/lib/Target/X86/X86InstrAMX.td | 31 ++++ llvm/lib/Target/X86/X86RegisterInfo.cpp | 6 +- 7 files changed, 272 insertions(+), 23 deletions(-) create mode 100644 clang/test/CodeGen/X86/amx_fp8_api.c diff --git a/clang/include/clang/Basic/BuiltinsX86_64.def b/clang/include/clang/Basic/BuiltinsX86_64.def index 25c10d39df32e2..8653fc217bdddb 100644 --- a/clang/include/clang/Basic/BuiltinsX86_64.def +++ b/clang/include/clang/Basic/BuiltinsX86_64.def @@ -141,6 +141,10 @@ TARGET_BUILTIN(__builtin_ia32_tcvtrowps2phl_internal, "V32xUsUsV256iUi", "n", "a TARGET_BUILTIN(__builtin_ia32_tilemovrow_internal, "V16iUsUsV256iUi", "n", "amx-avx512,avx10.2-512") TARGET_BUILTIN(__builtin_ia32_tmmultf32ps_internal, "V256iUsUsUsV256iV256iV256i", "n", "amx-tf32") TARGET_BUILTIN(__builtin_ia32_ttmmultf32ps_internal, "V256iUsUsUsV256iV256iV256i", "n", "amx-tf32,amx-transpose") +TARGET_BUILTIN(__builtin_ia32_tdpbf8ps_internal, "V256iUsUsUsV256iV256iV256i", "n", "amx-fp8") +TARGET_BUILTIN(__builtin_ia32_tdpbhf8ps_internal, "V256iUsUsUsV256iV256iV256i", "n", "amx-fp8") +TARGET_BUILTIN(__builtin_ia32_tdphbf8ps_internal, "V256iUsUsUsV256iV256iV256i", "n", "amx-fp8") +TARGET_BUILTIN(__builtin_ia32_tdphf8ps_internal, "V256iUsUsUsV256iV256iV256i", "n", "amx-fp8") // AMX TARGET_BUILTIN(__builtin_ia32_tile_loadconfig, "vvC*", "n", "amx-tile") diff --git a/clang/lib/Headers/amxfp8intrin.h b/clang/lib/Headers/amxfp8intrin.h index 0f5ddc87e5a752..4ada936a5d40af 100644 --- a/clang/lib/Headers/amxfp8intrin.h +++ b/clang/lib/Headers/amxfp8intrin.h @@ -15,81 +15,214 @@ #define __AMXFP8INTRIN_H #ifdef __x86_64__ -/// Peform the dot product of a BF8 value \a a by a BF8 value \a b accumulating -/// into a Single Precision (FP32) source/dest \a dst. +#define __DEFAULT_FN_ATTRS_FP8 \ + __attribute__((__always_inline__, __nodebug__, __target__("amx-fp8"))) + +static __inline__ _tile1024i __DEFAULT_FN_ATTRS_FP8 +_tile_dpbf8ps_internal(unsigned short m, unsigned short n, unsigned short k, + _tile1024i dst, _tile1024i src1, _tile1024i src2) { + return __builtin_ia32_tdpbf8ps_internal(m, n, k, dst, src1, src2); +} + +/// Perform the dot product of a BF8 value \a src1 by a BF8 value \a src2 +/// accumulating into a Single Precision (FP32) source/dest \a dst. /// /// \headerfile <immintrin.h> /// /// \code -/// void _tile_dpbf8ps (__tile dst, __tile a, __tile b) +/// void __tile_dpbf8ps (__tile1024i *dst, __tile1024i src1, __tile1024i src2) +/// \endcode +/// +/// \code{.operation} +/// FOR m := 0 TO dst.rows - 1 +/// temp1[(dst.colsb / 4 - 1) : 0] = 0 +/// FOR k := 0 TO src1.colsb / 4 - 1 +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// temp1[n] += +/// INT64(src1.row[m].float8[4*k+0]) * INT64(src2.row[k].float8[4*n+0]) +/// + INT64(src1.row[m].float8[4*k+1]) * INT64(src2.row[k].float8[4*n+1]) +/// + INT64(src1.row[m].float8[4*k+2]) * INT64(src2.row[k].float8[4*n+2]) +/// + INT64(src1.row[m].float8[4*k+3]) * INT64(src2.row[k].float8[4*n+3]) +/// ENDFOR +/// ENDFOR +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// tmp.row[m].fp32[n] = dst.row[m].fp32[n] + FP32(temp1[n]) +/// ENDFOR +/// write_row_and_zero(dst, m, tmp, dst.colsb) +/// zero_upper_rows(dst, dst.rows) +/// zero_tileconfig_start() /// \endcode /// /// This intrinsic corresponds to the \c TDPBF8PS instruction. /// /// \param dst /// The destination tile. Max size is 1024 Bytes. -/// \param a +/// \param src1 /// The 1st source tile. Max size is 1024 Bytes. -/// \param b +/// \param src2 /// The 2nd source tile. Max size is 1024 Bytes. -#define _tile_dpbf8ps(dst, a, b) __builtin_ia32_tdpbf8ps((dst), (a), (b)) +__DEFAULT_FN_ATTRS_FP8 static void +__tile_dpbf8ps(__tile1024i *dst, __tile1024i src1, __tile1024i src2) { + dst->tile = _tile_dpbf8ps_internal(src1.row, src2.col, src1.col, dst->tile, + src1.tile, src2.tile); +} + +static __inline__ _tile1024i __DEFAULT_FN_ATTRS_FP8 +_tile_dpbhf8ps_internal(unsigned short m, unsigned short n, unsigned short k, + _tile1024i dst, _tile1024i src1, _tile1024i src2) { + return __builtin_ia32_tdpbhf8ps_internal(m, n, k, dst, src1, src2); +} -/// Perform the dot product of a BF8 value \a a by an HF8 value \a b +/// Perform the dot product of a BF8 value \a src1 by an HF8 value \a src2 /// accumulating into a Single Precision (FP32) source/dest \a dst. /// /// \headerfile <immintrin.h> /// /// \code -/// void _tile_dpbhf8ps (__tile dst, __tile a, __tile b) +/// void __tile_dpbhf8ps (__tile1024i dst, __tile1024i src1, __tile1024i src2) +/// \endcode +/// +/// \code{.operation} +/// FOR m := 0 TO dst.rows - 1 +/// temp1[(dst.colsb / 4 - 1) : 0] = 0 +/// FOR k := 0 TO src1.colsb / 4 - 1 +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// temp1[n] += +/// INT64(src1.row[m].float8[4*k+0]) * INT64(src2.row[k].float8[4*n+0]) +/// + INT64(src1.row[m].float8[4*k+1]) * INT64(src2.row[k].float8[4*n+1]) +/// + INT64(src1.row[m].float8[4*k+2]) * INT64(src2.row[k].float8[4*n+2]) +/// + INT64(src1.row[m].float8[4*k+3]) * INT64(src2.row[k].float8[4*n+3]) +/// ENDFOR +/// ENDFOR +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// tmp.row[m].fp32[n] = dst.row[m].fp32[n] + FP32(temp1[n]) +/// ENDFOR +/// write_row_and_zero(dst, m, tmp, dst.colsb) +/// zero_upper_rows(dst, dst.rows) +/// zero_tileconfig_start() /// \endcode /// /// This intrinsic corresponds to the \c TDPBHF8PS instruction. /// /// \param dst /// The destination tile. Max size is 1024 Bytes. -/// \param a +/// \param src1 /// The 1st source tile. Max size is 1024 Bytes. -/// \param b +/// \param src2 /// The 2nd source tile. Max size is 1024 Bytes. -#define _tile_dpbhf8ps(dst, a, b) __builtin_ia32_tdpbhf8ps((dst), (a), (b)) +__DEFAULT_FN_ATTRS_FP8 static void +__tile_dpbhf8ps(__tile1024i *dst, __tile1024i src1, __tile1024i src2) { + dst->tile = _tile_dpbhf8ps_internal(src1.row, src2.col, src1.col, dst->tile, + src1.tile, src2.tile); +} -/// Perform the dot product of an HF8 value \a a by a BF8 value \a b +static __inline__ _tile1024i __DEFAULT_FN_ATTRS_FP8 +_tile_dphbf8ps_internal(unsigned short m, unsigned short n, unsigned short k, + _tile1024i dst, _tile1024i src1, _tile1024i src2) { + return __builtin_ia32_tdphbf8ps_internal(m, n, k, dst, src1, src2); +} + +/// Perform the dot product of an HF8 value \a src1 by a BF8 value \a src2 /// accumulating into a Single Precision (FP32) source/dest \a dst. /// /// \headerfile <immintrin.h> /// /// \code -/// void _tile_dphbf8ps (__tile dst, __tile a, __tile b) +/// void __tile_dphbf8ps (__tile1024i dst, __tile1024i src1, __tile1024i src2) +/// \endcode +/// +/// \code{.operation} +/// FOR m := 0 TO dst.rows - 1 +/// temp1[(dst.colsb / 4 - 1) : 0] = 0 +/// FOR k := 0 TO src1.colsb / 4 - 1 +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// temp1[n] += +/// INT64(src1.row[m].float8[4*k+0]) * INT64(src2.row[k].float8[4*n+0]) +/// + INT64(src1.row[m].float8[4*k+1]) * INT64(src2.row[k].float8[4*n+1]) +/// + INT64(src1.row[m].float8[4*k+2]) * INT64(src2.row[k].float8[4*n+2]) +/// + INT64(src1.row[m].float8[4*k+3]) * INT64(src2.row[k].float8[4*n+3]) +/// ENDFOR +/// ENDFOR +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// tmp.row[m].fp32[n] = dst.row[m].fp32[n] + FP32(temp1[n]) +/// ENDFOR +/// write_row_and_zero(dst, m, tmp, dst.colsb) +/// zero_upper_rows(dst, dst.rows) +/// zero_tileconfig_start() /// \endcode /// /// This intrinsic corresponds to the \c TDPHBF8PS instruction. /// /// \param dst /// The destination tile. Max size is 1024 Bytes. -/// \param a +/// \param src1 /// The 1st source tile. Max size is 1024 Bytes. -/// \param b +/// \param src2 /// The 2nd source tile. Max size is 1024 Bytes. -#define _tile_dphbf8ps(dst, a, b) __builtin_ia32_tdphbf8ps((dst), (a), (b)) -/// Perform the dot product of an HF8 value \a a by an HF8 value \a b +__DEFAULT_FN_ATTRS_FP8 static void +__tile_dphbf8ps(__tile1024i *dst, __tile1024i src1, __tile1024i src2) { + dst->tile = _tile_dphbf8ps_internal(src1.row, src2.col, src1.col, dst->tile, + src1.tile, src2.tile); +} + +static __inline__ _tile1024i __DEFAULT_FN_ATTRS_FP8 +_tile_dphf8ps_internal(unsigned short m, unsigned short n, unsigned short k, + _tile1024i dst, _tile1024i src1, _tile1024i src2) { + return __builtin_ia32_tdphf8ps_internal(m, n, k, dst, src1, src2); +} + +/// Perform the dot product of an HF8 value \a src1 by an HF8 value \a src2 /// accumulating into a Single Precision (FP32) source/dest \a dst. /// /// \headerfile <immintrin.h> /// /// \code -/// void _tile_dphf8ps (__tile dst, __tile a, __tile b) +/// void __tile_dphf8ps (__tile1024i dst, __tile1024i src1, __tile1024i src2) +/// \endcode +/// +/// \code{.operation} +/// FOR m := 0 TO dst.rows - 1 +/// temp1[(dst.colsb / 4 - 1) : 0] = 0 +/// FOR k := 0 TO src1.colsb / 4 - 1 +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// temp1[n] += +/// INT64(src1.row[m].float8[4*k+0]) * INT64(src2.row[k].float8[4*n+0]) +/// + INT64(src1.row[m].float8[4*k+1]) * INT64(src2.row[k].float8[4*n+1]) +/// + INT64(src1.row[m].float8[4*k+2]) * INT64(src2.row[k].float8[4*n+2]) +/// + INT64(src1.row[m].float8[4*k+3]) * INT64(src2.row[k].float8[4*n+3]) +/// ENDFOR +/// ENDFOR +/// FOR n := 0 TO dst.colsb / 4 - 1 +/// tmp.row[m].fp32[n] = dst.row[m].fp32[n] + FP32(temp1[n]) +/// ENDFOR +/// write_row_and_zero(dst, m, tmp, dst.colsb) +/// zero_upper_rows(dst, dst.rows) +/// zero_tileconfig_start() /// \endcode /// /// This intrinsic corresponds to the \c TDPHF8PS instruction. /// /// \param dst /// The destination tile. Max size is 1024 Bytes. -/// \param a +/// \param src1 /// The 1st source tile. Max size is 1024 Bytes. -/// \param b +/// \param src2 /// The 2nd source tile. Max size is 1024 Bytes. -#define _tile_dphf8ps(dst, a, b) __builtin_ia32_tdphf8ps((dst), (a), (b)) +__DEFAULT_FN_ATTRS_FP8 static void +__tile_dphf8ps(__tile1024i *dst, __tile1024i src1, __tile1024i src2) { + dst->tile = _tile_dphf8ps_internal(src1.row, src2.col, src1.col, dst->tile, + src1.tile, src2.tile); +} + +#define _tile_dpbf8ps(dst, src1, src2) \ + __builtin_ia32_tdpbf8ps((dst), (src1), (src2)) +#define _tile_dpbhf8ps(dst, src1, src2) \ + __builtin_ia32_tdpbhf8ps((dst), (src1), (src2)) +#define _tile_dphbf8ps(dst, src1, src2) \ + __builtin_ia32_tdphbf8ps((dst), (src1), (src2)) +#define _tile_dphf8ps(dst, src1, src2) \ + __builtin_ia32_tdphf8ps((dst), (src1), (src2)) #endif /* __x86_64__ */ #endif /* __AMXFP8INTRIN_H */ diff --git a/clang/test/CodeGen/X86/amx_fp8_api.c b/clang/test/CodeGen/X86/amx_fp8_api.c new file mode 100644 index 00000000000000..2a3af1b7f5cd9a --- /dev/null +++ b/clang/test/CodeGen/X86/amx_fp8_api.c @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 %s -ffreestanding -triple=x86_64-unknown-unknown -target-feature +amx-fp8 \ +// RUN: -emit-llvm -o - -Werror -pedantic | FileCheck %s +#include <immintrin.h> + +void test_tdpbf8ps(__tile1024i src1, __tile1024i src2, __tile1024i dst) { + //CHECK-LABEL: @test_tdpbf8ps + //CHECK-DAG: call x86_amx @llvm.x86.cast.vector.to.tile.v256i32(<256 x i32> {{%.*}}) + //CHECK-DAG: call x86_amx @llvm.x86.tdpbf8ps.internal + //CHECK-DAG: call <256 x i32> @llvm.x86.cast.tile.to.vector.v256i32(x86_amx {{%.*}}) + __tile_dpbf8ps(&dst, src1, src2); +} + +void test_tdpbhf8ps(__tile1024i src1, __tile1024i src2, __tile1024i dst) { + //CHECK-LABEL: @test_tdpbhf8ps + //CHECK-DAG: call x86_amx @llvm.x86.cast.vector.to.tile.v256i32(<256 x i32> {{%.*}}) + //CHECK-DAG: call x86_amx @llvm.x86.tdpbhf8ps.internal + //CHECK-DAG: call <256 x i32> @llvm.x86.cast.tile.to.vector.v256i32(x86_amx {{%.*}}) + __tile_dpbhf8ps(&dst, src1, src2); +} + +void test_tdphbf8ps(__tile1024i src1, __tile1024i src2, __tile1024i dst) { + //CHECK-LABEL: @test_tdphbf8ps + //CHECK-DAG: call x86_amx @llvm.x86.cast.vector.to.tile.v256i32(<256 x i32> {{%.*}}) + //CHECK-DAG: call x86_amx @llvm.x86.tdphbf8ps.internal + //CHECK-DAG: call <256 x i32> @llvm.x86.cast.tile.to.vector.v256i32(x86_amx {{%.*}}) + __tile_dphbf8ps(&dst, src1, src2); +} + +void test_tdphf8ps(__tile1024i src1, __tile1024i src2, __tile1024i dst) { + //CHECK-LABEL: @test_tdphf8ps + //CHECK-DAG: call x86_amx @llvm.x86.cast.vector.to.tile.v256i32(<256 x i32> {{%.*}}) + //CHECK-DAG: call x86_amx @llvm.x86.tdphf8ps.internal + //CHECK-DAG: call <256 x i32> @llvm.x86.cast.tile.to.vector.v256i32(x86_amx {{%.*}}) + __tile_dphf8ps(&dst, src1, src2); +} + diff --git a/llvm/include/llvm/IR/IntrinsicsX86.td b/llvm/include/llvm/IR/IntrinsicsX86.td index b2d6f44b7927a9..5211b82cc8d31b 100644 --- a/llvm/include/llvm/IR/IntrinsicsX86.td +++ b/llvm/include/llvm/IR/IntrinsicsX86.td @@ -6120,6 +6120,31 @@ let TargetPrefix = "x86" in { Intrinsic<[llvm_x86amx_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_x86amx_ty, llvm_x86amx_ty, llvm_x86amx_ty], []>; + + def int_x86_tdpbf8ps_internal : + ClangBuiltin<"__builtin_ia32_tdpbf8ps_internal">, + Intrinsic<[llvm_x86amx_ty], + [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, + llvm_x86amx_ty, llvm_x86amx_ty, llvm_x86amx_ty], + []>; + def int_x86_tdpbhf8ps_internal : + ClangBuiltin<"__builtin_ia32_tdpbhf8ps_internal">, + Intrinsic<[llvm_x86amx_ty], + [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, + llvm_x86amx_ty, llvm_x86amx_ty, llvm_x86amx_ty], + []>; + def int_x86_tdphbf8ps_internal : + ClangBuiltin<"__builtin_ia32_tdphbf8ps_internal">, + Intrinsic<[llvm_x86amx_ty], + [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, + llvm_x86amx_ty, llvm_x86amx_ty, llvm_x86amx_ty], + []>; + def int_x86_tdphf8ps_internal : + ClangBuiltin<"__builtin_ia32_tdphf8ps_internal">, + Intrinsic<[llvm_x86amx_ty], + [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, + llvm_x86amx_ty, llvm_x86amx_ty, llvm_x86amx_ty], + []>; } //===----------------------------------------------------------------------===// diff --git a/llvm/lib/Target/X86/X86ExpandPseudo.cpp b/llvm/lib/Target/X86/X86ExpandPseudo.cpp index 4f045d78f75fb2..b673a3766a6832 100644 --- a/llvm/lib/Target/X86/X86ExpandPseudo.cpp +++ b/llvm/lib/Target/X86/X86ExpandPseudo.cpp @@ -757,7 +757,11 @@ bool X86ExpandPseudo::expandMI(MachineBasicBlock &MBB, case X86::PTDPBF16PSV: case X86::PTDPFP16PSV: case X86::PTMMULTF32PSV: - case X86::PTTMMULTF32PSV: { + case X86::PTTMMULTF32PSV: + case X86::PTDPBF8PSV: + case X86::PTDPBHF8PSV: + case X86::PTDPHBF8PSV: + case X86::PTDPHF8PSV: { MI.untieRegOperand(4); for (unsigned i = 3; i > 0; --i) MI.removeOperand(i); @@ -777,6 +781,18 @@ bool X86ExpandPseudo::expandMI(MachineBasicBlock &MBB, case X86::PTTMMULTF32PSV: Opc = X86::TTMMULTF32PS; break; + case X86::PTDPBF8PSV: + Opc = X86::TDPBF8PS; + break; + case X86::PTDPBHF8PSV: + Opc = X86::TDPBHF8PS; + break; + case X86::PTDPHBF8PSV: + Opc = X86::TDPHBF8PS; + break; + case X86::PTDPHF8PSV: + Opc = X86::TDPHF8PS; + break; default: llvm_unreachable("Unexpected Opcode"); diff --git a/llvm/lib/Target/X86/X86InstrAMX.td b/llvm/lib/Target/X86/X86InstrAMX.td index 04527716e31627..da0077a990242b 100644 --- a/llvm/lib/Target/X86/X86InstrAMX.td +++ b/llvm/lib/Target/X86/X86InstrAMX.td @@ -304,6 +304,37 @@ let Predicates = [HasAMXFP8, In64BitMode] in { [(int_x86_tdphf8ps timm:$src1, timm:$src2, timm:$src3)]>; } + + let Constraints = "$src4 = $dst" in { + def PTDPBF8PSV : PseudoI<(outs TILE:$dst), + (ins GR16:$src1, GR16:$src2, GR16:$src3, + TILE:$src4, TILE:$src5, TILE:$src6), + [(set TILE:$dst, + (int_x86_tdpbf8ps_internal GR16:$src1, + GR16:$src2, GR16:$src3, TILE:$src4, + TILE:$src5, TILE:$src6))]>; + def PTDPBHF8PSV : PseudoI<(outs TILE:$dst), + (ins GR16:$src1, GR16:$src2, GR16:$src3, + TILE:$src4, TILE:$src5, TILE:$src6), + [(set TILE:$dst, + (int_x86_tdpbhf8ps_internal GR16:$src1, + GR16:$src2, GR16:$src3, TILE:$src4, + TILE:$src5, TILE:$src6))]>; + def PTDPHBF8PSV : PseudoI<(outs TILE:$dst), + (ins GR16:$src1, GR16:$src2, GR16:$src3, + TILE:$src4, TILE:$src5, TILE:$src6), + [(set TILE:$dst, + (int_x86_tdphbf8ps_internal GR16:$src1, + GR16:$src2, GR16:$src3, TILE:$src4, + TILE:$src5, TILE:$src6))]>; + def PTDPHF8PSV : PseudoI<(outs TILE:$dst), + (ins GR16:$src1, GR16:$src2, GR16:$src3, + TILE:$src4, TILE:$src5, TILE:$src6), + [(set TILE:$dst, + (int_x86_tdphf8ps_internal GR16:$src1, + GR16:$src2, GR16:$src3, TILE:$src4, + TILE:$src5, TILE:$src6))]>; + } } } diff --git a/llvm/lib/Target/X86/X86RegisterInfo.cpp b/llvm/lib/Target/X86/X86RegisterInfo.cpp index 09418c9bb74d34..25eb90cdd6d60c 100644 --- a/llvm/lib/Target/X86/X86RegisterInfo.cpp +++ b/llvm/lib/Target/X86/X86RegisterInfo.cpp @@ -1078,7 +1078,11 @@ static ShapeT getTileShape(Register VirtReg, VirtRegMap *VRM, case X86::PTCMMRLFP16PSV: case X86::PTTRANSPOSEDV: case X86::PTMMULTF32PSV: - case X86::PTTMMULTF32PSV: { + case X86::PTTMMULTF32PSV: + case X86::PTDPBF8PSV: + case X86::PTDPBHF8PSV: + case X86::PTDPHBF8PSV: + case X86::PTDPHF8PSV: { MachineOperand &MO1 = MI->getOperand(1); MachineOperand &MO2 = MI->getOperand(2); ShapeT Shape(&MO1, &MO2, MRI); >From 7301f5628cc495a18c729e75dda90ca523444b9b Mon Sep 17 00:00:00 2001 From: Feng Zou <feng....@intel.com> Date: Tue, 12 Nov 2024 15:35:54 +0800 Subject: [PATCH 2/2] Undefined function attribute and added missing IR test. --- clang/lib/Headers/amxfp8intrin.h | 2 + llvm/test/CodeGen/X86/amx-fp8-internal.ll | 69 +++++++++++++++++++++++ 2 files changed, 71 insertions(+) create mode 100644 llvm/test/CodeGen/X86/amx-fp8-internal.ll diff --git a/clang/lib/Headers/amxfp8intrin.h b/clang/lib/Headers/amxfp8intrin.h index 4ada936a5d40af..92e7989974e71a 100644 --- a/clang/lib/Headers/amxfp8intrin.h +++ b/clang/lib/Headers/amxfp8intrin.h @@ -224,5 +224,7 @@ __tile_dphf8ps(__tile1024i *dst, __tile1024i src1, __tile1024i src2) { #define _tile_dphf8ps(dst, src1, src2) \ __builtin_ia32_tdphf8ps((dst), (src1), (src2)) +#undef __DEFAULT_FN_ATTRS_FP8 + #endif /* __x86_64__ */ #endif /* __AMXFP8INTRIN_H */ diff --git a/llvm/test/CodeGen/X86/amx-fp8-internal.ll b/llvm/test/CodeGen/X86/amx-fp8-internal.ll new file mode 100644 index 00000000000000..ade71499b361a8 --- /dev/null +++ b/llvm/test/CodeGen/X86/amx-fp8-internal.ll @@ -0,0 +1,69 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py +; RUN: llc < %s -mtriple=x86_64-unknown-unknown -mattr=+amx-tile,+avx512f \ +; RUN: -mattr=+amx-fp8 -verify-machineinstrs | FileCheck %s + +define void @test_amx(i8* %pointer, i8* %base, i64 %stride) { +; CHECK-LABEL: test_amx: +; CHECK: # %bb.0: +; CHECK-NEXT: pushq %rbp +; CHECK-NEXT: .cfi_def_cfa_offset 16 +; CHECK-NEXT: subq $3952, %rsp # imm = 0xF70 +; CHECK-NEXT: .cfi_def_cfa_offset 3968 +; CHECK-NEXT: .cfi_offset %rbp, -16 +; CHECK-NEXT: vxorps %xmm0, %xmm0, %xmm0 +; CHECK-NEXT: vmovups %zmm0, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movb $1, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movb $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movw $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movb $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movw $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movb $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movw $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movb $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movw $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movb $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: movw $8, {{[0-9]+}}(%rsp) +; CHECK-NEXT: ldtilecfg {{[0-9]+}}(%rsp) +; CHECK-NEXT: movw $8, %ax +; CHECK-NEXT: tileloadd (%rsi,%rdx), %tmm0 +; CHECK-NEXT: tilezero %tmm1 +; CHECK-NEXT: tilezero %tmm2 +; CHECK-NEXT: movabsq $64, %rbp +; CHECK-NEXT: tilestored %tmm2, 896(%rsp,%rbp) # 1024-byte Folded Spill +; CHECK-NEXT: tileloadd 896(%rsp,%rbp), %tmm3 # 1024-byte Folded Reload +; CHECK-NEXT: tdpbf8ps %tmm1, %tmm0, %tmm3 +; CHECK-NEXT: tdpbhf8ps %tmm1, %tmm0, %tmm3 +; CHECK-NEXT: tilestored %tmm2, 1920(%rsp,%rbp) # 1024-byte Folded Spill +; CHECK-NEXT: tileloadd 1920(%rsp,%rbp), %tmm4 # 1024-byte Folded Reload +; CHECK-NEXT: tdphbf8ps %tmm1, %tmm0, %tmm4 +; CHECK-NEXT: tdphf8ps %tmm1, %tmm0, %tmm2 +; CHECK-NEXT: tilestored %tmm3, (%rdi,%rdx) +; CHECK-NEXT: addq $3952, %rsp # imm = 0xF70 +; CHECK-NEXT: .cfi_def_cfa_offset 16 +; CHECK-NEXT: popq %rbp +; CHECK-NEXT: .cfi_def_cfa_offset 8 +; CHECK-NEXT: tilerelease +; CHECK-NEXT: vzeroupper +; CHECK-NEXT: retq + + %a = call x86_amx @llvm.x86.tileloadd64.internal(i16 8, i16 8, i8* %base, i64 %stride) + %b = call x86_amx @llvm.x86.tilezero.internal(i16 8, i16 8) + %c = call x86_amx @llvm.x86.tilezero.internal(i16 8, i16 8) + + %c1 = call x86_amx @llvm.x86.tdpbf8ps.internal(i16 8, i16 8, i16 8, x86_amx %c, x86_amx %a, x86_amx %b) + %c2 = call x86_amx @llvm.x86.tdpbhf8ps.internal(i16 8, i16 8, i16 8, x86_amx %c1, x86_amx %a, x86_amx %b) + %c3 = call x86_amx @llvm.x86.tdphbf8ps.internal(i16 8, i16 8, i16 8, x86_amx %c, x86_amx %a, x86_amx %b) + %c4 = call x86_amx @llvm.x86.tdphf8ps.internal(i16 8, i16 8, i16 8, x86_amx %c, x86_amx %a, x86_amx %b) + + call void @llvm.x86.tilestored64.internal(i16 8, i16 8, i8* %pointer, i64 %stride, x86_amx %c2) + ret void +} + +declare x86_amx @llvm.x86.tilezero.internal(i16, i16) +declare x86_amx @llvm.x86.tileloadd64.internal(i16, i16, i8*, i64) +declare void @llvm.x86.tilestored64.internal(i16, i16, i8*, i64, x86_amx) + +declare x86_amx @llvm.x86.tdpbf8ps.internal(i16, i16, i16, x86_amx, x86_amx, x86_amx) +declare x86_amx @llvm.x86.tdpbhf8ps.internal(i16, i16, i16, x86_amx, x86_amx, x86_amx) +declare x86_amx @llvm.x86.tdphbf8ps.internal(i16, i16, i16, x86_amx, x86_amx, x86_amx) +declare x86_amx @llvm.x86.tdphf8ps.internal(i16, i16, i16, x86_amx, x86_amx, x86_amx) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
