Skip to content

Commit 87f8800

Browse files
VigneshwarJVigneshwarJ
committed
Revert "Added neg and abs modifier for src2"
This reverts commit 67a22cc.
1 parent 08a5f93 commit 87f8800

File tree

19 files changed

+164
-982
lines changed

19 files changed

+164
-982
lines changed

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -439,8 +439,8 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-conversion-
439439
//===----------------------------------------------------------------------===//
440440
// GFX950 only builtins.
441441
//===----------------------------------------------------------------------===//
442-
TARGET_BUILTIN(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4, "V4fV8ZiV8ZiV4fIiIiIbIbIiiIii", "nc", "gfx950-insts")
443-
TARGET_BUILTIN(__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4, "V16fV8ZiV8ZiV16fIiIiIbIbIiiIii", "nc", "gfx950-insts")
442+
TARGET_BUILTIN(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4, "V4fV8ZiV8ZiV4fIiIiIiiIii", "nc", "gfx950-insts")
443+
TARGET_BUILTIN(__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4, "V16fV8ZiV8ZiV16fIiIiIiiIii", "nc", "gfx950-insts")
444444

445445
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_f16, "V4fV8hV8hV4fIiIiIi", "nc", "gfx950-insts")
446446
TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_bf16, "V4fV8yV8yV4fIiIiIi", "nc", "gfx950-insts")

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -435,18 +435,18 @@ v16f test_mfma_f32_32x32x16_bf16(v8bf16 a, v8bf16 b, v16f c) {
435435

436436
// CHECK-GFX950-LABEL: @test_mfma_scale_f32_16x16x128_f8f6f4
437437
// CHECK-GFX950: [[EXTRACT_A:%.+]] = shufflevector <8 x i32> %a, <8 x i32> poison, <6 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5>
438-
// CHECK-GFX950: call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> [[EXTRACT_A]], <8 x i32> %b, <4 x float> %c, i32 3, i32 1, i1 false, i1 false, i32 2, i32 %scale_a, i32 3, i32 %scale_b)
438+
// CHECK-GFX950: call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v6i32.v8i32(<6 x i32> [[EXTRACT_A]], <8 x i32> %b, <4 x float> %c, i32 3, i32 1, i32 2, i32 %scale_a, i32 3, i32 %scale_b)
439439
void test_mfma_scale_f32_16x16x128_f8f6f4(global v4f* out, v8i a, v8i b, v4f c, int scale_a, int scale_b)
440440
{
441-
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 3, 1, false, false, 2, scale_a, 3, scale_b);
441+
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 3, 1, 2, scale_a, 3, scale_b);
442442
}
443443

444444
// CHECK-GFX950-LABEL: @test_mfma_scale_f32_32x32x64_f8f6f4
445445
// CHECK-GFX950: [[EXTRACT_A:%.+]] = shufflevector <8 x i32> %a, <8 x i32> poison, <6 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5>
446-
// CHECK-GFX950: call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v6i32.v8i32(<6 x i32> [[EXTRACT_A]], <8 x i32> %b, <16 x float> %c, i32 3, i32 1, i1 false, i1 false, i32 2, i32 %scale_a, i32 3, i32 %scale_b)
446+
// CHECK-GFX950: call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v6i32.v8i32(<6 x i32> [[EXTRACT_A]], <8 x i32> %b, <16 x float> %c, i32 3, i32 1, i32 2, i32 %scale_a, i32 3, i32 %scale_b)
447447
void test_mfma_scale_f32_32x32x64_f8f6f4(global v16f* out, v8i a, v8i b, v16f c, int scale_a, int scale_b)
448448
{
449-
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 3, 1, false, false, 2, scale_a, 3, scale_b);
449+
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 3, 1, 2, scale_a, 3, scale_b);
450450
}
451451

452452
// CHECK-GFX950-LABEL: @test_mfma_i32_16x16x64_i8(

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

Lines changed: 8 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -38,21 +38,17 @@ void test_mfma_f32_32x32x16_bf16(__global float16* out, bfloat8 a, bfloat8 b, fl
3838
}
3939

