forked from GPUOpen-Drivers/llvm-project
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Merged main:5abbf20f0fe5 into amd-gfx:2f1f2c28c29e
Local branch amd-gfx 2f1f2c2 Merged main:09f717b929ae into amd-gfx:4708702c57fa Remote branch main 5abbf20 [ARM] Additional test for Min loop. NFC
- Loading branch information
Showing
58 changed files
with
2,728 additions
and
149 deletions.
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,51 @@ | ||
| // RUN: %clang_cc1 %s -ffreestanding -triple=x86_64-unknown-unknown -target-feature +avx512f -target-feature +amx-int8 \ | ||
| // RUN: -target-feature +amx-bf16 -emit-llvm -o - -Werror -pedantic | FileCheck %s --check-prefixes=CHECK | ||
|
|
||
| #include <immintrin.h> | ||
|
|
||
| char buf[1024]; | ||
| #define STRIDE 32 | ||
|
|
||
| char buf2[1024]; | ||
|
|
||
| // This is an example code and integration test. | ||
| void test_api(int cond, short row, short col) { | ||
| //CHECK-LABEL: @test_api | ||
| //CHECK: call <256 x i32> @llvm.x86.tileloadd64.internal | ||
| //CHECK: call <256 x i32> @llvm.x86.tdpbssd.internal | ||
| //CHECK: call void @llvm.x86.tilestored64.internal | ||
| __tile1024i a = {row, 8}; | ||
| __tile1024i b = {8, col}; | ||
| __tile1024i c = {row, col}; | ||
|
|
||
| if (cond) { | ||
| __tile_loadd(&a, buf, STRIDE); | ||
| __tile_loadd(&b, buf, STRIDE); | ||
| __tile_loadd(&c, buf, STRIDE); | ||
| } else { | ||
| __tile_loadd(&a, buf2, STRIDE); | ||
| __tile_loadd(&b, buf2, STRIDE); | ||
| __tile_loadd(&c, buf2, STRIDE); | ||
| } | ||
| __tile_dpbsud(&c, a, b); | ||
| __tile_stored(buf, STRIDE, c); | ||
| } | ||
|
|
||
| void test_tile_loadd(short row, short col) { | ||
| //CHECK-LABEL: @test_tile_loadd | ||
| //CHECK: call <256 x i32> @llvm.x86.tileloadd64.internal | ||
| __tile1024i a = {row, col}; | ||
| __tile_loadd(&a, buf, STRIDE); | ||
| } | ||
|
|
||
| void test_tile_dpbsud(__tile1024i a, __tile1024i b, __tile1024i c) { | ||
| //CHECK-LABEL: @test_tile_dpbsud | ||
| //CHECK: call <256 x i32> @llvm.x86.tdpbssd.internal | ||
| __tile_dpbsud(&c, a, b); | ||
| } | ||
|
|
||
| void test_tile_stored(__tile1024i c) { | ||
| //CHECK-LABEL: @test_tile_stored | ||
| //CHECK: call void @llvm.x86.tilestored64.internal | ||
| __tile_stored(buf, STRIDE, c); | ||
| } |
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,97 @@ | ||
| //===- llvm/CodeGen/TileShapeInfo.h - ---------------------------*- 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 | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
| // | ||
| /// \file Shape utility for AMX. | ||
| /// AMX hardware requires to config the shape of tile data register before use. | ||
| /// The 2D shape includes row and column. In AMX intrinsics interface the shape | ||
| /// is passed as 1st and 2nd parameter and they are lowered as the 1st and 2nd | ||
| /// machine operand of AMX pseudo instructions. ShapeT class is to facilitate | ||
| /// tile config and register allocator. The row and column are machine operand | ||
| /// of AMX pseudo instructions. | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #ifndef LLVM_CODEGEN_TILESHAPEINFO_H | ||
| #define LLVM_CODEGEN_TILESHAPEINFO_H | ||
|
|
||
| #include "llvm/ADT/DenseMapInfo.h" | ||
| #include "llvm/CodeGen/MachineInstr.h" | ||
| #include "llvm/CodeGen/MachineOperand.h" | ||
| #include "llvm/CodeGen/MachineRegisterInfo.h" | ||
| #include "llvm/CodeGen/Register.h" | ||
| #include <utility> | ||
|
|
||
| namespace llvm { | ||
|
|
||
| class ShapeT { | ||
| public: | ||
| ShapeT(MachineOperand *Row, MachineOperand *Col, | ||
| const MachineRegisterInfo *MRI = nullptr) | ||
| : Row(Row), Col(Col) { | ||
| if (MRI) | ||
| deduceImm(MRI); | ||
| } | ||
| ShapeT() | ||
| : Row(nullptr), Col(nullptr), RowImm(InvalidImmShape), | ||
| ColImm(InvalidImmShape) {} | ||
| bool operator==(const ShapeT &Shape) { | ||
| MachineOperand *R = Shape.Row; | ||
| MachineOperand *C = Shape.Col; | ||
| if (!R || !C) | ||
| return false; | ||
| if (!Row || !Col) | ||
| return false; | ||
| if (Row->getReg() == R->getReg() && Col->getReg() == C->getReg()) | ||
| return true; | ||
| if ((RowImm != InvalidImmShape) && (ColImm != InvalidImmShape)) | ||
| return RowImm == Shape.getRowImm() && ColImm == Shape.getColImm(); | ||
| return false; | ||
| } | ||
|
|
||
| bool operator!=(const ShapeT &Shape) { return !(*this == Shape); } | ||
|
|
||
| MachineOperand *getRow() const { return Row; } | ||
|
|
||
| MachineOperand *getCol() const { return Col; } | ||
|
|
||
| int64_t getRowImm() const { return RowImm; } | ||
|
|
||
| int64_t getColImm() const { return ColImm; } | ||
|
|
||
| bool isValid() { return (Row != nullptr) && (Col != nullptr); } | ||
|
|
||
| void deduceImm(const MachineRegisterInfo *MRI) { | ||
| // All def must be the same value, otherwise it is invalid MIs. | ||
| // Find the immediate. | ||
| // TODO copy propagation. | ||
| auto GetImm = [&](Register Reg) { | ||
| int64_t Imm = InvalidImmShape; | ||
| for (const MachineOperand &DefMO : MRI->def_operands(Reg)) { | ||
| const auto *MI = DefMO.getParent(); | ||
| if (MI->isMoveImmediate()) { | ||
| Imm = MI->getOperand(1).getImm(); | ||
| break; | ||
| } | ||
| } | ||
| return Imm; | ||
| }; | ||
| RowImm = GetImm(Row->getReg()); | ||
| ColImm = GetImm(Col->getReg()); | ||
| } | ||
|
|
||
| private: | ||
| static constexpr int64_t InvalidImmShape = -1; | ||
| MachineOperand *Row; | ||
| MachineOperand *Col; | ||
| int64_t RowImm; | ||
| int64_t ColImm; | ||
| }; | ||
|
|
||
| } // namespace llvm | ||
|
|
||
| #endif |
Oops, something went wrong.