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 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 &Features, + DiagnosticsEngine &Diags, StringRef CPU, + const std::vector &FeatureVec) const override; + ArrayRef 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(Name) + static GPUKind parseR600Name(StringRef Name) { + return llvm::StringSwitch(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(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 &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) \ 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: 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: 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.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/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}} -}