4040
void test_mfma_scale_f32_16x16x128_f8f6f4(__global float4* out, int8 a, int8 b, float4 c, int X, int Y) {
41-
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, X, 0, false, false, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
42-
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, X, false, false, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
43-
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, 0, X, false, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
44-
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, 0, false, X, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
45-
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, 0, false, false, X, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
46-
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, 0, false, false, 0, Y, X, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
41+
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, X, 0, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
42+
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, X, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
43+
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, 0, X, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
44+
*out = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, 0, 0, 0, Y, X, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' must be a constant integer}}
4745
}
4846

4947
void test_mfma_scale_f32_32x32x64_f8f6f4(__global float16* out, int8 a, int8 b, float16 c, int X, int Y) {
50-
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, X, 0, false, false, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
51-
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, X, false, false, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
52-
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, 0, X, false, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
53-
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, 0, false, X, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
54-
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, 0, false, false, X, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
55-
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, 0, false, false, 0, Y, X, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
48+
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, X, 0, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
49+
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, X, 1, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
50+
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, 0, X, Y, 2, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
51+
*out = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, 0, 0, Y, X, Y); // expected-error{{argument to '__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' must be a constant integer}}
5652
}
5753

5854
void test_mfma_i32_16x16x64_i8(__global int4* out, int4 a, int4 b, int4 c, int X) {

‎clang/test/SemaOpenCL/builtins-amdgcn-error-gfx950.cl‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -50,8 +50,8 @@ void test(__global float4* out0, half8 a0, half8 b0, float4 c0,
5050
*out13 = __builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8(a13, b13, c13, 0, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8' needs target feature gfx950-insts}}
5151
*out13 = __builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8(a13, b13, c13, 0, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8' needs target feature gfx950-insts}}
5252
*out13 = __builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8(a13, b13, c13, 0, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8' needs target feature gfx950-insts}}
53-
*out14 = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a14, b14, c14, 0, 0, false, false, 0, d14, 0, e14); // expected-error{{'__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' needs target feature gfx950-insts}}
54-
*out15 = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a15, b15, c15, 0, 0, false, false, 0, d15, 0, e15); // expected-error{{'__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' needs target feature gfx950-insts}}
53+
*out14 = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a14, b14, c14, 0, 0, 0, d14, 0, e14); // expected-error{{'__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4' needs target feature gfx950-insts}}
54+
*out15 = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a15, b15, c15, 0, 0, 0, d15, 0, e15); // expected-error{{'__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4' needs target feature gfx950-insts}}
5555
*out16 = __builtin_amdgcn_permlane16_swap(a16, b16, false, false); // expected-error{{'__builtin_amdgcn_permlane16_swap' needs target feature permlane16-swap}}
5656
*out16 = __builtin_amdgcn_permlane32_swap(a16, b16, false, false); // expected-error{{'__builtin_amdgcn_permlane32_swap' needs target feature permlane32-swap}}
5757
*out17 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(*out17, a17, b17, c17, 0); // expected-error{{'__builtin_amdgcn_cvt_scalef32_sr_bf8_bf16' needs target feature bf8-cvt-scale-insts}}

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3261,17 +3261,17 @@ class AMDGPUMfmaScaleIntrinsic<LLVMType DestTy> :
32613261
[llvm_anyvector_ty, llvm_anyvector_ty, DestTy,
32623262
llvm_i32_ty, // cbsz
32633263
llvm_i32_ty, // blgp
3264-
llvm_i1_ty, // neg_src2
3265-
llvm_i1_ty, // abs_src2
3264+
// llvm_i1_ty, // TODO: neg_src2
3265+
// llvm_i1_ty, // TODO: abs_src2
3266+
// llvm_i1_ty, // TODO: clamp
32663267
llvm_i32_ty, // op_sel (A matrix scale, 2-bits) // TODO: Make i2?
32673268
llvm_i32_ty, // v_mfma_ld_scale_b32 src0 (A matrix scale)
32683269
llvm_i32_ty, // op_sel (B matrix scale, 2-bits) // TODO: Make i2?
32693270
llvm_i32_ty // v_mfma_ld_scale_b32 src1 (B matrix scale)
32703271
],
32713272
[IntrConvergent, IntrNoMem,
32723273
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>,
3273-
ImmArg<ArgIndex<5>>, ImmArg<ArgIndex<6>>,
3274-
ImmArg<ArgIndex<7>>, ImmArg<ArgIndex<9>>
3274+
ImmArg<ArgIndex<5>>, ImmArg<ArgIndex<7>>
32753275
]>;
32763276

32773277
defset list<Intrinsic> AMDGPUMFMAIntrinsics908 = {

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4891,8 +4891,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
48914891
? getAGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI)
48924892
: getVGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI);
48934893

