Skip to content

Commit e58d68f

Browse files
hctimhctim
committed
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 fc28f60.
1 parent 71e6e82 commit e58d68f

File tree

8 files changed

+26
-28
lines changed

8 files changed

+26
-28
lines changed

‎clang/lib/Basic/Targets/AMDGPU.cpp‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -192,7 +192,6 @@ bool AMDGPUTargetInfo::initFeatureMap(
192192
Features["gfx10-insts"] = true;
193193
Features["gfx10-3-insts"] = true;
194194
Features["s-memrealtime"] = true;
195-
Features["s-memtime-inst"] = true;
196195
break;
197196
case GK_GFX1012:
198197
case GK_GFX1011:

‎clang/test/CodeGenOpenCL/amdgpu-features.cl‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -58,9 +58,9 @@
5858
// 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"
5959
// 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"
6060
// 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"
61-
// 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"
62-
// 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"
63-
// 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"
64-
// 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"
61+
// 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"
62+
// 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"
63+
// 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"
64+
// 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"
6565

6666
kernel void test() {}

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

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1030 -verify -S -o - %s
3+
4+
void test_gfx1030_s_memtime()
5+
{
6+
__builtin_amdgcn_s_memtime(); // expected-error {{'__builtin_amdgcn_s_memtime' needs target feature s-memtime-inst}}
7+
}

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

Lines changed: 6 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -563,12 +563,6 @@ def FeatureSMemTimeInst : SubtargetFeature<"s-memtime-inst",
563563
"Has s_memtime instruction"
564564
>;
565565

566-
def FeatureShaderCyclesRegister : SubtargetFeature<"shader-cycles-register",
567-
"HasShaderCyclesRegister",
568-
"true",
569-
"Has SHADER_CYCLES hardware register"
570-
>;
571-
572566
def FeatureMadMacF32Insts : SubtargetFeature<"mad-mac-f32-insts",
573567
"HasMadMacF32Insts",
574568
"true",
@@ -783,7 +777,7 @@ def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10",
783777
FeatureNoSdstCMPX, FeatureVscnt, FeatureRegisterBanking,
784778
FeatureVOP3Literal, FeatureDPP8, FeatureExtendedImageInsts,
785779
FeatureNoDataDepHazard, FeaturePkFmacF16Inst,
786-
FeatureGFX10A16, FeatureSMemTimeInst, FeatureFastDenormalF32, FeatureG16,
780+
FeatureGFX10A16, FeatureFastDenormalF32, FeatureG16,
787781
FeatureUnalignedBufferAccess, FeatureUnalignedDSAccess
788782
]
789783
>;
@@ -994,6 +988,7 @@ def FeatureISAVersion10_1_0 : FeatureSet<
994988
FeatureScalarAtomics,
995989
FeatureScalarFlatScratchInsts,
996990
FeatureGetWaveIdInst,
991+
FeatureSMemTimeInst,
997992
FeatureMadMacF32Insts,
998993
FeatureDsSrc2Insts,
999994
FeatureLdsMisalignedBug,
@@ -1014,6 +1009,7 @@ def FeatureISAVersion10_1_1 : FeatureSet<
10141009
FeatureScalarAtomics,
10151010
FeatureScalarFlatScratchInsts,
10161011
FeatureGetWaveIdInst,
1012+
FeatureSMemTimeInst,
10171013
FeatureMadMacF32Insts,
10181014
FeatureDsSrc2Insts,
10191015
FeatureLdsMisalignedBug,
@@ -1034,6 +1030,7 @@ def FeatureISAVersion10_1_2 : FeatureSet<
10341030
FeatureScalarAtomics,
10351031
FeatureScalarFlatScratchInsts,
10361032
FeatureGetWaveIdInst,
1033+
FeatureSMemTimeInst,
10371034
FeatureMadMacF32Insts,
10381035
FeatureDsSrc2Insts,
10391036
FeatureLdsMisalignedBug,
@@ -1050,8 +1047,7 @@ def FeatureISAVersion10_3_0 : FeatureSet<
10501047
FeatureDot5Insts,
10511048
FeatureDot6Insts,
10521049
FeatureNSAEncoding,
1053-
FeatureWavefrontSize32,
1054-
FeatureShaderCyclesRegister]>;
1050+
FeatureWavefrontSize32]>;
10551051

10561052
//===----------------------------------------------------------------------===//
10571053

