[AMDGPU] Restore the s_memtime instruction in gfx1030

[foad]

Summary

gfx1030 added a new way to implement readcyclecounter using the
SHADER_CYCLES hardware register, but the s_memtime instruction still
exists, so the MC layer should still accept it and the
llvm.amdgcn.s.memtime intrinsic should still work.

[rampitec]

It also needs clang change in the clang/lib/Basic/Targets/AMDGPU.cpp and clang tests.

[foad]

Add clang changes.

[rampitec]

LGTM, thanks!

[hctim]

Hello, it looks like this change broke the MSan and ASan buildbots.

http://lab.llvm.org:8011/#/builders/5/builds/5327

FAIL: LLVM :: CodeGen/AMDGPU/GlobalISel/readcyclecounter.ll (39307 of 74613)
******************** TEST 'LLVM :: CodeGen/AMDGPU/GlobalISel/readcyclecounter.ll' FAILED ********************
Script:
--
: 'RUN: at line 2';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/llc -global-isel -march=amdgcn -mcpu=tonga -verify-machineinstrs < /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/CodeGen/AMDGPU/GlobalISel/../readcyclecounter.ll | /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/FileCheck -enable-var-scope -check-prefix=MEMTIME -check-prefix=SIVI -check-prefix=GCN /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/CodeGen/AMDGPU/GlobalISel/../readcyclecounter.ll
: 'RUN: at line 3';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/llc -global-isel -march=amdgcn -mcpu=gfx1010 -verify-machineinstrs < /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/CodeGen/AMDGPU/GlobalISel/../readcyclecounter.ll | /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/FileCheck -enable-var-scope -check-prefix=MEMTIME -check-prefix=GCN /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/CodeGen/AMDGPU/GlobalISel/../readcyclecounter.ll
--
Exit Code: 2
Command Output (stderr):
--
==75699==WARNING: MemorySanitizer: use-of-uninitialized-value
    #0 0x3518e32 in bool llvm::InstructionSelector::executeMatchTable<llvm::AMDGPUInstructionSelector const, llvm::PredicateBitsetImpl<62ul>, llvm::Optional<llvm::SmallVector<std::__1::function<void (llvm::MachineInstrBuilder&)>, 4u> > (llvm::AMDGPUInstructionSelector::*)(llvm::MachineOperand&) const, void (llvm::AMDGPUInstructionSelector::*)(llvm::MachineInstrBuilder&, llvm::MachineInstr const&, int) const>(llvm::AMDGPUInstructionSelector const&, llvm::SmallVector<llvm::MachineInstrBuilder, 4u>&, llvm::InstructionSelector::MatcherState&, llvm::InstructionSelector::ISelInfoTy<llvm::PredicateBitsetImpl<62ul>, llvm::Optional<llvm::SmallVector<std::__1::function<void (llvm::MachineInstrBuilder&)>, 4u> > (llvm::AMDGPUInstructionSelector::*)(llvm::MachineOperand&) const, void (llvm::AMDGPUInstructionSelector::*)(llvm::MachineInstrBuilder&, llvm::MachineInstr const&, int) const> const&, long const*, llvm::TargetInstrInfo const&, llvm::MachineRegisterInfo&, llvm::TargetRegisterInfo const&, llvm::RegisterBankInfo const&, llvm::PredicateBitsetImpl<62ul> const&, llvm::CodeGenCoverage&) const /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/include/llvm/CodeGen/GlobalISel/InstructionSelectorImpl.h:149:11
    #1 0x34fe428 in llvm::AMDGPUInstructionSelector::selectImpl(llvm::MachineInstr&, llvm::CodeGenCoverage&) const /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/lib/Target/AMDGPU/AMDGPUGenGlobalISel.inc:1218:7
    #2 0x3557c75 in llvm::AMDGPUInstructionSelector::select(llvm::MachineInstr&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
    #3 0xb45543c in llvm::InstructionSelect::runOnMachineFunction(llvm::MachineFunction&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/CodeGen/GlobalISel/InstructionSelect.cpp:186:18
    #4 0x7327a1f in llvm::MachineFunctionPass::runOnFunction(llvm::Function&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/CodeGen/MachineFunctionPass.cpp:72:13
    #5 0x826124c in llvm::FPPassManager::runOnFunction(llvm::Function&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1439:27
    #6 0x6482dab in RunPassOnSCC /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/Analysis/CallGraphSCCPass.cpp:178:25
    #7 0x6482dab in RunAllPassesOnSCC /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/Analysis/CallGraphSCCPass.cpp:476:9
    #8 0x6482dab in (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/Analysis/CallGraphSCCPass.cpp:541:18
    #9 0x8263b7a in runOnModule /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1554:27
    #10 0x8263b7a in llvm::legacy::PassManagerImpl::run(llvm::Module&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/IR/LegacyPassManager.cpp:542:44
    #11 0x8280619 in llvm::legacy::PassManager::run(llvm::Module&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/IR/LegacyPassManager.cpp:1681:14
    #12 0x28f8ca2 in compileModule(char**, llvm::LLVMContext&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llc/llc.cpp:687:8
    #13 0x28f2b80 in main /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llc/llc.cpp:388:22
    #14 0x7f1559fb809a in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x2409a)
    #15 0x2872809 in _start (/b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/llc+0x2872809)
