r272091 - AMDGPU: Verify subtarget specific builtins
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Tue Jun 7 18:56:43 PDT 2016
Author: arsenm
Date: Tue Jun 7 20:56:42 2016
New Revision: 272091
URL: http://llvm.org/viewvc/llvm-project?rev=272091&view=rev
Log:
AMDGPU: Verify subtarget specific builtins
Cleanup setup of subtarget features.
Added:
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
Removed:
cfe/trunk/test/SemaOpenCL/builtins-amdgcn.cl
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/Basic/Targets.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=272091&r1=272090&r2=272091&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Tue Jun 7 20:56:42 2016
@@ -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")
@@ -52,7 +56,8 @@ BUILTIN(__builtin_amdgcn_s_sleep, "vIi",
//===----------------------------------------------------------------------===//
// 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
@@ -64,3 +69,4 @@ BUILTIN(__builtin_amdgpu_ldexp, "ddi", "
BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc")
#undef BUILTIN
+#undef TARGET_BUILTIN
Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=272091&r1=272090&r2=272091&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Tue Jun 7 20:56:42 2016
@@ -1925,7 +1925,7 @@ static const char *const DataLayoutStrin
"-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[];
@@ -1949,22 +1949,26 @@ class AMDGPUTargetInfo : public TargetIn
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, const TargetOptions &)
- : 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;
}
@@ -2005,6 +2009,10 @@ public:
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);
@@ -2027,8 +2035,8 @@ public:
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)
@@ -2054,6 +2062,11 @@ public:
.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)
@@ -2070,43 +2083,15 @@ public:
.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;
}
void setSupportedOpenCLOpts() override {
@@ -2138,6 +2123,8 @@ public:
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[] = {
@@ -2197,6 +2184,57 @@ ArrayRef<const char *> AMDGPUTargetInfo:
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) \
@@ -4950,7 +4988,7 @@ public:
} else if (Feature == "+dsp") {
DSP = 1;
} else if (Feature == "+fp-only-sp") {
- HW_FP_remove |= HW_FP_DP;
+ HW_FP_remove |= HW_FP_DP;
} else if (Feature == "+strict-align") {
Unaligned = 0;
} else if (Feature == "+fp16") {
Added: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl?rev=272091&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl Tue Jun 7 20:56:42 2016
@@ -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}}
+}
+
Added: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl?rev=272091&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Tue Jun 7 20:56:42 2016
@@ -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();
+}
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=272091&r1=272090&r2=272091&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Tue Jun 7 20:56:42 2016
@@ -220,13 +220,6 @@ void test_s_memtime(global ulong* out)
*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)
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl?rev=272091&r1=272090&r2=272091&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-r600.cl Tue Jun 7 20:56:42 2016
@@ -1,7 +1,5 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu rv670 -S -emit-llvm -o - %s | FileCheck %s
-
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s
// CHECK-LABEL: @test_rsq_f32
// CHECK: call float @llvm.r600.rsq.f32
@@ -10,12 +8,14 @@ void test_rsq_f32(global float* out, flo
*out = __builtin_amdgpu_rsqf(a);
}
-// CHECK-LABEL: @test_rsq_f64
-// CHECK: call double @llvm.r600.rsq.f64
+#if cl_khr_fp64
+// XCHECK-LABEL: @test_rsq_f64
+// XCHECK: call double @llvm.r600.rsq.f64
void test_rsq_f64(global double* out, double a)
{
*out = __builtin_amdgpu_rsq(a);
}
+#endif
// CHECK-LABEL: @test_legacy_ldexp_f32
// CHECK: call float @llvm.AMDGPU.ldexp.f32
@@ -24,9 +24,11 @@ void test_legacy_ldexp_f32(global float*
*out = __builtin_amdgpu_ldexpf(a, b);
}
-// CHECK-LABEL: @test_legacy_ldexp_f64
-// CHECK: call double @llvm.AMDGPU.ldexp.f64
+#if cl_khr_fp64
+// XCHECK-LABEL: @test_legacy_ldexp_f64
+// XCHECK: call double @llvm.AMDGPU.ldexp.f64
void test_legacy_ldexp_f64(global double* out, double a, int b)
{
*out = __builtin_amdgpu_ldexp(a, b);
}
+#endif
Removed: cfe/trunk/test/SemaOpenCL/builtins-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn.cl?rev=272090&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn.cl (removed)
@@ -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}}
-}
More information about the cfe-commits
mailing list