@@ -1381,8 +1377,7 @@ def HasSMemRealTime : Predicate<"Subtarget->hasSMemRealTime()">,
13811377
def HasSMemTimeInst : Predicate<"Subtarget->hasSMemTimeInst()">,
13821378
AssemblerPredicate<(all_of FeatureSMemTimeInst)>;
13831379

1384-
def HasShaderCyclesRegister : Predicate<"Subtarget->hasShaderCyclesRegister()">,
1385-
AssemblerPredicate<(all_of FeatureShaderCyclesRegister)>;
1380+
def HasNoSMemTimeInst : Predicate<"!Subtarget->hasSMemTimeInst()">;
13861381

13871382
def HasPkFmacF16Inst : Predicate<"Subtarget->hasPkFmacF16Inst()">,
13881383
AssemblerPredicate<(all_of FeaturePkFmacF16Inst)>;

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

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -163,7 +163,6 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
163163
bool HasVscnt;
164164
bool HasGetWaveIdInst;
165165
bool HasSMemTimeInst;
166-
bool HasShaderCyclesRegister;
167166
bool HasRegisterBanking;
168167
bool HasVOP3Literal;
169168
bool HasNoDataDepHazard;
@@ -715,10 +714,6 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
715714
return HasSMemTimeInst;
716715
}
717716

718-
bool hasShaderCyclesRegister() const {
719-
return HasShaderCyclesRegister;
720-
}
721-
722717
bool hasRegisterBanking() const {
723718
return HasRegisterBanking;
724719
}

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

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -866,16 +866,14 @@ def : GCNPat <
866866
>;
867867
} // let OtherPredicates = [HasSMemTimeInst]
868868

869-
let OtherPredicates = [HasShaderCyclesRegister] in {
869+
let OtherPredicates = [HasNoSMemTimeInst] in {
870870
def : GCNPat <
871871
(i64 (readcyclecounter)),
872872
(REG_SEQUENCE SReg_64,
873873
(S_GETREG_B32 getHwRegImm<HWREG.SHADER_CYCLES, 0, -12>.ret), sub0,
874-
(S_MOV_B32 (i32 0)), sub1)> {
875-
// Prefer this to s_memtime because it has lower and more predictable latency.
876-
let AddedComplexity = 1;
877-
}
878-
} // let OtherPredicates = [HasShaderCyclesRegister]
874+
(S_MOV_B32 (i32 0)), sub1)
875+
>;
876+
} // let OtherPredicates = [HasNoSMemTimeInst]
879877

880878
//===----------------------------------------------------------------------===//
881879
// GFX10.

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck --check-prefixes=SIVI,GCN %s
22
; RUN: llc -march=amdgcn -mcpu=tonga -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck --check-prefixes=SIVI,GCN %s
33
; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
4-
; RUN: llc -march=amdgcn -mcpu=gfx1030 -mattr=-flat-for-global -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
4+
; RUN: not --crash llc -march=amdgcn -mcpu=gfx1030 -mattr=-flat-for-global -verify-machineinstrs < %s 2>&1 | FileCheck -check-prefix=GFX1030-ERR %s
55

66
declare i64 @llvm.amdgcn.s.memtime() #0
77

@@ -13,6 +13,7 @@ declare i64 @llvm.amdgcn.s.memtime() #0
1313
; SIVI-NOT: lgkmcnt
1414
; GCN: s_memtime s{{\[[0-9]+:[0-9]+\]}}
1515
; GCN: {{buffer|global}}_store_dwordx2
16+
; GFX1030-ERR: ERROR
1617
define amdgpu_kernel void @test_s_memtime(i64 addrspace(1)* %out) #0 {
1718
%cycle0 = call i64 @llvm.amdgcn.s.memtime()
1819
store volatile i64 %cycle0, i64 addrspace(1)* %out

‎llvm/test/MC/AMDGPU/gfx1030_err.s‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,9 @@ v_dot8c_i32_i4 v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
2121
s_get_waveid_in_workgroup s0
2222
// GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
2323

24+
s_memtime s[0:1]
25+
// GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
26+
2427
s_getreg_b32 s2, hwreg(HW_REG_XNACK_MASK)
2528
// GFX10: :[[@LINE-1]]:{{[0-9]+}}: error: specified hardware register is not supported on this GPU
2629

0 commit comments

Comments
 (0)