[llvm-branch-commits] [clang] [llvm] [X86][ACE] Add x86_bsr type and Block Scale Register support (PR #208706)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Fri Jul 10 05:01:56 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-backend-hexagon
Author: Ganesh (ganeshgit)
<details>
<summary>Changes</summary>
This patch adds comprehensive support for the Block Scale Register (BSR), a 1024-bit architectural register used by ACE (AI Compute Extensions) for storing scale factors in mixed-precision matrix operations.
- Add x86_bsr as an IR type and update similar to x86_amx
- Make x86_bsr explicit in ACE operations.
Intrinsics Infrastructure:
- Add x86_bsr intrinsics in IntrinsicsX86.td:
- bsrmovf: Store both halves to hardware BSR
- bsrmovh_set/bsrmovl_set: Store individual halves
- bsrmovh_get/bsrmovl_get: Read individual halves
- bsr_create: Create x86_bsr from two 512-bit vectors
- Conversion intrinsics between x86_bsr and vector types
X86 Backend:
- Add X86LowerBSRType pass to lower x86_bsr SSA values to implicit hardware register operations (similar to X86LowerAMXType)
- Handle BSR in register allocation and frame lowering
- Reserve BSR during register allocation to prevent conflicts
Header and API:
- Add __bsr struct type in acev1intrin.h bundling lo/hi 512-bit halves
- Add struct-based API: __bsr_make, __bsr_load, __bsr_store, __bsr_get_lo, __bsr_get_hi, __bsr_set_lo, __bsr_set_hi
This PR builds on PR https://github.com/llvm/llvm-project/pull/208408 and should be reviewed after/with it.
Co-authored-by: Umesh Kalvakuntla
---
Patch is 187.15 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/208706.diff
61 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsX86_64.td (+13-6)
- (modified) clang/lib/CodeGen/CGBuiltin.cpp (+18-8)
- (modified) clang/lib/Headers/acev1intrin.h (+239-5)
- (modified) clang/test/CodeGen/X86/ace-api.c (+87)
- (modified) llvm/include/llvm-c/Core.h (+10)
- (modified) llvm/include/llvm/Analysis/IR2Vec.h (+1)
- (modified) llvm/include/llvm/Bitcode/LLVMBitCodes.h (+2)
- (modified) llvm/include/llvm/CodeGen/ValueTypes.td (+1)
- (modified) llvm/include/llvm/IR/DataLayout.h (+2)
- (modified) llvm/include/llvm/IR/Intrinsics.h (+1)
- (modified) llvm/include/llvm/IR/Intrinsics.td (+3)
- (modified) llvm/include/llvm/IR/IntrinsicsX86.td (+94-20)
- (modified) llvm/include/llvm/IR/Type.h (+8-2)
- (modified) llvm/include/llvm/SandboxIR/Type.h (+3)
- (modified) llvm/lib/Analysis/ConstantFolding.cpp (+5-3)
- (modified) llvm/lib/AsmParser/LLLexer.cpp (+1)
- (modified) llvm/lib/Bitcode/Reader/BitcodeReader.cpp (+3)
- (modified) llvm/lib/Bitcode/Writer/BitcodeWriter.cpp (+3)
- (modified) llvm/lib/CodeGen/ValueTypes.cpp (+5)
- (modified) llvm/lib/IR/AsmWriter.cpp (+3)
- (modified) llvm/lib/IR/ConstantFold.cpp (+1-1)
- (modified) llvm/lib/IR/Core.cpp (+8)
- (modified) llvm/lib/IR/DataLayout.cpp (+2)
- (modified) llvm/lib/IR/Intrinsics.cpp (+10)
- (modified) llvm/lib/IR/LLVMContextImpl.cpp (+4-3)
- (modified) llvm/lib/IR/LLVMContextImpl.h (+1-1)
- (modified) llvm/lib/IR/Type.cpp (+16-2)
- (modified) llvm/lib/IR/TypedPointerType.cpp (+1-1)
- (modified) llvm/lib/IR/Verifier.cpp (+10-1)
- (modified) llvm/lib/Target/DirectX/DXILWriter/DXILBitcodeWriter.cpp (+1)
- (modified) llvm/lib/Target/Hexagon/HexagonTargetObjectFile.cpp (+1)
- (modified) llvm/lib/Target/X86/CMakeLists.txt (+1)
- (modified) llvm/lib/Target/X86/X86.h (+12)
- (modified) llvm/lib/Target/X86/X86ISelLowering.cpp (+4)
- (modified) llvm/lib/Target/X86/X86LowerAMXType.cpp (+7-2)
- (added) llvm/lib/Target/X86/X86LowerBSRType.cpp (+373)
- (modified) llvm/lib/Target/X86/X86RegisterInfo.cpp (+3)
- (modified) llvm/lib/Target/X86/X86RegisterInfo.td (+1-1)
- (modified) llvm/lib/Target/X86/X86TargetMachine.cpp (+3-1)
- (modified) llvm/lib/Transforms/InstCombine/InstCombineCasts.cpp (+1-1)
- (modified) llvm/lib/Transforms/InstCombine/InstCombineLoadStoreAlloca.cpp (+7-4)
- (added) llvm/test/Assembler/x86_bsr.ll (+40)
- (added) llvm/test/CodeGen/X86/ACE/ace-bsr-ordering.ll (+32)
- (added) llvm/test/CodeGen/X86/ACE/ace-internal-intrinsics.ll (+176)
- (modified) llvm/test/CodeGen/X86/O0-pipeline.ll (+1)
- (modified) llvm/test/CodeGen/X86/opt-pipeline.ll (+1)
- (modified) llvm/test/Instrumentation/Instrumentor/alloca_and_function.ll (+7-3)
- (modified) llvm/test/Instrumentation/Instrumentor/cast.ll (+30-30)
- (modified) llvm/test/Instrumentation/Instrumentor/cast_crash.ll (+2-2)
- (modified) llvm/test/Instrumentation/Instrumentor/compare.ll (+11-11)
- (modified) llvm/test/Instrumentation/Instrumentor/load_store.ll (+36-36)
- (modified) llvm/test/Instrumentation/Instrumentor/load_store_args.ll (+36-36)
- (modified) llvm/test/Instrumentation/Instrumentor/load_store_noreplace.ll (+36-36)
- (modified) llvm/test/Instrumentation/Instrumentor/module_and_globals.ll (+10-10)
- (modified) llvm/test/Instrumentation/Instrumentor/numeric.ll (+32-32)
- (modified) llvm/test/Instrumentation/Instrumentor/numeric_subtypeid.ll (+1-1)
- (modified) llvm/test/TableGen/CPtrWildcard.td (+2-2)
- (modified) llvm/test/TableGen/x86-fold-tables.inc (-5)
- (added) llvm/test/Verifier/x86_bsr.ll (+6)
- (modified) llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp (+2-2)
- (modified) llvm/utils/TableGen/X86ManualFoldTables.def (+9)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsX86_64.td b/clang/include/clang/Basic/BuiltinsX86_64.td
index 6f34256139144..983d4c2764862 100644
--- a/clang/include/clang/Basic/BuiltinsX86_64.td
+++ b/clang/include/clang/Basic/BuiltinsX86_64.td
@@ -446,11 +446,18 @@ let Features = "acev1", Attributes = [NoThrow] in {
def top4bsud_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, _Vector<256, int>, _Vector<64, signed char>, _Vector<64, signed char>)">;
}
-// Mixed precision internal
+// Mixed precision internal (last arg is BSR as v32i32, converted to x86_bsr by CGBuiltin)
let Features = "acev1", Attributes = [NoThrow] in {
- def top4mxhf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
- def top4mxbhf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
- def top4mxhbf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
- def top4mxbf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
- def top4mxbssps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
+ // BSR versions (8 args, with BSR) - arg 3 is immediate (i8)
+ def top4mxhf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>, _Vector<32, int>)">;
+ def top4mxbhf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>, _Vector<32, int>)">;
+ def top4mxhbf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>, _Vector<32, int>)">;
+ def top4mxbf8ps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>, _Vector<32, int>)">;
+ def top4mxbssps_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>, _Vector<32, int>)">;
+ // Non-BSR versions (7 args, no BSR) - used by __tile_ace_* macros in acev1intrin.h
+ def top4mxhf8ps_nobsr_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, _Constant unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
+ def top4mxbhf8ps_nobsr_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, _Constant unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
+ def top4mxhbf8ps_nobsr_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, _Constant unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
+ def top4mxbf8ps_nobsr_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, _Constant unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
+ def top4mxbssps_nobsr_internal : X86Builtin<"_Vector<256, int>(unsigned short, unsigned short, unsigned short, _Constant unsigned char, _Vector<256, int>, _Vector<16, int>, _Vector<16, int>)">;
}
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 475bfec6199fc..7ac5f34c09ee2 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -7047,11 +7047,17 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
}
}
- // Cast vector type (e.g., v256i32) to x86_amx, this only happen
- // in amx intrinsics.
+ // Cast vector type to x86_amx (v256i32) or x86_bsr (v32i32).
+ // Use CreateIntrinsicWithoutFolding to avoid constant folding issues
+ // with these special types that cannot have constant values.
if (PTy->isX86_AMXTy())
- ArgValue = Builder.CreateIntrinsic(Intrinsic::x86_cast_vector_to_tile,
- {ArgValue->getType()}, {ArgValue});
+ ArgValue = Builder.CreateIntrinsicWithoutFolding(
+ Intrinsic::x86_cast_vector_to_tile, {ArgValue->getType()},
+ {ArgValue});
+ else if (PTy->isX86_BSRTy())
+ ArgValue = Builder.CreateIntrinsicWithoutFolding(
+ Intrinsic::x86_cast_vector_to_bsr, {ArgValue->getType()},
+ {ArgValue});
else
ArgValue = Builder.CreateBitCast(ArgValue, PTy);
}
@@ -7076,11 +7082,15 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
}
}
- // Cast x86_amx to vector type (e.g., v256i32), this only happen
- // in amx intrinsics.
+ // Cast x86_amx (v256i32) or x86_bsr (v32i32) to vector type.
+ // Use CreateIntrinsicWithoutFolding to avoid constant folding issues
+ // with these special types that cannot have constant values.
if (V->getType()->isX86_AMXTy())
- V = Builder.CreateIntrinsic(Intrinsic::x86_cast_tile_to_vector, {RetTy},
- {V});
+ V = Builder.CreateIntrinsicWithoutFolding(
+ Intrinsic::x86_cast_tile_to_vector, {RetTy}, {V});
+ else if (V->getType()->isX86_BSRTy())
+ V = Builder.CreateIntrinsicWithoutFolding(
+ Intrinsic::x86_cast_bsr_to_vector, {RetTy}, {V});
else
V = Builder.CreateBitCast(V, RetTy);
}
diff --git a/clang/lib/Headers/acev1intrin.h b/clang/lib/Headers/acev1intrin.h
index 0491c7f0b383d..44cbd2295a307 100644
--- a/clang/lib/Headers/acev1intrin.h
+++ b/clang/lib/Headers/acev1intrin.h
@@ -19,6 +19,144 @@
#define __DEFAULT_FN_ATTRS_ACE \
__attribute__((__always_inline__, __nodebug__, __target__("acev1")))
+/// Vector type for combining two 512-bit halves into 1024-bit BSR value.
+typedef int __v32si __attribute__((__vector_size__(128)));
+
+/// Combine two 512-bit vector halves into a single 1024-bit vector.
+/// This is an internal helper for BSR intrinsics.
+///
+/// \param __lo
+/// The low 512-bit half (B-scales, BSR bits [511:0]).
+/// \param __hi
+/// The high 512-bit half (A-scales, BSR bits [1023:512]).
+/// \returns A 1024-bit vector with __lo in elements [0:15] and __hi in [16:31].
+static __inline__ __v32si __DEFAULT_FN_ATTRS_ACE
+__bsr_combine_v32(__v16si __lo, __v16si __hi) {
+ return __builtin_shufflevector(__lo, __hi, 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10,
+ 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22,
+ 23, 24, 25, 26, 27, 28, 29, 30, 31);
+}
+
+/// \struct __bsr
+/// \brief BSR (Block Scale Register) struct type.
+///
+/// The __bsr struct bundles the low and high 512-bit halves of the 1024-bit
+/// Block Scale Register into a single value type. This mirrors the __tile1024i
+/// pattern from AMX, but without shape metadata (BSR has fixed 1024-bit size).
+///
+/// BSR layout per ACE spec:
+/// - BSR[511:0] = B-scales (lo) - column/B-input scales
+/// - BSR[1023:512] = A-scales (hi) - row/A-input scales
+///
+/// Usage:
+/// \code
+/// __bsr scales = __bsr_make(lo_zmm, hi_zmm);
+/// __bsr_store(scales); // write to hardware BSR before compute
+/// \endcode
+typedef struct __bsr_str {
+ __m512i lo; ///< Low 512-bit half (B-scales, BSR bits [511:0])
+ __m512i hi; ///< High 512-bit half (A-scales, BSR bits [1023:512])
+} __bsr;
+
+/// Construct a BSR value from low and high 512-bit halves.
+///
+/// \headerfile <immintrin.h>
+///
+/// \param __lo
+/// The low 512-bit half (B-scales, BSR bits [511:0]).
+/// \param __hi
+/// The high 512-bit half (A-scales, BSR bits [1023:512]).
+/// \returns A __bsr struct containing both halves.
+static __inline__ __bsr __DEFAULT_FN_ATTRS_ACE __bsr_make(__m512i __lo,
+ __m512i __hi) {
+ __bsr __b;
+ __b.lo = __lo;
+ __b.hi = __hi;
+ return __b;
+}
+
+/// Extract the low 512-bit half (B-scales) from a BSR value.
+///
+/// \headerfile <immintrin.h>
+///
+/// \param __b
+/// The BSR value to extract from.
+/// \returns The low 512-bit half (B-scales, BSR bits [511:0]).
+static __inline__ __m512i __DEFAULT_FN_ATTRS_ACE __bsr_get_lo(__bsr __b) {
+ return __b.lo;
+}
+
+/// Extract the high 512-bit half (A-scales) from a BSR value.
+///
+/// \headerfile <immintrin.h>
+///
+/// \param __b
+/// The BSR value to extract from.
+/// \returns The high 512-bit half (A-scales, BSR bits [1023:512]).
+static __inline__ __m512i __DEFAULT_FN_ATTRS_ACE __bsr_get_hi(__bsr __b) {
+ return __b.hi;
+}
+
+/// Return a new BSR value with the low half (B-scales) replaced.
+///
+/// \headerfile <immintrin.h>
+///
+/// \param __b
+/// The original BSR value.
+/// \param __lo
+/// The new low 512-bit half (B-scales).
+/// \returns A new __bsr with the low half replaced.
+static __inline__ __bsr __DEFAULT_FN_ATTRS_ACE __bsr_set_lo(__bsr __b,
+ __m512i __lo) {
+ __b.lo = __lo;
+ return __b;
+}
+
+/// Return a new BSR value with the high half (A-scales) replaced.
+///
+/// \headerfile <immintrin.h>
+///
+/// \param __b
+/// The original BSR value.
+/// \param __hi
+/// The new high 512-bit half (A-scales).
+/// \returns A new __bsr with the high half replaced.
+static __inline__ __bsr __DEFAULT_FN_ATTRS_ACE __bsr_set_hi(__bsr __b,
+ __m512i __hi) {
+ __b.hi = __hi;
+ return __b;
+}
+
+/// Store a BSR value to the hardware Block Scale Register.
+///
+/// \headerfile <immintrin.h>
+///
+/// This intrinsic corresponds to the <c> BSRMOVF </c> instruction.
+///
+/// Per ACE spec: BSRMOVF writes A-scales (hi) to BSR[1023:512] and
+/// B-scales (lo) to BSR[511:0].
+///
+/// \param __b
+/// The BSR value to store to the hardware register.
+static __inline__ void __DEFAULT_FN_ATTRS_ACE __bsr_store(__bsr __b) {
+ __builtin_ia32_bsrmovf((__v16si)__b.hi, (__v16si)__b.lo);
+}
+
+/// Load the current hardware BSR state into a __bsr struct.
+///
+/// \headerfile <immintrin.h>
+///
+/// This intrinsic corresponds to the <c> BSRMOVH </c> and <c> BSRMOVL </c>
+/// instructions (read forms).
+///
+/// \returns A __bsr struct containing the current hardware BSR state.
+static __inline__ __bsr __DEFAULT_FN_ATTRS_ACE __bsr_load(void) {
+ __bsr __b;
+ __b.lo = (__m512i)__builtin_ia32_bsrmovl_get();
+ __b.hi = (__m512i)__builtin_ia32_bsrmovh_get();
+ return __b;
+}
+
/// Load tile configuration from a 64-byte memory location. For ACE
/// (Palette 2), the palette_id byte must be 2. Unlike AMX (Palette 1),
/// ACE tiles have fixed dimensions of 16 rows × 64 bytes, so per-tile
@@ -521,7 +659,7 @@ static __inline__ void __tile_ace_top2bf16ps(__acetile *dst, __m512bh src1,
/// \param imm
/// 8-bit immediate selecting BSR scale factors to apply.
#define __tile_ace_top4mxhf8ps(dst, src1, src2, imm) \
- (*(dst) = __builtin_ia32_top4mxhf8ps_internal( \
+ (*(dst) = __builtin_ia32_top4mxhf8ps_nobsr_internal( \
16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2)))
/// Compute 4-way mixed precision outer product with BF8/HF8 format.
@@ -541,7 +679,7 @@ static __inline__ void __tile_ace_top2bf16ps(__acetile *dst, __m512bh src1,
/// \param imm
/// 8-bit immediate selecting BSR scale factors to apply.
#define __tile_ace_top4mxbhf8ps(dst, src1, src2, imm) \
- (*(dst) = __builtin_ia32_top4mxbhf8ps_internal( \
+ (*(dst) = __builtin_ia32_top4mxbhf8ps_nobsr_internal( \
16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2)))
/// Compute 4-way mixed precision outer product with HF8/BF8 format.
@@ -561,7 +699,7 @@ static __inline__ void __tile_ace_top2bf16ps(__acetile *dst, __m512bh src1,
/// \param imm
/// 8-bit immediate selecting BSR scale factors to apply.
#define __tile_ace_top4mxhbf8ps(dst, src1, src2, imm) \
- (*(dst) = __builtin_ia32_top4mxhbf8ps_internal( \
+ (*(dst) = __builtin_ia32_top4mxhbf8ps_nobsr_internal( \
16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2)))
/// Compute 4-way mixed precision outer product with BF8 (E5M2) format.
@@ -581,7 +719,7 @@ static __inline__ void __tile_ace_top2bf16ps(__acetile *dst, __m512bh src1,
/// \param imm
/// 8-bit immediate selecting BSR scale factors to apply.
#define __tile_ace_top4mxbf8ps(dst, src1, src2, imm) \
- (*(dst) = __builtin_ia32_top4mxbf8ps_internal( \
+ (*(dst) = __builtin_ia32_top4mxbf8ps_nobsr_internal( \
16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2)))
/// Compute 4-way mixed precision outer product of MX INT8 with BSR scaling.
@@ -601,7 +739,7 @@ static __inline__ void __tile_ace_top2bf16ps(__acetile *dst, __m512bh src1,
/// \param imm
/// 8-bit immediate selecting BSR scale factors to apply.
#define __tile_ace_top4mxbssps(dst, src1, src2, imm) \
- (*(dst) = __builtin_ia32_top4mxbssps_internal( \
+ (*(dst) = __builtin_ia32_top4mxbssps_nobsr_internal( \
16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2)))
/// Write a ZMM vector as a column in an ACE tile.
@@ -760,6 +898,102 @@ static __inline__ __m512h __tile_ace_cvtrowps2phl(__acetile *src,
return __builtin_ia32_tcvtrowps2phl_internal(16, 64, *src, idx);
}
+/// Compute 4-way mixed precision outer product with HF8 (E4M3) format
+/// using explicit BSR scales. Multiplies HF8 values from src1 with HF8
+/// values from src2, applies scales from the __bsr struct, converts to
+/// FP32 and accumulates into the ACE tile.
+///
+/// \headerfile <immintrin.h>
+///
+/// This intrinsic corresponds to the <c> TOP4MXHF8PS </c> instruction.
+///
+/// \param dst
+/// Pointer to destination/accumulator __acetile.
+/// \param src1
+/// First source ZMM vector containing HF8 (E4M3) values.
+/// \param src2
+/// Second source ZMM vector containing HF8 (E4M3) values.
+/// \param scales
+/// __bsr struct containing A-scales (hi) and B-scales (lo).
+/// \param imm
+/// 8-bit immediate selecting BSR scale factors to apply.
+#define __tile_ace_top4mxhf8ps_bsr(dst, src1, src2, scales, imm) \
+ (*(dst) = __builtin_ia32_top4mxhf8ps_internal( \
+ 16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2), \
+ __bsr_combine_v32((__v16si)((scales).lo), (__v16si)((scales).hi))))
+
+/// Compute 4-way mixed precision outer product with BF8/HF8 format
+/// using explicit BSR scales. Multiplies BF8 (E5M2) values from src1
+/// with HF8 (E4M3) values from src2, applies scales from the __bsr struct,
+/// converts to FP32 and accumulates into the ACE tile.
+///
+/// \headerfile <immintrin.h>
+///
+/// This intrinsic corresponds to the <c> TOP4MXBHF8PS </c> instruction.
+///
+/// \param dst
+/// Pointer to destination/accumulator __acetile.
+/// \param src1
+/// First source ZMM vector containing BF8 (E5M2) values.
+/// \param src2
+/// Second source ZMM vector containing HF8 (E4M3) values.
+/// \param scales
+/// __bsr struct containing A-scales (hi) and B-scales (lo).
+/// \param imm
+/// 8-bit immediate selecting BSR scale factors to apply.
+#define __tile_ace_top4mxbhf8ps_bsr(dst, src1, src2, scales, imm) \
+ (*(dst) = __builtin_ia32_top4mxbhf8ps_internal( \
+ 16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2), \
+ __bsr_combine_v32((__v16si)((scales).lo), (__v16si)((scales).hi))))
+
+/// Compute 4-way mixed precision outer product with HF8/BF8 format
+/// using explicit BSR scales. Multiplies HF8 (E4M3) values from src1
+/// with BF8 (E5M2) values from src2, applies scales from the __bsr struct,
+/// converts to FP32 and accumulates into the ACE tile.
+///
+/// \headerfile <immintrin.h>
+///
+/// This intrinsic corresponds to the <c> TOP4MXHBF8PS </c> instruction.
+///
+/// \param dst
+/// Pointer to destination/accumulator __acetile.
+/// \param src1
+/// First source ZMM vector containing HF8 (E4M3) values.
+/// \param src2
+/// Second source ZMM vector containing BF8 (E5M2) values.
+/// \param scales
+/// __bsr struct containing A-scales (hi) and B-scales (lo).
+/// \param imm
+/// 8-bit immediate selecting BSR scale factors to apply.
+#define __tile_ace_top4mxhbf8ps_bsr(dst, src1, src2, scales, imm) \
+ (*(dst) = __builtin_ia32_top4mxhbf8ps_internal( \
+ 16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2), \
+ __bsr_combine_v32((__v16si)((scales).lo), (__v16si)((scales).hi))))
+
+/// Compute 4-way mixed precision outer product with BF8 (E5M2) format
+/// using explicit BSR scales. Multiplies BF8 values from both sources,
+/// applies scales from the __bsr struct, converts to FP32 and accumulates
+/// into the ACE tile.
+///
+/// \headerfile <immintrin.h>
+///
+/// This intrinsic corresponds to the <c> TOP4MXBF8PS </c> instruction.
+///
+/// \param dst
+/// Pointer to destination/accumulator __acetile.
+/// \param src1
+/// First source ZMM vector containing BF8 (E5M2) values.
+/// \param src2
+/// Second source ZMM vector containing BF8 (E5M2) values.
+/// \param scales
+/// __bsr struct containing A-scales (hi) and B-scales (lo).
+/// \param imm
+/// 8-bit immediate selecting BSR scale factors to apply.
+#define __tile_ace_top4mxbf8ps_bsr(dst, src1, src2, scales, imm) \
+ (*(dst) = __builtin_ia32_top4mxbf8ps_internal( \
+ 16, 64, 64, (imm), *(dst), (__v16si)(src1), (__v16si)(src2), \
+ __bsr_combine_v32((__v16si)((scales).lo), (__v16si)((scales).hi))))
+
#undef __DEFAULT_FN_ATTRS_ACE
#endif /* __x86_64__ */
diff --git a/clang/test/CodeGen/X86/ace-api.c b/clang/test/CodeGen/X86/ace-api.c
index 8f82fbf66e39a..bc4298289bf9e 100644
--- a/clang/test/CodeGen/X86/ace-api.c
+++ b/clang/test/CodeGen/X86/ace-api.c
@@ -143,3 +143,90 @@ void test_ace_workflow(__m512i *input, __m512i *output) {
// Extract result via getrow
output[0] = __tile_ace_getrow(&acc, 0);
}
+
+// Test BSR struct-based API functions
+
+// CHECK-LABEL: @test_bsr_make
+// CHECK: ret void
+void test_bsr_make(__m512i lo, __m512i hi) {
+ __bsr b = __bsr_make(lo, hi);
+ (void)b;
+}
+
+// CHECK-LABEL: @test_bsr_get_lo
+// CHECK: ret <8 x i64>
+__m512i test_bsr_get_lo(__m512i lo, __m512i hi) {
+ __bsr b = __bsr_make(lo, hi);
+ return __bsr_get_lo(b);
+}
+
+// CHECK-LABEL: @test_bsr_get_hi
+// CHECK: ret <8 x i64>
+__m512i test_bsr_get_hi(__m512i lo, __m512i hi) {
+ __bsr b = __bsr_make(lo, hi);
+ return __bsr_get_hi(b);
+}
+
+// CHECK-LABEL: @test_bsr_set_lo
+// CHECK: ret void
+void test_bsr_set_lo(__m512i lo, __m512i hi, __m512i new_lo) {
+ __bsr b = __bsr_make(lo, hi);
+ b = __bsr_set_lo(b, new_lo);
+ (void)b;
+}
+
+// CHECK-LABEL: @test_bsr_set_hi
+// CHECK: ret void
+void test_bsr_set_hi(__m512i lo, __m512i hi, __m512i new_hi) {
+ __bsr b = __bsr_make(lo, hi);
+ b = __bsr_set_hi(b, new_hi);
+ (void)b;
+}
+
+// CHECK-LABEL: @test_bsr_store
+// CHECK: call void @llvm.x86.bsrmovf(<16 x i32> %{{.*}}, <16 x i32> %{{.*}})
+void test_bsr_store(__m512i lo, __m512i hi) {
+ __bsr b = __bsr_make(lo, hi);
+ __bsr_store(b);
+}
+
+// CHECK-LABEL: @test_bsr_load
+// CHECK: call <16 x i32> @llvm.x86.bsrmovl.get()
+// CHECK: call <16 x i32> @llvm.x86.bsrmovh.get()
+__bsr test_bsr_load(void) {
+ return __bsr_load();
+}
+
+// Test BSR-based mixed-precision outer product macros with __...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/208706
More information about the llvm-branch-commits
mailing list