maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

Replace PTX back-end with NVPTX back-end in all places where Clang cares 

NV_CONTRIB

llvm-svn: 157403
This commit is contained in:
Justin Holewinski 2012-05-24 17:43:12 +00:00
parent 6b7b7e66d1
commit 83e9668133
10 changed files with 73 additions and 183 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

0
clang/include/clang/Basic/BuiltinsPTX.def → clang/include/clang/Basic/BuiltinsNVPTX.def
View File

6
clang/include/clang/Basic/TargetBuiltins.h
View File

@ -35,12 +35,12 @@ namespace clang {
};
}
/// PTX builtins
namespace PTX {
/// NVPTX builtins
namespace NVPTX {
enum {
LastTIBuiltin = clang::Builtin::FirstTSBuiltin-1,
#define BUILTIN(ID, TYPE, ATTRS) BI##ID,
#include "clang/Basic/BuiltinsPTX.def"
#include "clang/Basic/BuiltinsNVPTX.def"
LastTSBuiltin
};
}

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

@ -945,137 +945,6 @@ public:
};
} // end anonymous namespace.
namespace {
static const unsigned PTXAddrSpaceMap[] = {
0, // opencl_global
4, // opencl_local
1, // opencl_constant
0, // cuda_device
1, // cuda_constant
4, // cuda_shared
};
class PTXTargetInfo : public TargetInfo {
static const char * const GCCRegNames[];
static const Builtin::Info BuiltinInfo[];
std::vector<llvm::StringRef> AvailableFeatures;
public:
PTXTargetInfo(const std::string& triple) : TargetInfo(triple) {
BigEndian = false;
TLSSupported = false;
LongWidth = LongAlign = 64;
AddrSpaceMap = &PTXAddrSpaceMap;
// Define available target features
// These must be defined in sorted order!
AvailableFeatures.push_back("compute10");
AvailableFeatures.push_back("compute11");
AvailableFeatures.push_back("compute12");
AvailableFeatures.push_back("compute13");
AvailableFeatures.push_back("compute20");
AvailableFeatures.push_back("double");
AvailableFeatures.push_back("no-fma");
AvailableFeatures.push_back("ptx20");
AvailableFeatures.push_back("ptx21");
AvailableFeatures.push_back("ptx22");
AvailableFeatures.push_back("ptx23");
AvailableFeatures.push_back("sm10");
AvailableFeatures.push_back("sm11");
AvailableFeatures.push_back("sm12");
AvailableFeatures.push_back("sm13");
AvailableFeatures.push_back("sm20");
AvailableFeatures.push_back("sm21");
AvailableFeatures.push_back("sm22");
AvailableFeatures.push_back("sm23");
}
virtual void getTargetDefines(const LangOptions &Opts,
MacroBuilder &Builder) const {
Builder.defineMacro("__PTX__");
}
virtual void getTargetBuiltins(const Builtin::Info *&Records,
unsigned &NumRecords) const {
Records = BuiltinInfo;
NumRecords = clang::PTX::LastTSBuiltin-Builtin::FirstTSBuiltin;
}
virtual bool hasFeature(StringRef Feature) const {
return Feature == "ptx";
}
virtual void getGCCRegNames(const char * const *&Names,
unsigned &NumNames) const;
virtual void getGCCRegAliases(const GCCRegAlias *&Aliases,
unsigned &NumAliases) const {
// No aliases.
Aliases = 0;
NumAliases = 0;
}
virtual bool validateAsmConstraint(const char *&Name,
TargetInfo::ConstraintInfo &info) const {
// FIXME: implement
return true;
}
virtual const char *getClobbers() const {
// FIXME: Is this really right?
return "";
}
virtual const char *getVAListDeclaration() const {
// FIXME: implement
return "typedef char* __builtin_va_list;";
}
virtual bool setFeatureEnabled(llvm::StringMap<bool> &Features,
StringRef Name,
bool Enabled) const;
};
const Builtin::Info PTXTargetInfo::BuiltinInfo[] = {
#define BUILTIN(ID, TYPE, ATTRS) { #ID, TYPE, ATTRS, 0, ALL_LANGUAGES },
#define LIBBUILTIN(ID, TYPE, ATTRS, HEADER) { #ID, TYPE, ATTRS, HEADER,\
ALL_LANGUAGES },
#include "clang/Basic/BuiltinsPTX.def"
};
const char * const PTXTargetInfo::GCCRegNames[] = {
"r0"
};
void PTXTargetInfo::getGCCRegNames(const char * const *&Names,
unsigned &NumNames) const {
Names = GCCRegNames;
NumNames = llvm::array_lengthof(GCCRegNames);
}
bool PTXTargetInfo::setFeatureEnabled(llvm::StringMap<bool> &Features,
StringRef Name,
bool Enabled) const {
if(std::binary_search(AvailableFeatures.begin(), AvailableFeatures.end(),
Name)) {
Features[Name] = Enabled;
return true;
} else {
return false;
}
}
class PTX32TargetInfo : public PTXTargetInfo {
public:
PTX32TargetInfo(const std::string& triple) : PTXTargetInfo(triple) {
PointerWidth = PointerAlign = 32;
SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedInt;
DescriptionString
= "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64";
}
};
class PTX64TargetInfo : public PTXTargetInfo {
public:
PTX64TargetInfo(const std::string& triple) : PTXTargetInfo(triple) {
PointerWidth = PointerAlign = 64;
SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedLongLong;
DescriptionString
= "e-p:64:64-i64:64:64-f64:64:64-n1:8:16:32:64";
}
};
}
namespace {
static const unsigned NVPTXAddrSpaceMap[] = {
1, // opencl_global
@ -1087,25 +956,29 @@ namespace {
};
class NVPTXTargetInfo : public TargetInfo {
static const char * const GCCRegNames[];
static const Builtin::Info BuiltinInfo[];
std::vector<llvm::StringRef> AvailableFeatures;
public:
NVPTXTargetInfo(const std::string& triple) : TargetInfo(triple) {
BigEndian = false;
TLSSupported = false;
LongWidth = LongAlign = 64;
AddrSpaceMap = &NVPTXAddrSpaceMap;
// Define available target features
// These must be defined in sorted order!
}
virtual void getTargetDefines(const LangOptions &Opts,
MacroBuilder &Builder) const {
Builder.defineMacro("__PTX__");
Builder.defineMacro("__NVPTX__");
}
virtual void getTargetBuiltins(const Builtin::Info *&Records,
unsigned &NumRecords) const {
// FIXME: implement.
Records = 0;
NumRecords = 0;
Records = BuiltinInfo;
NumRecords = clang::NVPTX::LastTSBuiltin-Builtin::FirstTSBuiltin;
}
virtual bool hasFeature(StringRef Feature) const {
return Feature == "nvptx";
return Feature == "ptx" || Feature == "nvptx";
}
virtual void getGCCRegNames(const char * const *&Names,
@ -1130,8 +1003,18 @@ namespace {
return "typedef char* __builtin_va_list;";
}
virtual bool setCPU(const std::string &Name) {
return Name == "sm_10";
return Name == "sm_10" || Name == "sm_13" || Name == "sm_20";
}
virtual bool setFeatureEnabled(llvm::StringMap<bool> &Features,
StringRef Name,
bool Enabled) const;
};
const Builtin::Info NVPTXTargetInfo::BuiltinInfo[] = {
#define BUILTIN(ID, TYPE, ATTRS) { #ID, TYPE, ATTRS, 0, ALL_LANGUAGES },
#define LIBBUILTIN(ID, TYPE, ATTRS, HEADER) { #ID, TYPE, ATTRS, HEADER,\
ALL_LANGUAGES },
#include "clang/Basic/BuiltinsNVPTX.def"
};
const char * const NVPTXTargetInfo::GCCRegNames[] = {
@ -1144,6 +1027,18 @@ namespace {
NumNames = llvm::array_lengthof(GCCRegNames);
}
bool NVPTXTargetInfo::setFeatureEnabled(llvm::StringMap<bool> &Features,
StringRef Name,
bool Enabled) const {
if(std::binary_search(AvailableFeatures.begin(), AvailableFeatures.end(),
Name)) {
Features[Name] = Enabled;
return true;
} else {
return false;
}
}
class NVPTX32TargetInfo : public NVPTXTargetInfo {
public:
NVPTX32TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) {
@ -4139,11 +4034,6 @@ static TargetInfo *AllocateTarget(const std::string &T) {
return new PPC64TargetInfo(T);
}
case llvm::Triple::ptx32:
return new PTX32TargetInfo(T);
case llvm::Triple::ptx64:
return new PTX64TargetInfo(T);
case llvm::Triple::nvptx:
return new NVPTX32TargetInfo(T);
case llvm::Triple::nvptx64:

32
clang/lib/CodeGen/TargetInfo.cpp
View File

@ -2891,14 +2891,14 @@ llvm::Value *ARMABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty,
}
//===----------------------------------------------------------------------===//
// PTX ABI Implementation
// NVPTX ABI Implementation
//===----------------------------------------------------------------------===//
namespace {
class PTXABIInfo : public ABIInfo {
class NVPTXABIInfo : public ABIInfo {
public:
PTXABIInfo(CodeGenTypes &CGT) : ABIInfo(CGT) {}
NVPTXABIInfo(CodeGenTypes &CGT) : ABIInfo(CGT) {}
ABIArgInfo classifyReturnType(QualType RetTy) const;
ABIArgInfo classifyArgumentType(QualType Ty) const;
@ -2908,16 +2908,16 @@ public:
CodeGenFunction &CFG) const;
};
class PTXTargetCodeGenInfo : public TargetCodeGenInfo {
class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
public:
PTXTargetCodeGenInfo(CodeGenTypes &CGT)
: TargetCodeGenInfo(new PTXABIInfo(CGT)) {}
NVPTXTargetCodeGenInfo(CodeGenTypes &CGT)
: TargetCodeGenInfo(new NVPTXABIInfo(CGT)) {}
virtual void SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const;
};
ABIArgInfo PTXABIInfo::classifyReturnType(QualType RetTy) const {
ABIArgInfo NVPTXABIInfo::classifyReturnType(QualType RetTy) const {
if (RetTy->isVoidType())
return ABIArgInfo::getIgnore();
if (isAggregateTypeForABI(RetTy))
@ -2925,14 +2925,14 @@ ABIArgInfo PTXABIInfo::classifyReturnType(QualType RetTy) const {
return ABIArgInfo::getDirect();
}
ABIArgInfo PTXABIInfo::classifyArgumentType(QualType Ty) const {
ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType Ty) const {
if (isAggregateTypeForABI(Ty))
return ABIArgInfo::getIndirect(0);
return ABIArgInfo::getDirect();
}
void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
void NVPTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
FI.getReturnInfo() = classifyReturnType(FI.getReturnType());
for (CGFunctionInfo::arg_iterator it = FI.arg_begin(), ie = FI.arg_end();
it != ie; ++it)
@ -2943,6 +2943,8 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
return;
// Calling convention as default by an ABI.
// We're still using the PTX_Kernel/PTX_Device calling conventions here,
// but we should switch to NVVM metadata later on.
llvm::CallingConv::ID DefaultCC;
const LangOptions &LangOpts = getContext().getLangOpts();
if (LangOpts.OpenCL || LangOpts.CUDA) {
@ -2961,13 +2963,13 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const {
}
llvm::Value *PTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty,
llvm::Value *NVPTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty,
CodeGenFunction &CFG) const {
llvm_unreachable("PTX does not support varargs");
llvm_unreachable("NVPTX does not support varargs");
}
void PTXTargetCodeGenInfo::SetTargetAttributes(const Decl *D,
llvm::GlobalValue *GV,
void NVPTXTargetCodeGenInfo::
SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const{
const FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
if (!FD) return;
@ -3704,11 +3706,9 @@ const TargetCodeGenInfo &CodeGenModule::getTargetCodeGenInfo() {
case llvm::Triple::ppc64:
return *(TheTargetCodeGenInfo = new PPC64TargetCodeGenInfo(Types));
case llvm::Triple::ptx32:
case llvm::Triple::ptx64:
case llvm::Triple::nvptx:
case llvm::Triple::nvptx64:
return *(TheTargetCodeGenInfo = new PTXTargetCodeGenInfo(Types));
return *(TheTargetCodeGenInfo = new NVPTXTargetCodeGenInfo(Types));
case llvm::Triple::mblaze:
return *(TheTargetCodeGenInfo = new MBlazeTargetCodeGenInfo(Types));

4
clang/test/CodeGen/builtins-ptx.c → clang/test/CodeGen/builtins-nvptx.c
View File

@ -1,5 +1,5 @@
// RUN: %clang_cc1 -triple ptx32-unknown-unknown -emit-llvm -o %t %s
// RUN: %clang_cc1 -triple ptx64-unknown-unknown -emit-llvm -o %t %s
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -emit-llvm -o %t %s
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -emit-llvm -o %t %s
int read_tid() {

6
clang/test/CodeGen/ptx-cc.c → clang/test/CodeGen/nvptx-cc.c
View File

@ -1,7 +1,7 @@
// RUN: %clang_cc1 -triple ptx32-unknown-unknown -O3 -S -o %t %s -emit-llvm
// RUN: %clang_cc1 -triple ptx64-unknown-unknown -O3 -S -o %t %s -emit-llvm
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -O3 -S -o %t %s -emit-llvm
// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -O3 -S -o %t %s -emit-llvm
// Just make sure Clang uses the proper calling convention for the PTX back-end.
// Just make sure Clang uses the proper calling convention for the NVPTX back-end.
// If something is wrong, the back-end will fail.
void foo(float* a,
float* b) {

14
clang/test/CodeGenCUDA/address-spaces.cu
View File

@ -1,24 +1,24 @@
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple ptx32-unknown-unknown | FileCheck %s
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
#include "../SemaCUDA/cuda.h"
// CHECK: @i = global
// CHECK: @i = addrspace(1) global
__device__ int i;
// CHECK: @j = addrspace(1) global
// CHECK: @j = addrspace(4) global
__constant__ int j;
// CHECK: @k = addrspace(4) global
// CHECK: @k = addrspace(3) global
__shared__ int k;
__device__ void foo() {
// CHECK: load i32* @i
// CHECK: load i32* bitcast (i32 addrspace(1)* @i to i32*)
i++;
// CHECK: load i32* bitcast (i32 addrspace(1)* @j to i32*)
// CHECK: load i32* bitcast (i32 addrspace(4)* @j to i32*)
j++;
// CHECK: load i32* bitcast (i32 addrspace(4)* @k to i32*)
// CHECK: load i32* bitcast (i32 addrspace(3)* @k to i32*)
k++;
}

2
clang/test/CodeGenCUDA/ptx-kernels.cu
View File

@ -1,4 +1,4 @@
// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
#include "../SemaCUDA/cuda.h"

2
clang/test/CodeGenOpenCL/ptx-calls.cl
View File

@ -1,4 +1,4 @@
// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
void device_function() {
}

2
clang/test/CodeGenOpenCL/ptx-kernels.cl
View File

@ -1,4 +1,4 @@
// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -emit-llvm -o - | FileCheck %s
void device_function() {
}