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 @@ -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 &Features, + DiagnosticsEngine &Diags, StringRef CPU, + const std::vector &FeatureVec) const override; + ArrayRef 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 &Features, + DiagnosticsEngine &Diags, StringRef CPU, + const std::vector &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: 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}} -}