[AMDGPU] w/a for gfx908 mfma SrcC literal HW bug

[rampitec]

Summary

gfx908 ignores an mfma if SrcC is a literal.

[arsenm]

Can we have different opcodes with a different source type?

[rampitec]

In D66670#1643363, @arsenm wrote:

Can we have different opcodes with a different source type?

It will be a quite big change. What for?

[arsenm]

In D66670#1643370, @rampitec wrote:

In D66670#1643363, @arsenm wrote:

Can we have different opcodes with a different source type?

It will be a quite big change. What for?

Generally if the operand/register class changes there should be a separate opcode. It avoids needing all this special case checking, and makes folding code work naturally

[rampitec]

In D66670#1643378, @arsenm wrote:

In D66670#1643370, @rampitec wrote:

In D66670#1643363, @arsenm wrote:

Can we have different opcodes with a different source type?

It will be a quite big change. What for?

Generally if the operand/register class changes there should be a separate opcode. It avoids needing all this special case checking, and makes folding code work naturally

Anyway, you cannot have two opcodes with the same opcode without introducing a new encoding and a new decoder namespace. This bug will be fixed.

[arsenm]

In D66670#1643385, @rampitec wrote:

In D66670#1643378, @arsenm wrote:

In D66670#1643370, @rampitec wrote:

In D66670#1643363, @arsenm wrote:

Can we have different opcodes with a different source type?

It will be a quite big change. What for?

Generally if the operand/register class changes there should be a separate opcode. It avoids needing all this special case checking, and makes folding code work naturally

Anyway, you cannot have two opcodes with the same opcode without introducing a new encoding and a new decoder namespace. This bug will be fixed.

The encoding is the same, and a new decoder namespace shouldn't be an issue

lib/Target/AMDGPU/SIInstrInfo.cpp

2873–2876 ↗

