Skip to content

Commit dce84f8

Browse files
macurtis-amdmacurtis-amd
committed
Add builtin/intrinsic global_(load|store)_b128 (llvm#4455)
_(This is a consolidation of PR llvm#4213 which went through numerous revisions and accumulated quite a bit of noise.)_ Associated JIRA: [SWDEV-556587](https://ontrack-internal.amd.com/browse/SWDEV-556587) Add clang builtin for system-scoped store of 128bits New clang builtins: 1. `__builtin_amdgcn_global_load_b128` - Example usage: [clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl:57](https://github.com/AMD-Lightning-Internal/llvm-project/blob/amd/dev/macurtis/swdev/553136/draft-0910/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl#L57) - Generates intrinsic `llvm.amdgcn.global.load.b128` - [llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.b128.ll:23](https://github.com/AMD-Lightning-Internal/llvm-project/tree/amd/dev/macurtis/swdev/553136/draft-0910/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.load.b128.ll#L23) 2. `__builtin_amdgcn_global_store_b128` - Example usage: [clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl:18](https://github.com/AMD-Lightning-Internal/llvm-project/blob/amd/dev/macurtis/swdev/553136/draft-0910/clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl#L18) - Generates intrinsic `llvm.amdgcn.global.store.b128` - [llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.store.b128.ll:22](https://github.com/AMD-Lightning-Internal/llvm-project/tree/amd/dev/macurtis/swdev/553136/draft-0910/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.global.store.b128.ll#L22) Co-authored-by: Matthew Curtis <[email protected]>
1 parent 71ba070 commit dce84f8

23 files changed

+3865
-0
lines changed

‎clang/docs/LanguageExtensions.rst‎

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4901,6 +4901,52 @@ for a concrete target, and shall reflect the latter's identity and features.
49014901
Thus, it is possible to author high-level code, in e.g. HIP, that is target
49024902
adaptive in a dynamic fashion, contrary to macro based mechanisms.
49034903
4904+
__builtin_amdgcn_global_load_b128 and __builtin_amdgcn_global_store_b128
4905+
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
4906+
4907+
Signature:
4908+
4909+
.. code-block:: c
4910+
4911+
typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u;
4912+
typedef v4u __attribute__((address_space(1))) *global_ptr_to_v4u;
4913+
4914+
v4u __builtin_amdgcn_global_load_b128(
4915+
v4u __attribute__((address_space(1))) *src,
4916+
const char *scope);
4917+
4918+
void __builtin_amdgcn_global_store_b128(
4919+
v4u __attribute__((address_space(1))) *dst,
4920+
v4u data,
4921+
const char *scope);
4922+
4923+
Load or store a vector of 4 unsigned integers from or to global memory with
4924+
cache behavior specified by `scope` which must be a string literal.
4925+
4926+
Valid values for `scope` are:
4927+
4928+
===================== ==========================================================
4929+
scope architecture name
4930+
===================== ==========================================================
4931+
``"wavefront"`` wave
4932+
4933+
``"workgroup"`` group
4934+
4935+
``"agent"`` device
4936+
4937+
``""`` (empty string) system
4938+
===================== ==========================================================
4939+
4940+
These builtins are only supported on gfx942 and gfx950 devices.
4941+
4942+
For semantics on gfx942, see Tables 47 and 48 in section 9.1.10 "Memory Scope
4943+
and Temporal Controls" of the "AMD Instinct MI300" Instruction Set Architecture
4944+
Reference.
4945+
4946+
For semantics on gfx950, see Tables 49 and 50 in section 9.1.10 "Memory Scope
4947+
and Temporal Controls" of the CDNA4 Instruction Set Architecture Reference.
4948+
4949+
49044950
ARM/AArch64 Language Extensions
49054951
-------------------------------
49064952

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -165,6 +165,9 @@ BUILTIN(__builtin_amdgcn_raw_buffer_load_b128, "V4UiQbiiIi", "n")
165165

166166
TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_load_lds, "vQbv*3IUiiiIiIi", "t", "vmem-to-lds-load-insts")
167167

168+
TARGET_BUILTIN(__builtin_amdgcn_global_load_b128, "V4UiV4Ui*1cC*", "n", "gfx940-insts")
169+
TARGET_BUILTIN(__builtin_amdgcn_global_store_b128, "vV4Ui*1V4UicC*", "n", "gfx940-insts")
170+
168171
//===----------------------------------------------------------------------===//
169172
// Ballot builtins.
170173
//===----------------------------------------------------------------------===//

‎clang/include/clang/Sema/SemaAMDGPU.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,8 @@ class SemaAMDGPU : public SemaBase {
3030

3131
bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
3232

33+
bool checkScopedMemAccessFunctionCall(CallExpr *TheCall);
34+
3335
bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
3436
unsigned NumDataArgs);
3537

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20116,6 +20116,26 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
2011620116
llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy});
2011720117
return Builder.CreateCall(F, {Addr});
2011820118
}
20119+
case AMDGPU::BI__builtin_amdgcn_global_load_b128:
20120+
case AMDGPU::BI__builtin_amdgcn_global_store_b128: {
20121+
const bool IsStore =
20122+
BuiltinID == AMDGPU::BI__builtin_amdgcn_global_store_b128;
20123+
LLVMContext &Ctx = CGM.getLLVMContext();
20124+
SmallVector<Value *, 5> Args = {EmitScalarExpr(E->getArg(0))}; // addr
20125+
if (IsStore)
20126+
Args.push_back(EmitScalarExpr(E->getArg(1))); // data
20127+
const unsigned ScopeIdx = E->getNumArgs() - 1;
20128+
StringRef ScopeLit =
20129+
cast<StringLiteral>(E->getArg(ScopeIdx)->IgnoreParenCasts())
20130+
->getString();
20131+
llvm::MDNode *MD =
20132+
llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, ScopeLit)});
20133+
Args.push_back(llvm::MetadataAsValue::get(Ctx, MD)); // scope
20134+
llvm::Function *F =
20135+
CGM.getIntrinsic(IsStore ? Intrinsic::amdgcn_global_store_b128
20136+
: Intrinsic::amdgcn_global_load_b128);
20137+
return Builder.CreateCall(F, Args);
20138+
}
2011920139
case AMDGPU::BI__builtin_amdgcn_get_fpenv: {
2012020140
Function *F = CGM.getIntrinsic(Intrinsic::get_fpenv,
2012120141
{llvm::Type::getInt64Ty(getLLVMContext())});

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -84,6 +84,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
8484
case AMDGPU::BI__builtin_amdgcn_update_dpp: {
8585
return checkMovDPPFunctionCall(TheCall, 6, 2);
8686
}
87+
case AMDGPU::BI__builtin_amdgcn_global_load_b128:
88+
case AMDGPU::BI__builtin_amdgcn_global_store_b128:
89+
return checkScopedMemAccessFunctionCall(TheCall);
8790
default:
8891
return false;
8992
}
@@ -129,6 +132,19 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
129132
return false;
130133
}
131134

