Skip to content

Commit 6e3e14f

Browse files
rampitecrampitec
committed
[AMDGPU] Support gfx940 smfmac instructions
Differential Revision: https://reviews.llvm.org/D122191
1 parent dd67e69 commit 6e3e14f

24 files changed

+768
-6
lines changed

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -309,6 +309,12 @@ TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x32_i8, "V4iWiWiV4iIiIiIi", "nc",
309309
TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x16_i8, "V16iWiWiV16iIiIiIi", "nc", "mai-insts")
310310
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x8_xf32, "V4fV2fV2fV4fIiIiIi", "nc", "mai-insts")
311311
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4_xf32, "V16fV2fV2fV16fIiIiIi", "nc", "mai-insts")
312+
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x32_f16, "V4fV4hV8hV4fiIiIi", "nc", "mai-insts")
313+
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x16_f16, "V16fV4hV8hV16fiIiIi", "nc", "mai-insts")
314+
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x32_bf16, "V4fV4sV8sV4fiIiIi", "nc", "mai-insts")
315+
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x16_bf16, "V16fV4sV8sV16fiIiIi", "nc", "mai-insts")
316+
TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_16x16x64_i8, "V4iV2iV4iV4iiIiIi", "nc", "mai-insts")
317+
TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_32x32x32_i8, "V16iV2iV4iV16iiIiIi", "nc", "mai-insts")
312318

313319
#undef BUILTIN
314320
#undef TARGET_BUILTIN

