[clang] e58d68f - Revert "[AMDGPU] Restore the s_memtime instruction in gfx1030"
Mitch Phillips via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 5 18:25:28 PST 2021
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
More information about the cfe-commits
mailing list