(On Diff #216911)

I would expect this to be in RI.opCanUseInlineConstant

[rampitec]

In D66670#1643408, @arsenm wrote:

In D66670#1643385, @rampitec wrote:

In D66670#1643378, @arsenm wrote:

In D66670#1643370, @rampitec wrote:

In D66670#1643363, @arsenm wrote:

Can we have different opcodes with a different source type?

It will be a quite big change. What for?

Generally if the operand/register class changes there should be a separate opcode. It avoids needing all this special case checking, and makes folding code work naturally

Anyway, you cannot have two opcodes with the same opcode without introducing a new encoding and a new decoder namespace. This bug will be fixed.

The encoding is the same, and a new decoder namespace shouldn't be an issue

I still think the change will be too massive. Note, that encoding is still valid and legal.

[rampitec]

lib/Target/AMDGPU/SIInstrInfo.cpp

2873–2876 ↗

(On Diff #216911)

It does not have a subtarget.

[arsenm]

Can these be encoded as a constant literal instead? We would then just need to workaround the encoding rather than disallowing the operand

[rampitec]

In D66670#1643525, @arsenm wrote:

Can these be encoded as a constant literal instead? We would then just need to workaround the encoding rather than disallowing the operand

It is VOP3, so not in gfx9.

[rampitec]

Moved check into opCanUseInlineConstant().
Captured subtarget in SIRegisterInfo to do this.

[arsenm]

LGTM

llvm/lib/Target/AMDGPU/AMDGPU.td

 >;
 
 def FeatureLdsMisalignedBug : SubtargetFeature<"lds-misaligned-bug",
   "LDSMisalignedBug",
   "true",
   "Some GFX10 bug with misaligned multi-dword LDS access in WGP mode"
 >;
 
+def FeatureMFMAInlineLiteralBug : SubtargetFeature<"mfma-inline-literal-bug",
+  "HasMFMAInlineLiteralBug",
+  "true",
+  "MFMA cannot use inline literal as SrcC"
+>;
+
 def FeatureVcmpxPermlaneHazard : SubtargetFeature<"vcmpx-permlane-hazard",
   "HasVcmpxPermlaneHazard",
   "true",
   "TODO: describe me"
 >;
 
 def FeatureVMEMtoScalarWriteHazard : SubtargetFeature<"vmem-to-scalar-write-hazard",
   "HasVMEMtoScalarWriteHazard",
def FeatureISAVersion9_0_8 : FeatureSet<
    FeatureDot3Insts,
    FeatureDot4Insts,
    FeatureDot5Insts,
    FeatureDot6Insts,
    FeatureMAIInsts,
    FeaturePkFmacF16Inst,
    FeatureAtomicFaddInsts,
    FeatureSRAMECC,
+   FeatureMFMAInlineLiteralBug,
    FeatureCodeObjectV3]>;
 
 def FeatureISAVersion9_0_9 : FeatureSet<
   [FeatureGFX9,
    FeatureMadMixInsts,
    FeatureLDSBankCount32,
    FeatureXNACK,
    FeatureCodeObjectV3]>;

llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h

protected:
   bool FlatScratchInsts;
   bool ScalarFlatScratchInsts;
   bool AddNoCarryInsts;
   bool HasUnpackedD16VMem;
   bool R600ALUInst;
   bool CaymanISA;
   bool CFALUBug;
   bool LDSMisalignedBug;
+  bool HasMFMAInlineLiteralBug;
   bool HasVertexCache;
   short TexVTXClauseSize;
   bool ScalarizeGlobal;
 
   bool HasVcmpxPermlaneHazard;
   bool HasVMEMtoScalarWriteHazard;
   bool HasSMEMtoVectorWriteHazard;
   bool HasInstFwdPrefetchBug;
public:
   bool loadStoreOptEnabled() const {
     return EnableLoadStoreOpt;
   }
 
   bool hasSGPRInitBug() const {
     return SGPRInitBug;
   }
 
+  bool hasMFMAInlineLiteralBug() const {
+    return HasMFMAInlineLiteralBug;
+  }
+
   bool has12DWordStoreHazard() const {
     return getGeneration() != AMDGPUSubtarget::SOUTHERN_ISLANDS;
   }
 
   // \returns true if the subtarget supports DWORDX3 load/store instructions.
   bool hasDwordx3LoadStores() const {
     return CIInsts;
   }

llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp

GCNSubtarget::GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
     FlatAddressSpace(false),
     FlatInstOffsets(false),
     FlatGlobalInsts(false),
     FlatScratchInsts(false),
     ScalarFlatScratchInsts(false),
     AddNoCarryInsts(false),
     HasUnpackedD16VMem(false),
     LDSMisalignedBug(false),
+    HasMFMAInlineLiteralBug(false),
 
     ScalarizeGlobal(false),
 
     HasVcmpxPermlaneHazard(false),
     HasVMEMtoScalarWriteHazard(false),
     HasSMEMtoVectorWriteHazard(false),
     HasInstFwdPrefetchBug(false),
     HasVcmpxExecWARHazard(false),

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

static bool tryToFoldACImm(const SIInstrInfo *TII,
   if (!OpInfo || UseOpIdx >= Desc.getNumOperands())
     return false;
 
   uint8_t OpTy = OpInfo[UseOpIdx].OperandType;
   if (OpTy < AMDGPU::OPERAND_REG_INLINE_AC_FIRST ||
       OpTy > AMDGPU::OPERAND_REG_INLINE_AC_LAST)
     return false;
 
-  if (OpToFold.isImm() && TII->isInlineConstant(OpToFold, OpTy)) {
+  if (OpToFold.isImm() && TII->isInlineConstant(OpToFold, OpTy) &&
+      TII->isOperandLegal(*UseMI, UseOpIdx, &OpToFold)) {
     UseMI->getOperand(UseOpIdx).ChangeToImmediate(OpToFold.getImm());
     return true;
   }
 
   if (!OpToFold.isReg())
     return false;
 
   Register UseReg = OpToFold.getReg();
if (I == 1) {
 
       Imm = SubImm;
       continue;
     }
     if (Imm != SubImm)
       return false; // Can only fold splat constants
   }
 
+  if (!TII->isOperandLegal(*UseMI, UseOpIdx, Op))
+    return false;
+
   FoldList.push_back(FoldCandidate(UseMI, UseOpIdx, Op));
   return true;
 }
 
 void SIFoldOperands::foldOperand(
   MachineOperand &OpToFold,
   MachineInstr *UseMI,
   int UseOpIdx,

llvm/lib/Target/AMDGPU/SIRegisterInfo.h

 
 class GCNSubtarget;
 class LiveIntervals;
 class MachineRegisterInfo;
 class SIMachineFunctionInfo;
 
 class SIRegisterInfo final : public AMDGPURegisterInfo {
 private:
+  const GCNSubtarget &ST;
   unsigned SGPRSetID;
   unsigned VGPRSetID;
   unsigned AGPRSetID;
   BitVector SGPRPressureSets;
   BitVector VGPRPressureSets;
   BitVector AGPRPressureSets;
   bool SpillSGPRToVGPR;
   bool SpillSGPRToSMEM;
bool opCanUseLiteralConstant(unsigned OpType) const {
     // TODO: 64-bit operands have extending behavior from 32-bit literal.
     return OpType >= AMDGPU::OPERAND_REG_IMM_FIRST &&
            OpType <= AMDGPU::OPERAND_REG_IMM_LAST;
   }
 
   /// \returns True if operands defined with this operand type can accept
   /// an inline constant. i.e. An integer value in the range (-16, 64) or
   /// -4.0f, -2.0f, -1.0f, -0.5f, 0.0f, 0.5f, 1.0f, 2.0f, 4.0f.
-  bool opCanUseInlineConstant(unsigned OpType) const {
-    return OpType >= AMDGPU::OPERAND_SRC_FIRST &&
-           OpType <= AMDGPU::OPERAND_SRC_LAST;
-  }
+  bool opCanUseInlineConstant(unsigned OpType) const;
 
   unsigned findUnusedRegister(const MachineRegisterInfo &MRI,
                               const TargetRegisterClass *RC,
                               const MachineFunction &MF) const;
 
   unsigned getSGPRPressureSet() const { return SGPRSetID; };
   unsigned getVGPRPressureSet() const { return VGPRSetID; };
   unsigned getAGPRPressureSet() const { return AGPRSetID; };

llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp

 static cl::opt<bool> EnableSpillSGPRToVGPR(
   "amdgpu-spill-sgpr-to-vgpr",
   cl::desc("Enable spilling VGPRs to SGPRs"),
   cl::ReallyHidden,
   cl::init(true));
 
 SIRegisterInfo::SIRegisterInfo(const GCNSubtarget &ST) :
   AMDGPURegisterInfo(),
+  ST(ST),
   SGPRPressureSets(getNumRegPressureSets()),
   VGPRPressureSets(getNumRegPressureSets()),
   AGPRPressureSets(getNumRegPressureSets()),
   SpillSGPRToVGPR(false),
   SpillSGPRToSMEM(false),
   isWave32(ST.isWave32()) {
   if (EnableSpillSGPRToSMEM && ST.hasScalarStores())
     SpillSGPRToSMEM = true;
case 16:
       return &AMDGPU::VReg_512RegClass;
     case 32: /* fall-through */
     default:
       llvm_unreachable("Invalid sub-register class size");
     }
   }
 }
 
+bool SIRegisterInfo::opCanUseInlineConstant(unsigned OpType) const {
+  if (OpType >= AMDGPU::OPERAND_REG_INLINE_AC_FIRST &&
+      OpType <= AMDGPU::OPERAND_REG_INLINE_AC_LAST)
+    return !ST.hasMFMAInlineLiteralBug();
+
+  return OpType >= AMDGPU::OPERAND_SRC_FIRST &&
+         OpType <= AMDGPU::OPERAND_SRC_LAST;
+}
+
 bool SIRegisterInfo::shouldRewriteCopySrc(
   const TargetRegisterClass *DefRC,
   unsigned DefSubReg,
   const TargetRegisterClass *SrcRC,
   unsigned SrcSubReg) const {
   // We want to prefer the smallest register class possible, so we don't want to
   // stop and rewrite on anything that looks like a subregister
   // extract. Operations mostly don't care about the super register class, so we

llvm/test/CodeGen/AMDGPU/agpr-register-count.ll

 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)
 
 ; GCN-LABEL: {{^}}test_32_agprs:
-; GCN: v_mfma_f32_32x32x1f32 a[0:31], {{v[0-9]+}}, {{v[0-9]+}}, 0
+; GCN: v_mfma_f32_32x32x1f32 a[0:31], {{v[0-9]+}}, {{v[0-9]+}},
 ; GCN-NOT: v28
 ; GCN: NumVgprs: 32
 ; GCN: VGPRBlocks: 7
 define amdgpu_kernel void @test_32_agprs(<32 x float> addrspace(1)* %arg) {
 bb:
   %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> <float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, i32 0, i32 0, i32 0)
   store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
   ret void
 }

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

-; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck --check-prefixes=GCN,NOLIT-SRCC %s
+; RUN: llc -march=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug -verify-machineinstrs < %s | FileCheck --check-prefixes=GCN,LIT-SRCC %s
 
 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)
 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32, i32, i32)
 declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
 declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float, float, <16 x float>, i32, i32, i32)
 declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float, float, <4 x float>, i32, i32, i32)
 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half>, <4 x half>, <32 x float>, i32, i32, i32)
 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half>, <4 x half>, <16 x float>, i32, i32, i32)
bb:
   %mai.2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %mai.1, i32 0, i32 0, i32 0)
   store <4 x float> %mai.2, <4 x float> addrspace(1)* %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_imm_splat:
 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
-; GCN: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
+; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
+; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
+; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
+; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
+; NOLIT-SRCC: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9:]+}}]
+; LIT-SRCC: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
 ; GCN: v_accvgpr_read_b32
 ; GCN: v_accvgpr_read_b32
 ; GCN: v_accvgpr_read_b32
 ; GCN: v_accvgpr_read_b32
 ; GCN: global_store_dwordx4
 define amdgpu_kernel void @test_mfma_f32_4x4x1f32_imm_splat(<4 x float> addrspace(1)* %arg) {
 bb:
   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> <float 1.0, float 1.0, float 1.0, float 1.0>, i32 0, i32 0, i32 0)
   store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32_imm_splat:
 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