4894+
OpdsMapping[8] = getVGPROpMapping(MI.getOperand(8).getReg(), MRI, *TRI);
48944895
OpdsMapping[10] = getVGPROpMapping(MI.getOperand(10).getReg(), MRI, *TRI);
4895-
OpdsMapping[12] = getVGPROpMapping(MI.getOperand(12).getReg(), MRI, *TRI);
48964896
break;
48974897
}
48984898
case Intrinsic::amdgcn_smfmac_f32_16x16x32_f16:

‎llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp‎

Lines changed: 6 additions & 39 deletions
Original file line numberDiff line numberDiff line change
@@ -152,8 +152,6 @@ class AMDGPUOperand : public MCParsedAsmOperand {
152152
ImmTyOpSelHi,
153153
ImmTyNegLo,
154154
ImmTyNegHi,
155-
ImmTyNegLoSrc2,
156-
ImmTyNegHiSrc2,
157155
ImmTyIndexKey8bit,
158156
ImmTyIndexKey16bit,
159157
ImmTyDPP8,
@@ -418,8 +416,6 @@ class AMDGPUOperand : public MCParsedAsmOperand {
418416
bool isOpSelHi() const { return isImmTy(ImmTyOpSelHi); }
419417
bool isNegLo() const { return isImmTy(ImmTyNegLo); }
420418
bool isNegHi() const { return isImmTy(ImmTyNegHi); }
421-
bool isNegHiSrc2() const { return isImmTy(ImmTyNegHiSrc2); }
422-
bool isNegLoSrc2() const { return isImmTy(ImmTyNegLoSrc2); }
423419
bool isBitOp3() const { return isImmTy(ImmTyBitOp3) && isUInt<8>(getImm()); }
424420

425421
bool isRegOrImm() const {
@@ -1142,8 +1138,6 @@ class AMDGPUOperand : public MCParsedAsmOperand {
11421138
case ImmTyHigh: OS << "High"; break;
11431139
case ImmTyBLGP: OS << "BLGP"; break;
11441140
case ImmTyCBSZ: OS << "CBSZ"; break;
1145-
case ImmTyNegLoSrc2: OS << "NegSrc2"; break;
1146-
case ImmTyNegHiSrc2: OS << "AbsSrc2"; break;
11471141
case ImmTyABID: OS << "ABID"; break;
11481142
case ImmTyEndpgm: OS << "Endpgm"; break;
11491143
case ImmTyWaitVDST: OS << "WaitVDST"; break;
@@ -1638,7 +1632,7 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
16381632
ParseStatus parseOperandArrayWithPrefix(
16391633
const char *Prefix, OperandVector &Operands,
16401634
AMDGPUOperand::ImmTy ImmTy = AMDGPUOperand::ImmTyNone,
1641-
std::function<bool(int64_t &)> ConvertResult = nullptr);
1635+
bool (*ConvertResult)(int64_t &) = nullptr);
16421636

16431637
ParseStatus
16441638
parseNamedBit(StringRef Name, OperandVector &Operands,
@@ -1693,8 +1687,6 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
16931687
ParseStatus parseFlatOffset(OperandVector &Operands);
16941688
ParseStatus parseR128A16(OperandVector &Operands);
16951689
ParseStatus parseBLGP(OperandVector &Operands);
1696-
ParseStatus parseNegHiSrc2(OperandVector &Operands);
1697-
ParseStatus parseNegLoSrc2(OperandVector &Operands);
16981690
bool tryParseFmt(const char *Pref, int64_t MaxVal, int64_t &Val);
16991691
bool matchDfmtNfmt(int64_t &Dfmt, int64_t &Nfmt, StringRef FormatStr, SMLoc Loc);
17001692

@@ -6556,15 +6548,15 @@ ParseStatus AMDGPUAsmParser::parseIntWithPrefix(
65566548

65576549
ParseStatus AMDGPUAsmParser::parseOperandArrayWithPrefix(
65586550
const char *Prefix, OperandVector &Operands, AMDGPUOperand::ImmTy ImmTy,
6559-
std::function<bool(int64_t &)> ConvertResult) {
6551+
bool (*ConvertResult)(int64_t &)) {
65606552
SMLoc S = getLoc();
65616553
if (!trySkipId(Prefix, AsmToken::Colon))
65626554
return ParseStatus::NoMatch;
65636555

65646556
if (!skipToken(AsmToken::LBrac, "expected a left square bracket"))
65656557
return ParseStatus::Failure;
65666558

6567-
int64_t Val = 0;
6559+
unsigned Val = 0;
65686560
const unsigned MaxSize = 4;
65696561

65706562
// FIXME: How to verify the number of elements matches the number of src
@@ -6589,9 +6581,7 @@ ParseStatus AMDGPUAsmParser::parseOperandArrayWithPrefix(
65896581
if (!skipToken(AsmToken::Comma, "expected a comma"))
65906582
return ParseStatus::Failure;
65916583
}
6592-
if (ConvertResult && !ConvertResult(Val)) {
6593-
Error(S, "invalid " + StringRef(Prefix) + " value.");
6594-
}
6584+
65956585
Operands.push_back(AMDGPUOperand::CreateImm(this, Val, S, ImmTy));
65966586
return ParseStatus::Success;
65976587
}
@@ -7161,23 +7151,6 @@ ParseStatus AMDGPUAsmParser::parseBLGP(OperandVector &Operands) {
71617151
return Res;
71627152
}
71637153

7164-
static bool RightShift2Bits(int64_t &Neg) {
7165-
Neg >>= 2;
7166-
return true;
7167-
}
7168-
7169-
ParseStatus AMDGPUAsmParser::parseNegLoSrc2(OperandVector &Operands) {
7170-
return parseOperandArrayWithPrefix(
7171-
"neg_lo", Operands, AMDGPUOperand::ImmTyNegLoSrc2,
7172-
RightShift2Bits); // Extracting only neg_lo[2]
7173-
}
7174-
7175-
ParseStatus AMDGPUAsmParser::parseNegHiSrc2(OperandVector &Operands) {
7176-
return parseOperandArrayWithPrefix(
7177-
"neg_hi", Operands, AMDGPUOperand::ImmTyNegHiSrc2,
7178-
RightShift2Bits); // Extracting only neg_hi[2]
7179-
}
7180-
71817154
//===----------------------------------------------------------------------===//
71827155
// Exp
71837156
//===----------------------------------------------------------------------===//
@@ -8878,12 +8851,6 @@ void AMDGPUAsmParser::cvtScaledMFMA(MCInst &Inst,
88788851
addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyBLGP,
88798852
0, InsertPos);
88808853

8881-
// add neg and abs for src2
8882-
addOptionalImmOperand(Inst, Operands, OptionalIdx,
8883-
AMDGPUOperand::ImmTyNegLoSrc2, 0);
8884-
addOptionalImmOperand(Inst, Operands, OptionalIdx,
8885-
AMDGPUOperand::ImmTyNegHiSrc2, 0);
8886-
88878854
// Add dummy src_modifiers
88888855
Inst.addOperand(MCOperand::createImm(0));
88898856
Inst.addOperand(MCOperand::createImm(0));
@@ -8907,9 +8874,9 @@ void AMDGPUAsmParser::cvtScaledMFMA(MCInst &Inst,
89078874
for (unsigned J = 0; J < 2; ++J) {
89088875
unsigned ModVal = 0;
89098876
if (OpSel & (1 << J))
8910-
ModVal |= SISrcMods::OP_SEL_0; // 3rd bit is from opsel
8877+
ModVal |= SISrcMods::OP_SEL_0;
89118878
if (OpSelHi & (1 << J))
8912-
ModVal |= SISrcMods::OP_SEL_1; // 4th bit is from opsel_hi
8879+
ModVal |= SISrcMods::OP_SEL_1;
89138880

89148881
const int ModIdx = AMDGPU::getNamedOperandIdx(Opc, ModOps[J]);
89158882
Inst.getOperand(ModIdx).setImm(ModVal);

‎llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUInstPrinter.cpp‎

Lines changed: 0 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1297,20 +1297,6 @@ void AMDGPUInstPrinter::printNegHi(const MCInst *MI, unsigned OpNo,
12971297
printPackedModifier(MI, " neg_hi:[", SISrcMods::NEG_HI, O);
12981298
}
12991299

1300-
void AMDGPUInstPrinter::printNegLoSrc2(const MCInst *MI, unsigned OpNo,
1301-
const MCSubtargetInfo &STI,
1302-
raw_ostream &O) {
1303-
if (unsigned NegLo = !!(MI->getOperand(OpNo).getImm()))
1304-
O << " neg_lo:[0,0," << NegLo << ']';
1305-
}
1306-
1307-
void AMDGPUInstPrinter::printNegHiSrc2(const MCInst *MI, unsigned OpNo,
1308-
const MCSubtargetInfo &STI,
1309-
raw_ostream &O) {
1310-
if (unsigned NegHi = !!(MI->getOperand(OpNo).getImm()))
1311-
O << " neg_hi:[0,0," << NegHi << ']';
1312-
}
1313-
13141300
void AMDGPUInstPrinter::printIndexKey8bit(const MCInst *MI, unsigned OpNo,
13151301
const MCSubtargetInfo &STI,
13161302
raw_ostream &O) {

‎llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUInstPrinter.h‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -126,10 +126,6 @@ class AMDGPUInstPrinter : public MCInstPrinter {
126126
const MCSubtargetInfo &STI, raw_ostream &O);
127127
void printNegHi(const MCInst *MI, unsigned OpNo,
128128
const MCSubtargetInfo &STI, raw_ostream &O);
129-
void printNegLoSrc2(const MCInst *MI, unsigned OpNo,
130-
const MCSubtargetInfo &STI, raw_ostream &O);
131-
void printNegHiSrc2(const MCInst *MI, unsigned OpNo,
132-
const MCSubtargetInfo &STI, raw_ostream &O);
133129
void printIndexKey8bit(const MCInst *MI, unsigned OpNo,
134130
const MCSubtargetInfo &STI, raw_ostream &O);
135131
void printIndexKey16bit(const MCInst *MI, unsigned OpNo,

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

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -310,7 +310,7 @@ def SIdenorm_mode : SDNode<"AMDGPUISD::DENORM_MODE",
310310
class UnscaledMFMAOptimizationPat<SDPatternOperator intrin> : PatFrag<
311311
(ops node:$srca, node:$srcb, node:$srcc,
312312
node:$cbsz, node:$blgp),
313-
(intrin $srca, $srcb, $srcc, $cbsz, $blgp, 0, 0,
313+
(intrin $srca, $srcb, $srcc, $cbsz, $blgp,
314314
srcvalue, 0, srcvalue, 0)
315315
>;
316316

@@ -1244,9 +1244,6 @@ def op_sel_hi0 : ArrayOperand0<"op_sel_hi", "OpSelHi">;
12441244
def neg_lo0 : ArrayOperand0<"neg_lo", "NegLo">;
12451245
def neg_hi0 : ArrayOperand0<"neg_hi", "NegHi">;
12461246

1247-
def neg_lo_src2 : CustomOperand<i1, 1, "NegLoSrc2">;
1248-
def neg_hi_src2 : CustomOperand<i1, 1, "NegHiSrc2">;
1249-
12501247
def IndexKey16bit : CustomOperand<i32, 1>;
12511248
def IndexKey8bit : CustomOperand<i32, 1>;
12521249

0 commit comments

Comments
 (0)