SUMMARY: MemorySanitizer: use-of-uninitialized-value /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/include/llvm/CodeGen/GlobalISel/InstructionSelectorImpl.h:149:11 in bool llvm::InstructionSelector::executeMatchTable<llvm::AMDGPUInstructionSelector const, llvm::PredicateBitsetImpl<62ul>, llvm::Optional<llvm::SmallVector<std::__1::function<void (llvm::MachineInstrBuilder&)>, 4u> > (llvm::AMDGPUInstructionSelector::*)(llvm::MachineOperand&) const, void (llvm::AMDGPUInstructionSelector::*)(llvm::MachineInstrBuilder&, llvm::MachineInstr const&, int) const>(llvm::AMDGPUInstructionSelector const&, llvm::SmallVector<llvm::MachineInstrBuilder, 4u>&, llvm::InstructionSelector::MatcherState&, llvm::InstructionSelector::ISelInfoTy<llvm::PredicateBitsetImpl<62ul>, llvm::Optional<llvm::SmallVector<std::__1::function<void (llvm::MachineInstrBuilder&)>, 4u> > (llvm::AMDGPUInstructionSelector::*)(llvm::MachineOperand&) const, void (llvm::AMDGPUInstructionSelector::*)(llvm::MachineInstrBuilder&, llvm::MachineInstr c
Exiting
FileCheck error: '<stdin>' is empty.
FileCheck command line:  /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/FileCheck -enable-var-scope -check-prefix=MEMTIME -check-prefix=SIVI -check-prefix=GCN /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/CodeGen/AMDGPU/GlobalISel/../readcyclecounter.ll

clang/lib/Basic/Targets/AMDGPU.cpp

case GK_GFX1030:
       Features["flat-address-space"] = true;
       Features["16-bit-insts"] = true;
       Features["dpp"] = true;
       Features["gfx8-insts"] = true;
       Features["gfx9-insts"] = true;
       Features["gfx10-insts"] = true;
       Features["gfx10-3-insts"] = true;
       Features["s-memrealtime"] = true;
+      Features["s-memtime-inst"] = true;
       break;
     case GK_GFX1012:
     case GK_GFX1011:
       Features["dot1-insts"] = true;
       Features["dot2-insts"] = true;
       Features["dot5-insts"] = true;
       Features["dot6-insts"] = true;
       LLVM_FALLTHROUGH;

clang/test/CodeGenOpenCL/amdgpu-features.cl

 // GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
 // GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime,+s-memtime-inst"
 // GFX909: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
 // GFX90A: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+mai-insts,+s-memrealtime,+s-memtime-inst"
 // GFX90C: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
 // GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
 // GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
 // GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
-// GFX1030: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1031: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1032: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1033: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
+// GFX1030: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
+// GFX1031: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
+// GFX1032: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
+// GFX1033: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-3-insts,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst"
 
 kernel void test() {}

clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl

-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1030 -verify -S -o - %s
-
-void test_gfx1030_s_memtime()
-{
-  __builtin_amdgcn_s_memtime(); // expected-error {{'__builtin_amdgcn_s_memtime' needs target feature s-memtime-inst}}
-}

llvm/lib/Target/AMDGPU/AMDGPU.td

 >;
 
 def FeatureSMemTimeInst : SubtargetFeature<"s-memtime-inst",
   "HasSMemTimeInst",
   "true",
   "Has s_memtime instruction"
 >;
 
+def FeatureShaderCyclesRegister : SubtargetFeature<"shader-cycles-register",
+  "HasShaderCyclesRegister",
+  "true",
+  "Has SHADER_CYCLES hardware register"
+>;
+
 def FeatureMadMacF32Insts : SubtargetFeature<"mad-mac-f32-insts",
   "HasMadMacF32Insts",
   "true",
   "Has v_mad_f32/v_mac_f32/v_madak_f32/v_madmk_f32 instructions"
 >;
 
 def FeatureDsSrc2Insts : SubtargetFeature<"ds-src2-insts",
   "HasDsSrc2Insts",
def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10",
    FeatureApertureRegs, FeatureGFX9Insts, FeatureGFX10Insts, FeatureVOP3P,
    FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp,
    FeatureSDWA, FeatureSDWAOmod, FeatureSDWAScalar, FeatureSDWASdst,
    FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts,
    FeatureAddNoCarryInsts, FeatureFmaMixInsts, FeatureGFX8Insts,
    FeatureNoSdstCMPX, FeatureVscnt, FeatureRegisterBanking,
    FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,
    FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
-   FeatureGFX10A16, FeatureFastDenormalF32, FeatureG16,
+   FeatureGFX10A16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureG16,
    FeatureUnalignedBufferAccess, FeatureUnalignedDSAccess
   ]
 >;
 
 class FeatureSet<list<SubtargetFeature> Features_> {
   list<SubtargetFeature> Features = Features_;
 }
 
!listconcat(FeatureGroup.GFX10_1_Bugs,
      FeatureLDSBankCount32,
      FeatureDLInsts,
      FeatureNSAEncoding,
      FeatureWavefrontSize32,
      FeatureScalarStores,
      FeatureScalarAtomics,
      FeatureScalarFlatScratchInsts,
      FeatureGetWaveIdInst,
-     FeatureSMemTimeInst,
      FeatureMadMacF32Insts,
      FeatureDsSrc2Insts,
      FeatureLdsMisalignedBug,
      FeatureSupportsXNACK])>;
 
 def FeatureISAVersion10_1_1 : FeatureSet<
   !listconcat(FeatureGroup.GFX10_1_Bugs,
     [FeatureGFX10,
      FeatureLDSBankCount32,
      FeatureDLInsts,
      FeatureDot1Insts,
      FeatureDot2Insts,
      FeatureDot5Insts,
      FeatureDot6Insts,
      FeatureNSAEncoding,
      FeatureWavefrontSize32,
      FeatureScalarStores,
      FeatureScalarAtomics,
      FeatureScalarFlatScratchInsts,
      FeatureGetWaveIdInst,
-     FeatureSMemTimeInst,
      FeatureMadMacF32Insts,
      FeatureDsSrc2Insts,
      FeatureLdsMisalignedBug,
      FeatureSupportsXNACK])>;
 
 def FeatureISAVersion10_1_2 : FeatureSet<
   !listconcat(FeatureGroup.GFX10_1_Bugs,
     [FeatureGFX10,
      FeatureLDSBankCount32,
      FeatureDLInsts,
      FeatureDot1Insts,
      FeatureDot2Insts,
      FeatureDot5Insts,
      FeatureDot6Insts,
      FeatureNSAEncoding,
      FeatureWavefrontSize32,
      FeatureScalarStores,
      FeatureScalarAtomics,
      FeatureScalarFlatScratchInsts,
      FeatureGetWaveIdInst,
-     FeatureSMemTimeInst,
      FeatureMadMacF32Insts,
      FeatureDsSrc2Insts,
      FeatureLdsMisalignedBug,
      FeatureSupportsXNACK])>;
 
 def FeatureISAVersion10_3_0 : FeatureSet<
   [FeatureGFX10,
    FeatureGFX10_BEncoding,
    FeatureGFX10_3Insts,
    FeatureLDSBankCount32,
    FeatureDLInsts,
    FeatureDot1Insts,
    FeatureDot2Insts,
    FeatureDot5Insts,
    FeatureDot6Insts,
    FeatureNSAEncoding,
-   FeatureWavefrontSize32]>;
+   FeatureWavefrontSize32,
+   FeatureShaderCyclesRegister]>;
 
 //===----------------------------------------------------------------------===//
 
 def AMDGPUInstrInfo : InstrInfo {
   let guessInstructionProperties = 1;
   let noNamedPositionallyEncodedOperands = 1;
 }
 
