diff --git a/clang/include/clang/Basic/BuiltinsPTX.def b/clang/include/clang/Basic/BuiltinsNVPTX.def similarity index 100% rename from clang/include/clang/Basic/BuiltinsPTX.def rename to clang/include/clang/Basic/BuiltinsNVPTX.def diff --git a/clang/include/clang/Basic/TargetBuiltins.h b/clang/include/clang/Basic/TargetBuiltins.h index 7c04bf7edf5b..e9b9f8552114 100644 --- a/clang/include/clang/Basic/TargetBuiltins.h +++ b/clang/include/clang/Basic/TargetBuiltins.h @@ -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 }; } diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp index e9752700f052..e38ca550c34b 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -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 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 &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 &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 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 &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,28 +1027,40 @@ namespace { NumNames = llvm::array_lengthof(GCCRegNames); } + bool NVPTXTargetInfo::setFeatureEnabled(llvm::StringMap &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) { + NVPTX32TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) { PointerWidth = PointerAlign = 32; - SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedInt; + SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedInt; DescriptionString = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-" "f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-" "n16:32:64"; - } + } }; class NVPTX64TargetInfo : public NVPTXTargetInfo { public: - NVPTX64TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) { + NVPTX64TargetInfo(const std::string& triple) : NVPTXTargetInfo(triple) { PointerWidth = PointerAlign = 64; - SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedLongLong; + SizeType = PtrDiffType = IntPtrType = TargetInfo::UnsignedLongLong; DescriptionString = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-" "f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-" "n16:32:64"; - } + } }; } @@ -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: diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index 97ca23813019..357b3fe53747 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -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,14 +2963,14 @@ void PTXABIInfo::computeInfo(CGFunctionInfo &FI) const { } -llvm::Value *PTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty, - CodeGenFunction &CFG) const { - llvm_unreachable("PTX does not support varargs"); +llvm::Value *NVPTXABIInfo::EmitVAArg(llvm::Value *VAListAddr, QualType Ty, + CodeGenFunction &CFG) const { + llvm_unreachable("NVPTX does not support varargs"); } -void PTXTargetCodeGenInfo::SetTargetAttributes(const Decl *D, - llvm::GlobalValue *GV, - CodeGen::CodeGenModule &M) const{ +void NVPTXTargetCodeGenInfo:: +SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV, + CodeGen::CodeGenModule &M) const{ const FunctionDecl *FD = dyn_cast(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)); diff --git a/clang/test/CodeGen/builtins-ptx.c b/clang/test/CodeGen/builtins-nvptx.c similarity index 92% rename from clang/test/CodeGen/builtins-ptx.c rename to clang/test/CodeGen/builtins-nvptx.c index 6dd10188e9fb..4a094bbd7e33 100644 --- a/clang/test/CodeGen/builtins-ptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -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() { diff --git a/clang/test/CodeGen/ptx-cc.c b/clang/test/CodeGen/nvptx-cc.c similarity index 53% rename from clang/test/CodeGen/ptx-cc.c rename to clang/test/CodeGen/nvptx-cc.c index 2212d4260b35..1c0d943f956b 100644 --- a/clang/test/CodeGen/ptx-cc.c +++ b/clang/test/CodeGen/nvptx-cc.c @@ -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) { diff --git a/clang/test/CodeGenCUDA/address-spaces.cu b/clang/test/CodeGenCUDA/address-spaces.cu index 2da61ec95a6d..61d4d6b6ba48 100644 --- a/clang/test/CodeGenCUDA/address-spaces.cu +++ b/clang/test/CodeGenCUDA/address-spaces.cu @@ -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++; } diff --git a/clang/test/CodeGenCUDA/ptx-kernels.cu b/clang/test/CodeGenCUDA/ptx-kernels.cu index ecca8519af63..f0bf2952a126 100644 --- a/clang/test/CodeGenCUDA/ptx-kernels.cu +++ b/clang/test/CodeGenCUDA/ptx-kernels.cu @@ -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" diff --git a/clang/test/CodeGenOpenCL/ptx-calls.cl b/clang/test/CodeGenOpenCL/ptx-calls.cl index 6f336405c301..34a21c6c1da9 100644 --- a/clang/test/CodeGenOpenCL/ptx-calls.cl +++ b/clang/test/CodeGenOpenCL/ptx-calls.cl @@ -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() { } diff --git a/clang/test/CodeGenOpenCL/ptx-kernels.cl b/clang/test/CodeGenOpenCL/ptx-kernels.cl index 4d6fa1084d4b..1d7e497b7c01 100644 --- a/clang/test/CodeGenOpenCL/ptx-kernels.cl +++ b/clang/test/CodeGenOpenCL/ptx-kernels.cl @@ -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() { }