Author: Mitch Phillips Date: 2021-03-05T18:24:59-08:00 New Revision: e58d68fcd06ddc7743e0419c0b364df3d44121b6
URL: https://github.com/llvm/llvm-project/commit/e58d68fcd06ddc7743e0419c0b364df3d44121b6 DIFF: https://github.com/llvm/llvm-project/commit/e58d68fcd06ddc7743e0419c0b364df3d44121b6.diff LOG: Revert "[AMDGPU] Restore the s_memtime instruction in gfx1030" Broke the ASan/MSan buildbots. See more comments in the original patch, https://reviews.llvm.org/D97928. Build failure at http://lab.llvm.org:8011/#/builders/5/builds/5327 This reverts commit fc28f600e558c1344618bda149a068d6162b6f0b. Added: clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl Modified: clang/lib/Basic/Targets/AMDGPU.cpp clang/test/CodeGenOpenCL/amdgpu-features.cl llvm/lib/Target/AMDGPU/AMDGPU.td llvm/lib/Target/AMDGPU/GCNSubtarget.h llvm/lib/Target/AMDGPU/SMInstructions.td llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll llvm/test/MC/AMDGPU/gfx1030_err.s Removed: ################################################################################ diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index a84422e412ff..0f1211d6c409 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -192,7 +192,6 @@ bool AMDGPUTargetInfo::initFeatureMap( 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: diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl index 930c53705d84..f387c93bd6e0 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-features.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl @@ -58,9 +58,9 @@ // 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,+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" +// 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" kernel void test() {} diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl new file mode 100644 index 000000000000..34149148a5f3 --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl @@ -0,0 +1,7 @@ +// 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}} +} diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index e2e9c42dc985..452d3bb6c9a9 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -563,12 +563,6 @@ def FeatureSMemTimeInst : SubtargetFeature<"s-memtime-inst", "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", @@ -783,7 +777,7 @@ def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10", FeatureNoSdstCMPX, FeatureVscnt, FeatureRegisterBanking, FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts, FeatureNoDataDepHazard, FeaturePkFmacF16Inst, - FeatureGFX10A16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureG16, + FeatureGFX10A16, FeatureFastDenormalF32, FeatureG16, FeatureUnalignedBufferAccess, FeatureUnalignedDSAccess ] >; @@ -994,6 +988,7 @@ def FeatureISAVersion10_1_0 : FeatureSet< FeatureScalarAtomics, FeatureScalarFlatScratchInsts, FeatureGetWaveIdInst, + FeatureSMemTimeInst, FeatureMadMacF32Insts, FeatureDsSrc2Insts, FeatureLdsMisalignedBug, @@ -1014,6 +1009,7 @@ def FeatureISAVersion10_1_1 : FeatureSet< FeatureScalarAtomics, FeatureScalarFlatScratchInsts, FeatureGetWaveIdInst, + FeatureSMemTimeInst, FeatureMadMacF32Insts, FeatureDsSrc2Insts, FeatureLdsMisalignedBug, @@ -1034,6 +1030,7 @@ def FeatureISAVersion10_1_2 : FeatureSet< FeatureScalarAtomics, FeatureScalarFlatScratchInsts, FeatureGetWaveIdInst, + FeatureSMemTimeInst, FeatureMadMacF32Insts, FeatureDsSrc2Insts, FeatureLdsMisalignedBug, @@ -1050,8 +1047,7 @@ def FeatureISAVersion10_3_0 : FeatureSet< FeatureDot5Insts, FeatureDot6Insts, FeatureNSAEncoding, - FeatureWavefrontSize32, - FeatureShaderCyclesRegister]>; + FeatureWavefrontSize32]>; //===----------------------------------------------------------------------===// @@ -1381,8 +1377,7 @@ def HasSMemRealTime : Predicate<"Subtarget->hasSMemRealTime()">, def HasSMemTimeInst : Predicate<"Subtarget->hasSMemTimeInst()">, AssemblerPredicate<(all_of FeatureSMemTimeInst)>; -def HasShaderCyclesRegister : Predicate<"Subtarget->hasShaderCyclesRegister()">, - AssemblerPredicate<(all_of FeatureShaderCyclesRegister)>; +def HasNoSMemTimeInst : Predicate<"!Subtarget->hasSMemTimeInst()">; def HasPkFmacF16Inst : Predicate<"Subtarget->hasPkFmacF16Inst()">, AssemblerPredicate<(all_of FeaturePkFmacF16Inst)>; diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h index fc45f7ee11dc..c2d3491c23f8 100644 --- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h +++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h @@ -163,7 +163,6 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo, bool HasVscnt; bool HasGetWaveIdInst; bool HasSMemTimeInst; - bool HasShaderCyclesRegister; bool HasRegisterBanking; bool HasVOP3Literal; bool HasNoDataDepHazard; @@ -715,10 +714,6 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo, return HasSMemTimeInst; } - bool hasShaderCyclesRegister() const { - return HasShaderCyclesRegister; - } - bool hasRegisterBanking() const { return HasRegisterBanking; } diff --git a/llvm/lib/Target/AMDGPU/SMInstructions.td b/llvm/lib/Target/AMDGPU/SMInstructions.td index d2b5652a0db8..19afd72b3211 100644 --- a/llvm/lib/Target/AMDGPU/SMInstructions.td +++ b/llvm/lib/Target/AMDGPU/SMInstructions.td @@ -866,16 +866,14 @@ def : GCNPat < >; } // let OtherPredicates = [HasSMemTimeInst] -let OtherPredicates = [HasShaderCyclesRegister] in { +let OtherPredicates = [HasNoSMemTimeInst] 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)> { - // Prefer this to s_memtime because it has lower and more predictable latency. - let AddedComplexity = 1; -} -} // let OtherPredicates = [HasShaderCyclesRegister] + (S_MOV_B32 (i32 0)), sub1) +>; +} // let OtherPredicates = [HasNoSMemTimeInst] //===----------------------------------------------------------------------===// // GFX10. diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll index cd2eada5b82e..c9f0591a5dcc 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll @@ -1,7 +1,7 @@ ; 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: llc -march=amdgcn -mcpu=gfx1030 -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 declare i64 @llvm.amdgcn.s.memtime() #0 @@ -13,6 +13,7 @@ declare i64 @llvm.amdgcn.s.memtime() #0 ; 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 diff --git a/llvm/test/MC/AMDGPU/gfx1030_err.s b/llvm/test/MC/AMDGPU/gfx1030_err.s index 756983ef49b5..dbee18bd2d91 100644 --- a/llvm/test/MC/AMDGPU/gfx1030_err.s +++ b/llvm/test/MC/AMDGPU/gfx1030_err.s @@ -21,6 +21,9 @@ v_dot8c_i32_i4 v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