def HasMAIInsts : Predicate<"Subtarget->hasMAIInsts()">,
   AssemblerPredicate<(all_of FeatureMAIInsts)>;
 
 def HasSMemRealTime : Predicate<"Subtarget->hasSMemRealTime()">,
   AssemblerPredicate<(all_of FeatureSMemRealTime)>;
 
 def HasSMemTimeInst : Predicate<"Subtarget->hasSMemTimeInst()">,
   AssemblerPredicate<(all_of FeatureSMemTimeInst)>;
 
-def HasNoSMemTimeInst : Predicate<"!Subtarget->hasSMemTimeInst()">;
+def HasShaderCyclesRegister : Predicate<"Subtarget->hasShaderCyclesRegister()">,
+  AssemblerPredicate<(all_of FeatureShaderCyclesRegister)>;
 
 def HasPkFmacF16Inst : Predicate<"Subtarget->hasPkFmacF16Inst()">,
   AssemblerPredicate<(all_of FeaturePkFmacF16Inst)>;
 
 def HasMadMacF32Insts : Predicate<"Subtarget->hasMadMacF32Insts()">,
   AssemblerPredicate<(all_of FeatureMadMacF32Insts)>;
 
 def HasFmaLegacy32 : Predicate<"Subtarget->hasGFX10_3Insts()">,

