From 2e5f745694ecf1d67a2892dcbb5cce601d580bf6 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Wed, 8 Jun 2016 01:56:42 +0000 Subject: [PATCH] AMDGPU: Verify subtarget specific builtins Cleanup setup of subtarget features. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@272091 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/BuiltinsAMDGPU.def | 8 +- lib/Basic/Targets.cpp | 136 +++++++++++++------- test/CodeGenOpenCL/builtins-amdgcn-error.cl | 18 +++ test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 12 ++ test/CodeGenOpenCL/builtins-amdgcn.cl | 7 - test/CodeGenOpenCL/builtins-r600.cl | 16 ++- test/SemaOpenCL/builtins-amdgcn.cl | 6 - 7 files changed, 133 insertions(+), 70 deletions(-) create mode 100644 test/CodeGenOpenCL/builtins-amdgcn-error.cl create mode 100644 test/CodeGenOpenCL/builtins-amdgcn-vi.cl delete mode 100644 test/SemaOpenCL/builtins-amdgcn.cl diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index d712789366..d75142a0ca 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/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") @@ -52,7 +56,8 @@ BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n") //===----------------------------------------------------------------------===// // 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", "nc") BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc") #undef BUILTIN +#undef TARGET_BUILTIN diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp index 76cbb5e41a..134b84b8dc 100644 --- a/lib/Basic/Targets.cpp +++ b/lib/Basic/Targets.cpp @@ -1925,7 +1925,7 @@ static const char *const DataLayoutStringSI = "-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 TargetInfo { 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 &Features, + DiagnosticsEngine &Diags, StringRef CPU, + const std::vector &FeatureVec) const override; + ArrayRef 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(Name) + static GPUKind parseR600Name(StringRef Name) { + return llvm::StringSwitch(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(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 AMDGPUTargetInfo::getGCCRegNames() const { return llvm::makeArrayRef(GCCRegNames); } +bool AMDGPUTargetInfo::initFeatureMap( + llvm::StringMap &Features, + DiagnosticsEngine &Diags, StringRef CPU, + const std::vector &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") { diff --git a/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/test/CodeGenOpenCL/builtins-amdgcn-error.cl new file mode 100644 index 0000000000..89c3e490ec --- /dev/null +++ b/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}} +} + diff --git a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl new file mode 100644 index 0000000000..cda87a8e1e --- /dev/null +++ b/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(); +} diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 11448c3744..36164f898b 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -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) diff --git a/test/CodeGenOpenCL/builtins-r600.cl b/test/CodeGenOpenCL/builtins-r600.cl index 6369f3d83d..9ebcb1fe9d 100644 --- a/test/CodeGenOpenCL/builtins-r600.cl +++ b/test/CodeGenOpenCL/builtins-r600.cl @@ -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, float a) *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, float a, int b) *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 diff --git a/test/SemaOpenCL/builtins-amdgcn.cl b/test/SemaOpenCL/builtins-amdgcn.cl deleted file mode 100644 index bb342d058c..0000000000 --- a/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}} -} -- 2.40.0