‎clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,13 +10,16 @@ typedef float v4f __attribute__((ext_vector_type(4)));
1010
typedef float v16f __attribute__((ext_vector_type(16)));
1111
typedef float v32f __attribute__((ext_vector_type(32)));
1212
typedef half v4h __attribute__((ext_vector_type(4)));
13+
typedef half v8h __attribute__((ext_vector_type(8)));
1314
typedef half v16h __attribute__((ext_vector_type(16)));
1415
typedef half v32h __attribute__((ext_vector_type(32)));
16+
typedef int v2i __attribute__((ext_vector_type(2)));
1517
typedef int v4i __attribute__((ext_vector_type(4)));
1618
typedef int v16i __attribute__((ext_vector_type(16)));
1719
typedef int v32i __attribute__((ext_vector_type(32)));
1820
typedef short v2s __attribute__((ext_vector_type(2)));
1921
typedef short v4s __attribute__((ext_vector_type(4)));
22+
typedef short v8s __attribute__((ext_vector_type(8)));
2023
typedef short v16s __attribute__((ext_vector_type(16)));
2124
typedef short v32s __attribute__((ext_vector_type(32)));
2225
typedef double v4d __attribute__((ext_vector_type(4)));
@@ -247,4 +250,46 @@ void test_mfma_f32_32x32x4_xf32(global v16f* out, v2f a, v2f b, v16f c)
247250
{
248251
*out = __builtin_amdgcn_mfma_f32_32x32x4_xf32(a, b, c, 0, 0, 0);
249252
}
253+
254+
// CHECK-GFX940-LABEL: @test_smfmac_f32_16x16x32_f16
255+
// CHECK-GFX940: call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x32.f16(<4 x half> %a, <8 x half> %b, <4 x float> %c, i32 %idx, i32 0, i32 0)
256+
void test_smfmac_f32_16x16x32_f16(global v4f* out, v4h a, v8h b, v4f c, int idx)
257+
{
258+
*out = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a, b, c, idx, 0, 0);
259+
}
260+
261+
// CHECK-GFX940-LABEL: @test_smfmac_f32_32x32x16_f16
262+
// CHECK-GFX940: call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x16.f16(<4 x half> %a, <8 x half> %b, <16 x float> %c, i32 %idx, i32 0, i32 0)
263+
void test_smfmac_f32_32x32x16_f16(global v16f* out, v4h a, v8h b, v16f c, int idx)
264+
{
265+
*out = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a, b, c, idx, 0, 0);
266+
}
267+
268+
// CHECK-GFX940-LABEL: @test_smfmac_f32_16x16x32_bf16
269+
// CHECK-GFX940: call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x32.bf16(<4 x i16> %a, <8 x i16> %b, <4 x float> %c, i32 %idx, i32 0, i32 0)
270+
void test_smfmac_f32_16x16x32_bf16(global v4f* out, v4s a, v8s b, v4f c, int idx)
271+
{
272+
*out = __builtin_amdgcn_smfmac_f32_16x16x32_bf16(a, b, c, idx, 0, 0);
273+
}
274+
275+
// CHECK-GFX940-LABEL: @test_smfmac_f32_32x32x16_bf16
276+
// CHECK-GFX940: call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x16.bf16(<4 x i16> %a, <8 x i16> %b, <16 x float> %c, i32 %idx, i32 0, i32 0)
277+
void test_smfmac_f32_32x32x16_bf16(global v16f* out, v4s a, v8s b, v16f c, int idx)
278+
{
279+
*out = __builtin_amdgcn_smfmac_f32_32x32x16_bf16(a, b, c, idx, 0, 0);
280+
}
281+
282+
// CHECK-GFX940-LABEL: @test_smfmac_i32_16x16x64_i8
283+
// CHECK-GFX940: call <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x64.i8(<2 x i32> %a, <4 x i32> %b, <4 x i32> %c, i32 %idx, i32 0, i32 0)
284+
void test_smfmac_i32_16x16x64_i8(global v4i* out, v2i a, v4i b, v4i c, int idx)
285+
{
286+
*out = __builtin_amdgcn_smfmac_i32_16x16x64_i8(a, b, c, idx, 0, 0);
287+
}
288+
289+
// CHECK-GFX940-LABEL: @test_smfmac_i32_32x32x32_i8
290+
// CHECK-GFX940: call <16 x i32> @llvm.amdgcn.smfmac.i32.32x32x32.i8(<2 x i32> %a, <4 x i32> %b, <16 x i32> %c, i32 %idx, i32 0, i32 0)
291+
void test_smfmac_i32_32x32x32_i8(global v16i* out, v2i a, v4i b, v16i c, int idx)
292+
{
293+
*out = __builtin_amdgcn_smfmac_i32_32x32x32_i8(a, b, c, idx, 0, 0);
294+
}
250295
#endif // MFMA_GFX940_TESTS

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx940-param.cl

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,8 +3,13 @@
33
typedef float v2f __attribute__((ext_vector_type(2)));
44
typedef float v4f __attribute__((ext_vector_type(4)));
55
typedef float v16f __attribute__((ext_vector_type(16)));
6+
typedef int v2i __attribute__((ext_vector_type(2)));
67
typedef int v4i __attribute__((ext_vector_type(4)));
78
typedef int v16i __attribute__((ext_vector_type(16)));
9+
typedef half v4h __attribute__((ext_vector_type(4)));
10+
typedef half v8h __attribute__((ext_vector_type(8)));
11+
typedef short v4s __attribute__((ext_vector_type(4)));
12+
typedef short v8s __attribute__((ext_vector_type(8)));
813