135+
bool SemaAMDGPU::checkScopedMemAccessFunctionCall(CallExpr *TheCall) {
136+
bool Fail = false;
137+
// Last argument is a string literal
138+
Expr *Arg = TheCall->getArg(TheCall->getNumArgs() - 1);
139+
auto Scope = dyn_cast<StringLiteral>(Arg->IgnoreParenCasts());
140+
if (!Scope) {
141+
Fail = true;
142+
Diag(TheCall->getBeginLoc(), diag::err_expr_not_string_literal)
143+
<< Arg->getSourceRange();
144+
}
145+
return Fail;
146+
}
147+
132148
bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
133149
unsigned NumDataArgs) {
134150
assert(NumDataArgs <= 2);

‎clang/test/CodeGenOpenCL/builtins-amdgcn-global-load-store.cl‎

Lines changed: 99 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,99 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals smart
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx950 -emit-llvm -o - %s | FileCheck %s -check-prefixes=GFX,GFX950
4+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx9-4-generic -emit-llvm -o - %s | FileCheck %s -check-prefixes=GFX,GFX9_4_GENERIC
5+
6+
typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u32;
7+
typedef v4u32 __global *global_ptr_to_v4u32;
8+
9+
//------------------------------------------------------------------------------
10+
// Store
11+
//------------------------------------------------------------------------------
12+
// GFX-LABEL: @test_amdgcn_global_store_b128_00(
13+
// GFX-NEXT: entry:
14+
// GFX-NEXT: tail call void @llvm.amdgcn.global.store.b128(ptr addrspace(1) [[PTR:%.*]], <4 x i32> [[DATA:%.*]], metadata [[META4:![0-9]+]])
15+
// GFX-NEXT: ret void
16+
//
17+
void test_amdgcn_global_store_b128_00(global_ptr_to_v4u32 ptr, v4u32 data) {
18+
__builtin_amdgcn_global_store_b128(ptr, data, "wavefront");
19+
}
20+
21+
// GFX-LABEL: @test_amdgcn_global_store_b128_01(
22+
// GFX-NEXT: entry:
23+
// GFX-NEXT: tail call void @llvm.amdgcn.global.store.b128(ptr addrspace(1) [[PTR:%.*]], <4 x i32> [[DATA:%.*]], metadata [[META5:![0-9]+]])
24+
// GFX-NEXT: ret void
25+
//
26+
void test_amdgcn_global_store_b128_01(global_ptr_to_v4u32 ptr, v4u32 data) {
27+
__builtin_amdgcn_global_store_b128(ptr, data, "workgroup");
28+
}
29+
30+
// GFX-LABEL: @test_amdgcn_global_store_b128_10(
31+
// GFX-NEXT: entry:
32+
// GFX-NEXT: tail call void @llvm.amdgcn.global.store.b128(ptr addrspace(1) [[PTR:%.*]], <4 x i32> [[DATA:%.*]], metadata [[META6:![0-9]+]])
33+
// GFX-NEXT: ret void
34+
//
35+
void test_amdgcn_global_store_b128_10(global_ptr_to_v4u32 ptr, v4u32 data) {
36+
__builtin_amdgcn_global_store_b128(ptr, data, "agent");
37+
}
38+
39+
// GFX-LABEL: @test_amdgcn_global_store_b128_11(
40+
// GFX-NEXT: entry:
41+
// GFX-NEXT: tail call void @llvm.amdgcn.global.store.b128(ptr addrspace(1) [[PTR:%.*]], <4 x i32> [[DATA:%.*]], metadata [[META7:![0-9]+]])
42+
// GFX-NEXT: ret void
43+
//
44+
void test_amdgcn_global_store_b128_11(global_ptr_to_v4u32 ptr, v4u32 data) {
45+
__builtin_amdgcn_global_store_b128(ptr, data, "");
46+
}
47+
48+
//------------------------------------------------------------------------------
49+
// Load
50+
//------------------------------------------------------------------------------
51+
// GFX-LABEL: @test_amdgcn_global_load_b128_00(
52+
// GFX-NEXT: entry:
53+
// GFX-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.b128(ptr addrspace(1) [[PTR:%.*]], metadata [[META4]])
54+
// GFX-NEXT: ret <4 x i32> [[TMP0]]
55+
//
56+
v4u32 test_amdgcn_global_load_b128_00(global_ptr_to_v4u32 ptr) {
57+
return __builtin_amdgcn_global_load_b128(ptr, "wavefront");
58+
}
59+
60+
// GFX-LABEL: @test_amdgcn_global_load_b128_01(
61+
// GFX-NEXT: entry:
62+
// GFX-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.b128(ptr addrspace(1) [[PTR:%.*]], metadata [[META5]])
63+
// GFX-NEXT: ret <4 x i32> [[TMP0]]
64+
//
65+
v4u32 test_amdgcn_global_load_b128_01(global_ptr_to_v4u32 ptr) {
66+
return __builtin_amdgcn_global_load_b128(ptr, "workgroup");
67+
}
68+
69+
// GFX-LABEL: @test_amdgcn_global_load_b128_10(
70+
// GFX-NEXT: entry:
71+
// GFX-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.b128(ptr addrspace(1) [[PTR:%.*]], metadata [[META6]])
72+
// GFX-NEXT: ret <4 x i32> [[TMP0]]
73+
//
74+
v4u32 test_amdgcn_global_load_b128_10(global_ptr_to_v4u32 ptr) {
75+
return __builtin_amdgcn_global_load_b128(ptr, "agent");
76+
}
77+
78+
// GFX-LABEL: @test_amdgcn_global_load_b128_11(
79+
// GFX-NEXT: entry:
80+
// GFX-NEXT: [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.global.load.b128(ptr addrspace(1) [[PTR:%.*]], metadata [[META7]])
81+
// GFX-NEXT: ret <4 x i32> [[TMP0]]
82+
//
83+
v4u32 test_amdgcn_global_load_b128_11(global_ptr_to_v4u32 ptr) {
84+
return __builtin_amdgcn_global_load_b128(ptr, "");
85+
}
86+
//.
87+
// GFX950: [[META4]] = !{!"wavefront"}
88+
// GFX950: [[META5]] = !{!"workgroup"}
89+
// GFX950: [[META6]] = !{!"agent"}
90+
// GFX950: [[META7]] = !{!""}
91+
//.
92+
// GFX9_4_GENERIC: [[META4]] = !{!"wavefront"}
93+
// GFX9_4_GENERIC: [[META5]] = !{!"workgroup"}
94+
// GFX9_4_GENERIC: [[META6]] = !{!"agent"}
95+
// GFX9_4_GENERIC: [[META7]] = !{!""}
96+
//.
97+
//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
98+
// GFX950: {{.*}}
99+
// GFX9_4_GENERIC: {{.*}}

‎clang/test/SemaOpenCL/builtins-amdgcn-global-load-store-error.cl‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx950 -S -verify -o - %s
2+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx9-4-generic -S -verify -o - %s
3+
// REQUIRES: amdgpu-registered-target
4+
5+
typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u32;
6+
typedef v4u32 __global *global_ptr_to_v4u32;
7+
8+
void test_amdgcn_global_store_b128_00(v4u32 *ptr, v4u32 data, const char* scope) {
9+
__builtin_amdgcn_global_store_b128(ptr, data, ""); //expected-error{{passing '__private v4u32 *__private' to parameter of type '__attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int __global *' changes address space of pointer}}
10+
}
11+
12+
void test_amdgcn_global_store_b128_01(global_ptr_to_v4u32 ptr, v4u32 data, const char* scope) {
13+
__builtin_amdgcn_global_store_b128(ptr, data, scope); //expected-error{{expression is not a string literal}}
14+
}
15+
16+
v4u32 test_amdgcn_global_load_b128_00(v4u32 *ptr, const char* scope) {
17+
return __builtin_amdgcn_global_load_b128(ptr, ""); //expected-error{{passing '__private v4u32 *__private' to parameter of type '__attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int __global *' changes address space of pointer}}
18+
}
19+
20+
v4u32 test_amdgcn_global_load_b128_01(global_ptr_to_v4u32 ptr, const char* scope) {
21+
return __builtin_amdgcn_global_load_b128(ptr, scope); //expected-error{{expression is not a string literal}}
22+
}

‎clang/test/SemaOpenCL/builtins-amdgcn-global-load-store-target-error.cl‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
// We test loads and stores separately because clang only seems to exit after
2+
// the first 'target feature' error.
3+
4+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx9-generic -DTEST_LOAD -S -verify -o - %s
5+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx10-1-generic -DTEST_LOAD -S -verify -o - %s
6+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx10-3-generic -DTEST_LOAD -S -verify -o - %s
7+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx11-generic -DTEST_LOAD -S -verify -o - %s
8+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx12-generic -DTEST_LOAD -S -verify -o - %s
9+
10+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx9-generic -DTEST_STORE -S -verify -o - %s
11+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx10-1-generic -DTEST_STORE -S -verify -o - %s
12+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx10-3-generic -DTEST_STORE -S -verify -o - %s
13+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx11-generic -DTEST_STORE -S -verify -o - %s
14+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx12-generic -DTEST_STORE -S -verify -o - %s
15+
// REQUIRES: amdgpu-registered-target
16+
17+
typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u32;
18+
typedef v4u32 __global *global_ptr_to_v4u32;
19+
20+
#ifdef TEST_LOAD
21+
v4u32 test_amdgcn_global_load_b128_01(global_ptr_to_v4u32 ptr, const char* scope) {
22+
return __builtin_amdgcn_global_load_b128(ptr, ""); // expected-error{{'__builtin_amdgcn_global_load_b128' needs target feature gfx940-insts}}
23+
}
24+
#endif
25+
26+
#ifdef TEST_STORE
27+
void test_amdgcn_global_store_b128_01(global_ptr_to_v4u32 ptr, v4u32 data, const char* scope) {
28+
__builtin_amdgcn_global_store_b128(ptr, data, ""); // expected-error{{'__builtin_amdgcn_global_store_b128' needs target feature gfx940-insts}}
29+
}
30+
#endif

‎llvm/docs/AMDGPUUsage.rst‎

Lines changed: 79 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1436,6 +1436,85 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
14361436
- `v_mov_b32 <dest> <old>`
14371437
- `v_mov_b32 <dest> <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>`
14381438

1439+
llvm.amdgcn.global.load.b128 This intrinsic is supported on gfx942, gfx950.
1440+
1441+
Signature:
1442+
1443+
.. code-block:: llvm
1444+
1445+
<4 x i32> @llvm.amdgcn.raw.load.store.b128(
1446+
ptr addrspace(1), ; source
1447+
metadata) ; scope - e.g. '!0' where '!0 = !{!"wavegroup"}'
1448+
1449+
Reads the value from the source address with cache behavior
1450+
specified by the scope.
1451+
1452+
For gfc942 and gfx950 devices, this emits a
1453+
``global_load_dwordx4`` instruction with the appropriate
1454+
``SC0`` and ``SC1`` bits set.
1455+
1456+
Valid values for scope are
1457+
1458+
===================== =============================================================
1459+
scope architecture name
1460+
===================== =============================================================
1461+
``"wavefront"`` wave
1462+
1463+
``"workgroup"`` group
1464+
1465+
``"agent"`` device
1466+
1467+
``""`` (empty string) system
1468+
===================== =============================================================
1469+
1470+
For semantics on gfx942, see Table 47 in section 9.1.10
1471+
"Memory Scope and Temporal Controls" of the "AMD Instinct
1472+
MI300" Instruction Set Architecture Reference.
1473+
1474+
For semantics on gfx950, see Table 49 in section 9.1.10
1475+
"Memory Scope and Temporal Controls" of the CDNA4
1476+
Instruction Set Architecture Reference.
1477+
1478+
llvm.amdgcn.global.store.b128 This intrinsic is supported on gfx942, gfx950.
1479+
1480+
Signature:
1481+
1482+
.. code-block:: llvm
1483+
1484+
void @llvm.amdgcn.global.store.b128(
1485+
ptr addrspace(1), ; destination
1486+
<4 x i32>, ; value
1487+
metadata) ; scope - e.g. '!0' where '!0 = !{!"wavegroup"}'
1488+
1489+
Writes the value to the destination address with cache
1490+
behavior specified by the scope.
1491+
1492+
For gfc942 and gfx950 devices, this emits a
1493+
``global_store_dwordx4`` instruction with the appropriate
1494+
``SC0`` and ``SC1`` bits set.
1495+
1496+
Valid values for scope are
1497+
1498+
===================== =============================================================
1499+
scope architecture name
1500+
===================== =============================================================
1501+
``"wavefront"`` wave
1502+
1503+
``"workgroup"`` group
1504+
1505+
``"agent"`` device
1506+
1507+
``""`` (empty string) system
1508+
===================== =============================================================
1509+
1510+
For semantics on gfx942, see Table 48 in section 9.1.10
1511+
"Memory Scope and Temporal Controls" of the "AMD Instinct
1512+
MI300" Instruction Set Architecture Reference.
1513+
1514+
For semantics on gfx950, see Table 50 in section 9.1.10
1515+
"Memory Scope and Temporal Controls" of the CDNA4
1516+
Instruction Set Architecture Reference.
1517+
14391518
============================================== ==========================================================
14401519

14411520
.. TODO::

‎llvm/include/llvm/CodeGen/GlobalISel/GIMatchTableExecutor.h‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -239,6 +239,12 @@ enum {
239239
/// - SizeInBits(ULEB128) - The size of the pointer value in bits.
240240
GIM_CheckPointerToAny,
241241

242+
/// Check the machine type of the specified operand
243+
/// - InsnID(ULEB128) - Instruction ID
244+
/// - OpIdx(ULEB128) - Operand index
245+
/// - MachineOperandType(ULEB128) - Expected type
246+
GIM_CheckMachineOperandType,
247+
242248
/// Check the register bank for the specified operand
243249
/// - InsnID(ULEB128) - Instruction ID
244250
/// - OpIdx(ULEB128) - Operand index

0 commit comments

Comments
 (0)