-; GCN: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
+; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
+; NOLIT-SRCC: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9:]+}}]
+; LIT-SRCC: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
bb:
   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> <float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0>, i32 0, i32 0, i32 0)
   store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x8f16_imm_splat:
 ; GCN-DAG: v_mov_b32_e32 v[[TWO:[0-9]+]], 0x40004000
 ; GCN-DAG: v_mov_b32_e32 v[[ONE:[0-9]+]], 0x3c003c00
-; GCN: v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], v{{\[}}[[ONE]]:{{[0-9]+}}], v{{\[}}[[TWO]]:{{[0-9]+}}], 1.0
+; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
+; NOLIT-SRCC: v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], v{{\[}}[[ONE]]:{{[0-9]+}}], v{{\[}}[[TWO]]:{{[0-9]+}}], a[{{[0-9:]+}}]
+; LIT-SRCC: v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], v{{\[}}[[ONE]]:{{[0-9]+}}], v{{\[}}[[TWO]]:{{[0-9]+}}], 1.0
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
bb:
   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> <half 1.0, half 1.0, half 1.0, half 1.0>, <4 x half> <half 2.0, half 2.0, half 2.0, half 2.0>, <16 x float> <float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0>, i32 0, i32 0, i32 0)
   store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
   ret void
 }
 
 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_imm_splat:
 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
-; GCN: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 0
+; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
+; NOLIT-SRCC: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9:]+}}]
+; LIT-SRCC: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 0
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32
 ; GCN-DAG: v_accvgpr_read_b32