914
void test_mfma_i32_16x16x32i8(global v4i* out, long a, long b, v4i c, int d)
1015
{
@@ -33,3 +38,39 @@ void test_mfma_f32_32x32x4xf32(global v16f* out, v2f a, v2f b, v16f c, int d)
3338
*out = __builtin_amdgcn_mfma_f32_32x32x4_xf32(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4_xf32' must be a constant integer}}
3439
*out = __builtin_amdgcn_mfma_f32_32x32x4_xf32(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4_xf32' must be a constant integer}}
3540
}
41+
42+
void test_smfmac_f32_16x16x32_f16(global v4f* out, v4h a, v8h b, v4f c, int idx, int d)
43+
{
44+
*out = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_16x16x32_f16' must be a constant integer}}
45+
*out = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_16x16x32_f16' must be a constant integer}}
46+
}
47+
48+
void test_smfmac_f32_32x32x16_f16(global v16f* out, v4h a, v8h b, v16f c, int idx, int d)
49+
{
50+
*out = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_32x32x16_f16' must be a constant integer}}
51+
*out = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_32x32x16_f16' must be a constant integer}}
52+
}
53+
54+
void test_smfmac_f32_16x16x32_bf16(global v4f* out, v4s a, v8s b, v4f c, int idx, int d)
55+
{
56+
*out = __builtin_amdgcn_smfmac_f32_16x16x32_bf16(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_16x16x32_bf16' must be a constant integer}}
57+
*out = __builtin_amdgcn_smfmac_f32_16x16x32_bf16(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_16x16x32_bf16' must be a constant integer}}
58+
}
59+
60+
void test_smfmac_f32_32x32x16_bf16(global v16f* out, v4s a, v8s b, v16f c, int idx, int d)
61+
{
62+
*out = __builtin_amdgcn_smfmac_f32_32x32x16_bf16(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_32x32x16_bf16' must be a constant integer}}
63+
*out = __builtin_amdgcn_smfmac_f32_32x32x16_bf16(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_f32_32x32x16_bf16' must be a constant integer}}
64+
}
65+
66+
void test_smfmac_i32_16x16x64_i8(global v4i* out, v2i a, v4i b, v4i c, int idx, int d)
67+
{
68+
*out = __builtin_amdgcn_smfmac_i32_16x16x64_i8(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_i32_16x16x64_i8' must be a constant integer}}
69+
*out = __builtin_amdgcn_smfmac_i32_16x16x64_i8(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_i32_16x16x64_i8' must be a constant integer}}
70+
}
71+
72+
void test_smfmac_i32_32x32x32_i8(global v16i* out, v2i a, v4i b, v16i c, int idx, int d)
73+
{
74+
*out = __builtin_amdgcn_smfmac_i32_32x32x32_i8(a, b, c, idx, d, 0); // expected-error{{argument to '__builtin_amdgcn_smfmac_i32_32x32x32_i8' must be a constant integer}}
75+
*out = __builtin_amdgcn_smfmac_i32_32x32x32_i8(a, b, c, idx, 0, d); // expected-error{{argument to '__builtin_amdgcn_smfmac_i32_32x32x32_i8' must be a constant integer}}
76+
}

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2004,6 +2004,22 @@ def int_amdgcn_mfma_i32_32x32x16_i8 : AMDGPUMfmaIntrinsic<llvm_v16i32_ty, ll
20042004
def int_amdgcn_mfma_f32_16x16x8_xf32 : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_v2f32_ty>;
20052005
def int_amdgcn_mfma_f32_32x32x4_xf32 : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_v2f32_ty>;
20062006

2007+
// llvm.amdgcn.smfmac.?32.* vdst, srcA, srcB, srcC, index, cbsz, abid
2008+
class AMDGPUMSmfmacIntrinsic<LLVMType DestTy, LLVMType SrcA, LLVMType SrcB> :
2009+
GCCBuiltin<!subst("int", "__builtin", NAME)>,
2010+
Intrinsic<[DestTy],
2011+
[SrcA, SrcB, DestTy, llvm_i32_ty,
2012+
llvm_i32_ty, llvm_i32_ty],
2013+
[IntrConvergent, IntrNoMem, IntrWillReturn,
2014+
ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
2015+
2016+
def int_amdgcn_smfmac_f32_16x16x32_f16 : AMDGPUMSmfmacIntrinsic<llvm_v4f32_ty, llvm_v4f16_ty, llvm_v8f16_ty>;
2017+
def int_amdgcn_smfmac_f32_32x32x16_f16 : AMDGPUMSmfmacIntrinsic<llvm_v16f32_ty, llvm_v4f16_ty, llvm_v8f16_ty>;
2018+
def int_amdgcn_smfmac_f32_16x16x32_bf16 : AMDGPUMSmfmacIntrinsic<llvm_v4f32_ty, llvm_v4i16_ty, llvm_v8i16_ty>;
2019+
def int_amdgcn_smfmac_f32_32x32x16_bf16 : AMDGPUMSmfmacIntrinsic<llvm_v16f32_ty, llvm_v4i16_ty, llvm_v8i16_ty>;
2020+
def int_amdgcn_smfmac_i32_16x16x64_i8 : AMDGPUMSmfmacIntrinsic<llvm_v4i32_ty, llvm_v2i32_ty, llvm_v4i32_ty>;
2021+
def int_amdgcn_smfmac_i32_32x32x32_i8 : AMDGPUMSmfmacIntrinsic<llvm_v16i32_ty, llvm_v2i32_ty, llvm_v4i32_ty>;
2022+
20072023
//===----------------------------------------------------------------------===//
20082024
// Special Intrinsics for backend internal use only. No frontend
20092025
// should emit calls to these.

‎llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -971,6 +971,13 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC(MachineInstr &I) const {
971971
return selectGroupStaticSize(I);
972972
case Intrinsic::returnaddress:
973973
return selectReturnAddress(I);
974+
case Intrinsic::amdgcn_smfmac_f32_16x16x32_f16:
975+
case Intrinsic::amdgcn_smfmac_f32_32x32x16_f16:
976+
case Intrinsic::amdgcn_smfmac_f32_16x16x32_bf16:
977+
case Intrinsic::amdgcn_smfmac_f32_32x32x16_bf16:
978+
case Intrinsic::amdgcn_smfmac_i32_16x16x64_i8:
979+
case Intrinsic::amdgcn_smfmac_i32_32x32x32_i8:
980+
return selectSMFMACIntrin(I);
974981
default:
975982
return selectImpl(I, *CoverageInfo);
976983
}
@@ -3054,6 +3061,41 @@ bool AMDGPUInstructionSelector::selectBVHIntrinsic(MachineInstr &MI) const{
30543061
return true;
30553062
}
30563063

3064+
bool AMDGPUInstructionSelector::selectSMFMACIntrin(MachineInstr &MI) const {
3065+
unsigned Opc;
3066+
switch (MI.getIntrinsicID()) {
3067+
case Intrinsic::amdgcn_smfmac_f32_16x16x32_f16:
3068+
Opc = AMDGPU::V_SMFMAC_F32_16X16X32_F16_e64;
3069+
break;
3070+
case Intrinsic::amdgcn_smfmac_f32_32x32x16_f16:
3071+
Opc = AMDGPU::V_SMFMAC_F32_32X32X16_F16_e64;
3072+
break;
3073+
case Intrinsic::amdgcn_smfmac_f32_16x16x32_bf16:
3074+
Opc = AMDGPU::V_SMFMAC_F32_16X16X32_BF16_e64;
3075+
break;
3076+
case Intrinsic::amdgcn_smfmac_f32_32x32x16_bf16:
3077+
Opc = AMDGPU::V_SMFMAC_F32_32X32X16_BF16_e64;
3078+
break;
3079+
case Intrinsic::amdgcn_smfmac_i32_16x16x64_i8:
3080+
Opc = AMDGPU::V_SMFMAC_I32_16X16X64_I8_e64;
3081+
break;
3082+
case Intrinsic::amdgcn_smfmac_i32_32x32x32_i8:
3083+
Opc = AMDGPU::V_SMFMAC_I32_32X32X32_I8_e64;
3084+
break;
3085+
default:
3086+
llvm_unreachable("unhandled smfmac intrinsic");
3087+
}
3088+
3089+
auto VDst_In = MI.getOperand(4);
3090+
3091+
MI.setDesc(TII.get(Opc));
3092+
MI.removeOperand(4); // VDst_In
3093+
MI.removeOperand(1); // Intrinsic ID
3094+
MI.addOperand(VDst_In); // Readd VDst_In to the end
3095+
MI.addImplicitDefUseOperands(*MI.getParent()->getParent());
3096+
return true;
3097+
}
3098+
30573099
bool AMDGPUInstructionSelector::selectWaveAddress(MachineInstr &MI) const {
30583100
Register DstReg = MI.getOperand(0).getReg();
30593101
Register SrcReg = MI.getOperand(1).getReg();

‎llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -144,6 +144,7 @@ class AMDGPUInstructionSelector final : public InstructionSelector {
144144
bool selectGlobalAtomicFadd(MachineInstr &I, MachineOperand &AddrOp,
145145
MachineOperand &DataOp) const;
146146
bool selectBVHIntrinsic(MachineInstr &I) const;
147+
bool selectSMFMACIntrin(MachineInstr &I) const;
147148
bool selectWaveAddress(MachineInstr &I) const;
148149

149150
std::pair<Register, unsigned> selectVOP3ModsImpl(MachineOperand &Root,

‎llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4267,6 +4267,20 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
42674267
: getVGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI);
42684268
break;
42694269
}
4270+
case Intrinsic::amdgcn_smfmac_f32_16x16x32_f16:
4271+
case Intrinsic::amdgcn_smfmac_f32_32x32x16_f16:
4272+
case Intrinsic::amdgcn_smfmac_f32_16x16x32_bf16:
4273+
case Intrinsic::amdgcn_smfmac_f32_32x32x16_bf16:
4274+
case Intrinsic::amdgcn_smfmac_i32_16x16x64_i8:
4275+
case Intrinsic::amdgcn_smfmac_i32_32x32x32_i8: {
4276+
// vdst, srcA, srcB, srcC, idx
4277+
OpdsMapping[0] = getAGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI);
4278+
OpdsMapping[2] = getVGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);
4279+
OpdsMapping[3] = getVGPROpMapping(MI.getOperand(3).getReg(), MRI, *TRI);
4280+
OpdsMapping[4] = getAGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI);
4281+
OpdsMapping[5] = getVGPROpMapping(MI.getOperand(5).getReg(), MRI, *TRI);
4282+
break;
4283+
}
42704284
case Intrinsic::amdgcn_interp_p1:
42714285
case Intrinsic::amdgcn_interp_p2:
42724286
case Intrinsic::amdgcn_interp_mov:

‎llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -299,6 +299,12 @@ def : SourceOfDivergence<int_amdgcn_mfma_i32_16x16x32_i8>;
299299
def : SourceOfDivergence<int_amdgcn_mfma_i32_32x32x16_i8>;
300300
def : SourceOfDivergence<int_amdgcn_mfma_f32_16x16x8_xf32>;
301301
def : SourceOfDivergence<int_amdgcn_mfma_f32_32x32x4_xf32>;
302+
def : SourceOfDivergence<int_amdgcn_smfmac_f32_16x16x32_f16>;
303+
def : SourceOfDivergence<int_amdgcn_smfmac_f32_32x32x16_f16>;
304+
def : SourceOfDivergence<int_amdgcn_smfmac_f32_16x16x32_bf16>;
305+
def : SourceOfDivergence<int_amdgcn_smfmac_f32_32x32x16_bf16>;
306+
def : SourceOfDivergence<int_amdgcn_smfmac_i32_16x16x64_i8>;
307+
def : SourceOfDivergence<int_amdgcn_smfmac_i32_32x32x32_i8>;
302308

303309
// The dummy boolean output is divergent from the IR's perspective,
304310
// but the mask results are uniform. These produce a divergent and

‎llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -147,6 +147,8 @@ DECODE_OPERAND_REG(AReg_512)
147147
DECODE_OPERAND_REG(AReg_1024)
148148
DECODE_OPERAND_REG(AV_32)
149149
DECODE_OPERAND_REG(AV_64)
150+
DECODE_OPERAND_REG(AV_128)
151+
DECODE_OPERAND_REG(AV_512)
150152

151153
static DecodeStatus decodeOperand_VSrc16(MCInst &Inst,
152154
unsigned Imm,
@@ -996,6 +998,14 @@ MCOperand AMDGPUDisassembler::decodeOperand_AV_64(unsigned Val) const {
996998
return decodeSrcOp(OPW64, Val);
997999
}
9981000

1001+
MCOperand AMDGPUDisassembler::decodeOperand_AV_128(unsigned Val) const {
1002+
return decodeSrcOp(OPW128, Val);
1003+
}
1004+
1005+
MCOperand AMDGPUDisassembler::decodeOperand_AV_512(unsigned Val) const {
1006+
return decodeSrcOp(OPW512, Val);
1007+
}
1008+
9991009
MCOperand AMDGPUDisassembler::decodeOperand_VReg_64(unsigned Val) const {
10001010
return createRegOperand(AMDGPU::VReg_64RegClassID, Val);
10011011
}

‎llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -127,6 +127,8 @@ class AMDGPUDisassembler : public MCDisassembler {
127127
MCOperand decodeOperand_AReg_1024(unsigned Val) const;
128128
MCOperand decodeOperand_AV_32(unsigned Val) const;
129129
MCOperand decodeOperand_AV_64(unsigned Val) const;
130+
MCOperand decodeOperand_AV_128(unsigned Val) const;
131+
MCOperand decodeOperand_AV_512(unsigned Val) const;
130132

131133
enum OpWidthTy {
132134
OPW32,

0 commit comments

Comments
 (0)