llvm/lib/Target/AMDGPU/GCNSubtarget.h

protected:
   // This should not be used directly. 'TargetID' tracks the dynamic settings
   // for SRAMECC.
   bool EnableSRAMECC;
 
   bool HasNoSdstCMPX;
   bool HasVscnt;
   bool HasGetWaveIdInst;
   bool HasSMemTimeInst;
+  bool HasShaderCyclesRegister;
   bool HasRegisterBanking;
   bool HasVOP3Literal;
   bool HasNoDataDepHazard;
   bool FlatAddressSpace;
   bool FlatInstOffsets;
   bool FlatGlobalInsts;
   bool FlatScratchInsts;
   bool ScalarFlatScratchInsts;
public:
   bool hasGetWaveIdInst() const {
     return HasGetWaveIdInst;
   }
 
   bool hasSMemTimeInst() const {
     return HasSMemTimeInst;
   }
 
+  bool hasShaderCyclesRegister() const {
+    return HasShaderCyclesRegister;
+  }
+
   bool hasRegisterBanking() const {
     return HasRegisterBanking;
   }
 
   bool hasVOP3Literal() const {
     return HasVOP3Literal;
   }
 

llvm/lib/Target/AMDGPU/SMInstructions.td

 
 let OtherPredicates = [HasSMemTimeInst] in {
 def : GCNPat <
   (i64 (readcyclecounter)),
   (S_MEMTIME)
 >;
 } // let OtherPredicates = [HasSMemTimeInst]
 
-let OtherPredicates = [HasNoSMemTimeInst] in {
+let OtherPredicates = [HasShaderCyclesRegister] in {
 def : GCNPat <
   (i64 (readcyclecounter)),
   (REG_SEQUENCE SReg_64,
     (S_GETREG_B32 getHwRegImm<HWREG.SHADER_CYCLES, 0, -12>.ret), sub0,
-    (S_MOV_B32 (i32 0)), sub1)
->;
-} // let OtherPredicates = [HasNoSMemTimeInst]
+    (S_MOV_B32 (i32 0)), sub1)> {
+  // Prefer this to s_memtime because it has lower and more predictable latency.
+  let AddedComplexity = 1;
+}
+} // let OtherPredicates = [HasShaderCyclesRegister]
 
 //===----------------------------------------------------------------------===//
 // GFX10.
 //===----------------------------------------------------------------------===//
 
 class SMEM_Real_gfx10<bits<8> op, SM_Pseudo ps> :
     SM_Real<ps>, SIMCInstr<ps.PseudoInstr, SIEncodingFamily.GFX10>, Enc64 {
   bit glc;

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll

 ; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck --check-prefixes=SIVI,GCN %s
 ; RUN: llc -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck --check-prefixes=SIVI,GCN %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
-; RUN: not --crash llc -march=amdgcn -mcpu=gfx1030 -mattr=-flat-for-global -verify-machineinstrs < %s 2>&1 | FileCheck -check-prefix=GFX1030-ERR %s
+; RUN: llc -march=amdgcn -mcpu=gfx1030 -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 declare i64 @llvm.amdgcn.s.memtime() #0
 
 ; GCN-LABEL: {{^}}test_s_memtime:
 ; GCN-DAG: s_memtime s{{\[[0-9]+:[0-9]+\]}}
 ; GCN-DAG: s_load_dwordx2
 ; GCN: lgkmcnt
 ; GCN: {{buffer|global}}_store_dwordx2
 ; SIVI-NOT: lgkmcnt
 ; GCN: s_memtime s{{\[[0-9]+:[0-9]+\]}}
 ; GCN: {{buffer|global}}_store_dwordx2
-; GFX1030-ERR: ERROR
 define amdgpu_kernel void @test_s_memtime(i64 addrspace(1)* %out) #0 {
   %cycle0 = call i64 @llvm.amdgcn.s.memtime()
   store volatile i64 %cycle0, i64 addrspace(1)* %out
 
   %cycle1 = call i64 @llvm.amdgcn.s.memtime()
   store volatile i64 %cycle1, i64 addrspace(1)* %out
   ret void
 }
 
 attributes #0 = { nounwind }

llvm/test/MC/AMDGPU/gfx1030_err.s

 // GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 
 v_dot8c_i32_i4 v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
 // GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 
 s_get_waveid_in_workgroup s0
 // GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 
-s_memtime s[0:1]
-// GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
-
 s_getreg_b32 s2, hwreg(HW_REG_XNACK_MASK)
 // GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: specified hardware register is not supported on this GPU
 
 v_mac_f32 v0, v1, v2
 // GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 
 v_mad_f32 v0, v1, v2, v3
 // GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU