arsenm updated this revision to Diff 51636.
arsenm added a comment.
Try to move more code into initFeatureMap.
I'm not sure how the booleans for features in the class are for now. X86 seems
to have them, but it seems they are only used with user specified features? The
only ones that matter right now don't depend on the subtarget, so I've left
that in the constructor
http://reviews.llvm.org/D17516
Files:
include/clang/Basic/BuiltinsAMDGPU.def
lib/Basic/Targets.cpp
lib/Driver/Tools.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/Driver/Tools.cpp
===================================================================
--- lib/Driver/Tools.cpp
+++ lib/Driver/Tools.cpp
@@ -1839,8 +1839,9 @@
case llvm::Triple::systemz:
return getSystemZTargetCPU(Args);
- case llvm::Triple::r600:
case llvm::Triple::amdgcn:
+ return getAMDGCNTargetGPU(Args);
+ case llvm::Triple::r600:
return getR600TargetGPU(Args);
case llvm::Triple::wasm32:
Index: lib/Basic/Targets.cpp
===================================================================
--- lib/Basic/Targets.cpp
+++ lib/Basic/Targets.cpp
@@ -1752,7 +1752,7 @@
"-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
"-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64";
-class AMDGPUTargetInfo : public TargetInfo {
+class AMDGPUTargetInfo final : public TargetInfo {
static const Builtin::Info BuiltinInfo[];
static const char * const GCCRegNames[];
@@ -1776,23 +1776,26 @@
bool hasFMAF:1;
bool hasLDEXPF:1;
+ static bool isAMDGCN(const llvm::Triple &TT) {
+ return TT.getArch() == llvm::Triple::amdgcn;
+ }
+
public:
AMDGPUTargetInfo(const llvm::Triple &Triple)
- : TargetInfo(Triple) {
-
- if (Triple.getArch() == llvm::Triple::amdgcn) {
- resetDataLayout(DataLayoutStringSI);
- GPU = GK_SOUTHERN_ISLANDS;
+ : TargetInfo(Triple) ,
+ GPU(isAMDGCN(Triple) ? GK_SOUTHERN_ISLANDS : GK_R600),
+ hasFP64(false),
+ hasFMAF(false),
+ hasLDEXPF(false) {
+ if (getTriple().getArch() == llvm::Triple::amdgcn) {
hasFP64 = true;
hasFMAF = true;
hasLDEXPF = true;
- } else {
- resetDataLayout(DataLayoutStringR600);
- GPU = GK_R600;
- hasFP64 = false;
- hasFMAF = false;
- hasLDEXPF = false;
}
+
+ resetDataLayout(getTriple().getArch() == llvm::Triple::amdgcn ?
+ DataLayoutStringSI : DataLayoutStringR600);
+
AddrSpaceMap = &AMDGPUAddrSpaceMap;
UseAddrSpaceMapMangling = true;
}
@@ -1833,6 +1836,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);
@@ -1871,8 +1878,8 @@
return TargetInfo::CharPtrBuiltinVaList;
}
- bool setCPU(const std::string &Name) override {
- GPU = llvm::StringSwitch<GPUKind>(Name)
+ static GPUKind parseR600Name(StringRef Name) {
+ return llvm::StringSwitch<GPUKind>(Name)
.Case("r600" , GK_R600)
.Case("rv610", GK_R600)
.Case("rv620", GK_R600)
@@ -1898,6 +1905,11 @@
.Case("caicos", GK_NORTHERN_ISLANDS)
.Case("cayman", GK_CAYMAN)
.Case("aruba", GK_CAYMAN)
+ .Default(GK_NONE);
+ }
+
+ static GPUKind parseAMDGCNName(StringRef Name) {
+ return llvm::StringSwitch<GPUKind>(Name)
.Case("tahiti", GK_SOUTHERN_ISLANDS)
.Case("pitcairn", GK_SOUTHERN_ISLANDS)
.Case("verde", GK_SOUTHERN_ISLANDS)
@@ -1914,49 +1926,23 @@
.Case("fiji", GK_VOLCANIC_ISLANDS)
.Case("stoney", GK_VOLCANIC_ISLANDS)
.Default(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:
- resetDataLayout(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:
- resetDataLayout(DataLayoutStringR600);
- hasFP64 = true;
- hasFMAF = true;
- hasLDEXPF = false;
- break;
- case GK_SOUTHERN_ISLANDS:
- case GK_SEA_ISLANDS:
- case GK_VOLCANIC_ISLANDS:
- resetDataLayout(DataLayoutStringSI);
- hasFP64 = true;
- hasFMAF = true;
- hasLDEXPF = true;
- break;
- }
+ bool setCPU(const std::string &Name) override {
+ if (getTriple().getArch() == llvm::Triple::amdgcn)
+ GPU = parseAMDGCNName(Name);
+ else
+ GPU = parseR600Name(Name);
- return true;
+ return GPU != GK_NONE;
}
};
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[] = {
@@ -2016,6 +2002,57 @@
return llvm::makeArrayRef(GCCRegNames);
}
+bool AMDGPUTargetInfo::initFeatureMap(
+ llvm::StringMap<bool> &Features,
+ DiagnosticsEngine &Diags, StringRef CPU,
+ const std::vector<std::string> &FeatureVec) const {
+
+ // XXX - What does the member GPU mean if device name string passed here?
+ if (getTriple().getArch() == llvm::Triple::amdgcn) {
+ if (CPU.empty())
+ CPU = "tahiti";
+
+ switch (parseAMDGCNName(CPU)) {
+ case GK_SOUTHERN_ISLANDS:
+ case GK_SEA_ISLANDS:
+ break;
+
+ case GK_VOLCANIC_ISLANDS:
+ Features["s-memrealtime"] = true;
+ Features["16-bit-insts"] = true;
+ break;
+
+ case GK_NONE:
+ return false;
+ default:
+ llvm_unreachable("unhandled subtarget");
+ }
+ } else {
+ if (CPU.empty())
+ CPU = "r600";
+
+ switch (parseR600Name(CPU)) {
+ 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:
+ Features["fp64"] = true;
+ break;
+ case GK_NONE:
+ return false;
+ default:
+ llvm_unreachable("unhandled subtarget");
+ }
+ }
+
+ return TargetInfo::initFeatureMap(Features, Diags, CPU, FeatureVec);
+}
+
// 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
