arsenm created this revision.
arsenm added reviewers: tstellarAMD, echristo.
arsenm added a subscriber: cfe-commits.
Cleanup setup of subtarget features.
http://reviews.llvm.org/D17516
Files:
include/clang/Basic/BuiltinsAMDGPU.def
lib/Basic/Targets.cpp
test/CodeGenOpenCL/builtins-amdgcn-error.cl
test/CodeGenOpenCL/builtins-amdgcn-vi.cl
test/CodeGenOpenCL/builtins-amdgcn.cl
test/SemaOpenCL/builtins-amdgcn.cl
Index: test/SemaOpenCL/builtins-amdgcn.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn.cl
+++ /dev/null
@@ -1,6 +0,0 @@
-// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -fsyntax-only -verify %s
-
-void test_s_sleep(int x)
-{
- __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
-}
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -178,13 +178,6 @@
*out = __builtin_amdgcn_s_memtime();
}
-// CHECK-LABEL: @test_s_memrealtime
-// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
-void test_s_memrealtime(global ulong* out)
-{
- *out = __builtin_amdgcn_s_memrealtime();
-}
-
// CHECK-LABEL: @test_s_sleep
// CHECK: call void @llvm.amdgcn.s.sleep(i32 1)
// CHECK: call void @llvm.amdgcn.s.sleep(i32 15)
Index: test/CodeGenOpenCL/builtins-amdgcn-vi.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -0,0 +1,12 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
+
+typedef unsigned long ulong;
+
+
+// CHECK-LABEL: @test_s_memrealtime
+// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
+void test_s_memrealtime(global ulong* out)
+{
+ *out = __builtin_amdgcn_s_memrealtime();
+}
Index: test/CodeGenOpenCL/builtins-amdgcn-error.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCL/builtins-amdgcn-error.cl
@@ -0,0 +1,18 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -target-cpu tahiti -verify -S -o - %s
+
+// FIXME: We only get one error if the functions are the other order in the
+// file.
+
+typedef unsigned long ulong;
+
+ulong test_s_memrealtime()
+{
+ return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
+}
+
+void test_s_sleep(int x)
+{
+ __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
+}
+
Index: lib/Basic/Targets.cpp
===================================================================
--- lib/Basic/Targets.cpp
+++ lib/Basic/Targets.cpp
@@ -1801,23 +1801,32 @@
bool hasFMAF:1;
bool hasLDEXPF:1;
+ bool hasSMemRealTime : 1;
+ bool Has16BitInsts : 1;
+
+ static bool isAMDGCN(const llvm::Triple &TT) {
+ return TT.getArch() == llvm::Triple::amdgcn;
+ }
+
public:
AMDGPUTargetInfo(const llvm::Triple &Triple)
- : TargetInfo(Triple) {
+ : TargetInfo(Triple) ,
+ GPU(isAMDGCN(Triple) ? GK_SOUTHERN_ISLANDS : GK_R600),
+ hasFP64(false),
+ hasFMAF(false),
+ hasLDEXPF(false),
+ hasSMemRealTime(false),
+ Has16BitInsts(false) {
if (Triple.getArch() == llvm::Triple::amdgcn) {
DataLayoutString = DataLayoutStringSI;
- GPU = GK_SOUTHERN_ISLANDS;
hasFP64 = true;
hasFMAF = true;
hasLDEXPF = true;
} else {
DataLayoutString = DataLayoutStringR600;
- GPU = GK_R600;
- hasFP64 = false;
- hasFMAF = false;
- hasLDEXPF = false;
}
+
AddrSpaceMap = &AMDGPUAddrSpaceMap;
UseAddrSpaceMapMangling = true;
}
@@ -1858,6 +1867,10 @@
return false;
}
+ bool initFeatureMap(llvm::StringMap<bool> &Features,
+ DiagnosticsEngine &Diags, StringRef CPU,
+ const std::vector<std::string> &FeatureVec) const override;
+
ArrayRef<Builtin::Info> getTargetBuiltins() const override {
return llvm::makeArrayRef(BuiltinInfo,
clang::AMDGPU::LastTSBuiltin - Builtin::FirstTSBuiltin);
@@ -1929,39 +1942,45 @@
.Case("carrizo", GK_VOLCANIC_ISLANDS)
.Default(GK_NONE);
- if (GPU == GK_NONE) {
+ if (GPU == GK_NONE)
return false;
- }
// Set the correct data layout
- switch (GPU) {
- case GK_NONE:
- case GK_R600:
- case GK_R700:
- case GK_EVERGREEN:
- case GK_NORTHERN_ISLANDS:
- DataLayoutString = DataLayoutStringR600;
- hasFP64 = false;
- hasFMAF = false;
- hasLDEXPF = false;
- break;
- case GK_R600_DOUBLE_OPS:
- case GK_R700_DOUBLE_OPS:
- case GK_EVERGREEN_DOUBLE_OPS:
- case GK_CAYMAN:
- DataLayoutString = DataLayoutStringR600DoubleOps;
- hasFP64 = true;
- hasFMAF = true;
- hasLDEXPF = false;
- break;
- case GK_SOUTHERN_ISLANDS:
- case GK_SEA_ISLANDS:
- case GK_VOLCANIC_ISLANDS:
- DataLayoutString = DataLayoutStringSI;
- hasFP64 = true;
- hasFMAF = true;
- hasLDEXPF = true;
- break;
+ //if (isAMDGCN) {
+ if (getTriple().getArch() == llvm::Triple::amdgcn) {
+ switch (GPU) {
+ case GK_SOUTHERN_ISLANDS:
+ case GK_SEA_ISLANDS:
+ break;
+
+ case GK_VOLCANIC_ISLANDS:
+ hasSMemRealTime = true;
+ Has16BitInsts = true;
+ break;
+
+ default:
+ llvm_unreachable("unhandled subtarget");
+ }
+ } else {
+ switch (GPU) {
+ case GK_NONE:
+ case GK_R600:
+ case GK_R700:
+ case GK_EVERGREEN:
+ case GK_NORTHERN_ISLANDS:
+ break;
+ case GK_R600_DOUBLE_OPS:
+ case GK_R700_DOUBLE_OPS:
+ case GK_EVERGREEN_DOUBLE_OPS:
+ case GK_CAYMAN:
+ // FIXME: This should be its own target machine.
+ DataLayoutString = DataLayoutStringR600DoubleOps;
+ hasFP64 = true;
+ hasFMAF = true;
+ break;
+ default:
+ llvm_unreachable("unhandled subtarget");
+ }
}
return true;
@@ -1971,6 +1990,8 @@
const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
#define BUILTIN(ID, TYPE, ATTRS) \
{ #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr },
+#define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) \
+ { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, FEATURE },
#include "clang/Basic/BuiltinsAMDGPU.def"
};
const char * const AMDGPUTargetInfo::GCCRegNames[] = {
@@ -2030,6 +2051,20 @@
return llvm::makeArrayRef(GCCRegNames);
}
+bool AMDGPUTargetInfo::initFeatureMap(
+ llvm::StringMap<bool> &Features,
+ DiagnosticsEngine &Diags, StringRef CPU,
+ const std::vector<std::string> &FeatureVec) const {
+
+ if (Has16BitInsts)
+ Features["16-bit-insts"] = true;
+
+ if (hasSMemRealTime)
+ Features["s-memrealtime"] = true;
+
+ return true;
+}
+
// Namespace for x86 abstract base class
const Builtin::Info BuiltinInfo[] = {
#define BUILTIN(ID, TYPE, ATTRS) \
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -14,6 +14,10 @@
// The format of this database matches clang/Basic/Builtins.def.
+#if defined(BUILTIN) && !defined(TARGET_BUILTIN)
+# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
+#endif
+
BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
@@ -46,7 +50,8 @@
//===----------------------------------------------------------------------===//
// VI+ only builtins.
//===----------------------------------------------------------------------===//
-BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n")
+
+TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
//===----------------------------------------------------------------------===//
// Legacy names with amdgpu prefix
@@ -58,3 +63,4 @@
BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc")
#undef BUILTIN
+#undef TARGET_BUILTIN
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
