replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

AMDGPU: Verify subtarget specific builtins 

Cleanup setup of subtarget features.

llvm-svn: 272091
This commit is contained in:
Matt Arsenault 2016-06-08 01:56:42 +00:00
parent 76f425211e
commit 250024f905
7 changed files with 133 additions and 70 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -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

132
clang/lib/Basic/Targets.cpp
View File

@ -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<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::getGCCRegNames() const {
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) \

18
clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl Normal file
View File

@ -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}}
}

12
clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Normal file
View File

@ -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();
}

7
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@ -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)

16
clang/test/CodeGenOpenCL/builtins-r600.cl
View File

@ -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

6
clang/test/SemaOpenCL/builtins-amdgcn.cl
View File

@ -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}}
}