-
Notifications
You must be signed in to change notification settings - Fork 12.3k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[X86][AMX] Support AMX-FP8 (#113850)
- Loading branch information
Showing
24 changed files
with
384 additions
and
1 deletion.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,95 @@ | ||
/*===------------- amxfp8intrin.h - AMX intrinsics -*- C++ -*----------------=== | ||
* | ||
* Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
* See https://llvm.org/LICENSE.txt for license information. | ||
* SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
* | ||
*===------------------------------------------------------------------------=== | ||
*/ | ||
|
||
#ifndef __IMMINTRIN_H | ||
#error "Never use <amxfp8intrin.h> directly; include <immintrin.h> instead." | ||
#endif /* __IMMINTRIN_H */ | ||
|
||
#ifndef __AMXFP8INTRIN_H | ||
#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. | ||
/// | ||
/// \headerfile <immintrin.h> | ||
/// | ||
/// \code | ||
/// void _tile_dpbf8ps (__tile dst, __tile a, __tile b) | ||
/// \endcode | ||
/// | ||
/// This intrinsic corresponds to the \c TDPBF8PS instruction. | ||
/// | ||
/// \param dst | ||
/// The destination tile. Max size is 1024 Bytes. | ||
/// \param a | ||
/// The 1st source tile. Max size is 1024 Bytes. | ||
/// \param b | ||
/// The 2nd source tile. Max size is 1024 Bytes. | ||
#define _tile_dpbf8ps(dst, a, b) __builtin_ia32_tdpbf8ps((dst), (a), (b)) | ||
|
||
/// Perform the dot product of a BF8 value \a a by an HF8 value \a b | ||
/// accumulating into a Single Precision (FP32) source/dest \a dst. | ||
/// | ||
/// \headerfile <immintrin.h> | ||
/// | ||
/// \code | ||
/// void _tile_dpbhf8ps (__tile dst, __tile a, __tile b) | ||
/// \endcode | ||
/// | ||
/// This intrinsic corresponds to the \c TDPBHF8PS instruction. | ||
/// | ||
/// \param dst | ||
/// The destination tile. Max size is 1024 Bytes. | ||
/// \param a | ||
/// The 1st source tile. Max size is 1024 Bytes. | ||
/// \param b | ||
/// The 2nd source tile. Max size is 1024 Bytes. | ||
#define _tile_dpbhf8ps(dst, a, b) __builtin_ia32_tdpbhf8ps((dst), (a), (b)) | ||
|
||
/// Perform the dot product of an HF8 value \a a by a BF8 value \a b | ||
/// accumulating into a Single Precision (FP32) source/dest \a dst. | ||
/// | ||
/// \headerfile <immintrin.h> | ||
/// | ||
/// \code | ||
/// void _tile_dphbf8ps (__tile dst, __tile a, __tile b) | ||
/// \endcode | ||
/// | ||
/// This intrinsic corresponds to the \c TDPHBF8PS instruction. | ||
/// | ||
/// \param dst | ||
/// The destination tile. Max size is 1024 Bytes. | ||
/// \param a | ||
/// The 1st source tile. Max size is 1024 Bytes. | ||
/// \param b | ||
/// 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 | ||
/// accumulating into a Single Precision (FP32) source/dest \a dst. | ||
/// | ||
/// \headerfile <immintrin.h> | ||
/// | ||
/// \code | ||
/// void _tile_dphf8ps (__tile dst, __tile a, __tile b) | ||
/// \endcode | ||
/// | ||
/// This intrinsic corresponds to the \c TDPHF8PS instruction. | ||
/// | ||
/// \param dst | ||
/// The destination tile. Max size is 1024 Bytes. | ||
/// \param a | ||
/// The 1st source tile. Max size is 1024 Bytes. | ||
/// \param b | ||
/// The 2nd source tile. Max size is 1024 Bytes. | ||
#define _tile_dphf8ps(dst, a, b) __builtin_ia32_tdphf8ps((dst), (a), (b)) | ||
|
||
#endif /* __x86_64__ */ | ||
#endif /* __AMXFP8INTRIN_H */ |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,27 @@ | ||
// 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_amx(void *data) { | ||
//CHECK-LABEL: @test_amx | ||
//CHECK: call void @llvm.x86.tdpbf8ps(i8 1, i8 2, i8 3) | ||
_tile_dpbf8ps(1, 2, 3); | ||
} | ||
|
||
void test_amx2(void *data) { | ||
//CHECK-LABEL: @test_amx2 | ||
//CHECK: call void @llvm.x86.tdpbhf8ps(i8 1, i8 2, i8 3) | ||
_tile_dpbhf8ps(1, 2, 3); | ||
} | ||
|
||
void test_amx3(void *data) { | ||
//CHECK-LABEL: @test_amx3 | ||
//CHECK: call void @llvm.x86.tdphbf8ps(i8 1, i8 2, i8 3) | ||
_tile_dphbf8ps(1, 2, 3); | ||
} | ||
|
||
void test_amx4(void *data) { | ||
//CHECK-LABEL: @test_amx4 | ||
//CHECK: call void @llvm.x86.tdphf8ps(i8 1, i8 2, i8 3) | ||
_tile_dphf8ps(1, 2, 3); | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,10 @@ | ||
// RUN: %clang_cc1 %s -ffreestanding -triple=x86_64-unknown-unknown -target-feature +amx-tile -target-feature +amx-fp8 -verify | ||
|
||
#include <immintrin.h> | ||
|
||
void test_amx(void *data) { | ||
_tile_dpbf8ps(4, 3, 3); // expected-error {{tile arguments must refer to different tiles}} | ||
_tile_dpbhf8ps(4, 3, 3); // expected-error {{tile arguments must refer to different tiles}} | ||
_tile_dphbf8ps(4, 3, 3); // expected-error {{tile arguments must refer to different tiles}} | ||
_tile_dphf8ps(4, 3, 3); // expected-error {{tile arguments must refer to different tiles}} | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,32 @@ | ||
// RUN: %clang_cc1 %s -ffreestanding -triple=x86_64-unknown-unknown -target-feature +amx-fp8 -emit-llvm -o - -Wall -Werror -pedantic | FileCheck %s | ||
|
||
void f_tilemul(short a) | ||
{ | ||
//CHECK: call void asm sideeffect "tileloadd 0(%rsi,%r13,4), %tmm0 \0A\09tileloadd 0(%rdx,%r14,4), %tmm6 \0A\09tdpbf8ps %tmm6, %tmm0, %tmm7 \0A\09tilestored %tmm7, 0(%r12,%r15,4) \0A\09", "~{memory},~{tmm0},~{tmm6},~{tmm7},~{dirflag},~{fpsr},~{flags}"() | ||
__asm__ volatile ("tileloadd 0(%%rsi,%%r13,4), %%tmm0 \n\t" | ||
"tileloadd 0(%%rdx,%%r14,4), %%tmm6 \n\t" | ||
"tdpbf8ps %%tmm6, %%tmm0, %%tmm7 \n\t" | ||
"tilestored %%tmm7, 0(%%r12,%%r15,4) \n\t" | ||
::: "memory", "tmm0", "tmm6", "tmm7"); | ||
|
||
//CHECK: call void asm sideeffect "tileloadd 0(%rsi,%r13,4), %tmm0 \0A\09tileloadd 0(%rdx,%r14,4), %tmm6 \0A\09tdpbhf8ps %tmm6, %tmm0, %tmm7 \0A\09tilestored %tmm7, 0(%r12,%r15,4) \0A\09", "~{memory},~{tmm0},~{tmm6},~{tmm7},~{dirflag},~{fpsr},~{flags}"() | ||
__asm__ volatile ("tileloadd 0(%%rsi,%%r13,4), %%tmm0 \n\t" | ||
"tileloadd 0(%%rdx,%%r14,4), %%tmm6 \n\t" | ||
"tdpbhf8ps %%tmm6, %%tmm0, %%tmm7 \n\t" | ||
"tilestored %%tmm7, 0(%%r12,%%r15,4) \n\t" | ||
::: "memory", "tmm0", "tmm6", "tmm7"); | ||
|
||
//CHECK: call void asm sideeffect "tileloadd 0(%rsi,%r13,4), %tmm0 \0A\09tileloadd 0(%rdx,%r14,4), %tmm6 \0A\09tdphbf8ps %tmm6, %tmm0, %tmm7 \0A\09tilestored %tmm7, 0(%r12,%r15,4) \0A\09", "~{memory},~{tmm0},~{tmm6},~{tmm7},~{dirflag},~{fpsr},~{flags}"() | ||
__asm__ volatile ("tileloadd 0(%%rsi,%%r13,4), %%tmm0 \n\t" | ||
"tileloadd 0(%%rdx,%%r14,4), %%tmm6 \n\t" | ||
"tdphbf8ps %%tmm6, %%tmm0, %%tmm7 \n\t" | ||
"tilestored %%tmm7, 0(%%r12,%%r15,4) \n\t" | ||
::: "memory", "tmm0", "tmm6", "tmm7"); | ||
|
||
//CHECK: call void asm sideeffect "tileloadd 0(%rsi,%r13,4), %tmm0 \0A\09tileloadd 0(%rdx,%r14,4), %tmm6 \0A\09tdphf8ps %tmm6, %tmm0, %tmm7 \0A\09tilestored %tmm7, 0(%r12,%r15,4) \0A\09", "~{memory},~{tmm0},~{tmm6},~{tmm7},~{dirflag},~{fpsr},~{flags}"() | ||
__asm__ volatile ("tileloadd 0(%%rsi,%%r13,4), %%tmm0 \n\t" | ||
"tileloadd 0(%%rdx,%%r14,4), %%tmm6 \n\t" | ||
"tdphf8ps %%tmm6, %%tmm0, %%tmm7 \n\t" | ||
"tilestored %%tmm7, 0(%%r12,%%r15,4) \n\t" | ||
::: "memory", "tmm0", "tmm6", "tmm7"); | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.