From 83e9668133232be86a2b394651ecaf503e2d86fd Mon Sep 17 00:00:00 2001 From: Justin Holewinski Date: Thu, 24 May 2012 17:43:12 +0000 Subject: [PATCH] Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403 --- .../Basic/{BuiltinsPTX.def => BuiltinsNVPTX.def} | 0 clang/include/clang/Basic/TargetBuiltins.h | 6 +- clang/lib/Basic/Targets.cpp | 170 ++++----------------- clang/lib/CodeGen/TargetInfo.cpp | 36 ++--- .../CodeGen/{builtins-ptx.c => builtins-nvptx.c} | 4 +- clang/test/CodeGen/{ptx-cc.c => nvptx-cc.c} | 6 +- clang/test/CodeGenCUDA/address-spaces.cu | 14 +- clang/test/CodeGenCUDA/ptx-kernels.cu | 2 +- clang/test/CodeGenOpenCL/ptx-calls.cl | 2 +- clang/test/CodeGenOpenCL/ptx-kernels.cl | 2 +- 10 files changed, 66 insertions(+), 176 deletions(-) rename clang/include/clang/Basic/{BuiltinsPTX.def => BuiltinsNVPTX.def} (100%) rename clang/test/CodeGen/{builtins-ptx.c => builtins-nvptx.c} (92%) rename clang/test/CodeGen/{ptx-cc.c => nvptx-cc.c} (53%) 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 7c04bf7..e9b9f85 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 e975270..e38ca550 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -946,57 +946,39 @@ 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 + static const unsigned NVPTXAddrSpaceMap[] = { + 1, // opencl_global + 3, // opencl_local + 4, // opencl_constant + 1, // cuda_device + 4, // cuda_constant + 3, // cuda_shared }; - class PTXTargetInfo : public TargetInfo { + class NVPTXTargetInfo : public TargetInfo { static const char * const GCCRegNames[]; static const Builtin::Info BuiltinInfo[]; std::vector AvailableFeatures; public: - PTXTargetInfo(const std::string& triple) : TargetInfo(triple) { + NVPTXTargetInfo(const std::string& triple) : TargetInfo(triple) { BigEndian = false; TLSSupported = false; LongWidth = LongAlign = 64; - AddrSpaceMap = &PTXAddrSpaceMap; + AddrSpaceMap = &NVPTXAddrSpaceMap; // 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"); + // 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 { Records = BuiltinInfo; - NumRecords = clang::PTX::LastTSBuiltin-Builtin::FirstTSBuiltin; + NumRecords = clang::NVPTX::LastTSBuiltin-Builtin::FirstTSBuiltin; } virtual bool hasFeature(StringRef Feature) const { - return Feature == "ptx"; + return Feature == "ptx" || Feature == "nvptx"; } virtual void getGCCRegNames(const char * const *&Names, @@ -1020,32 +1002,34 @@ namespace { // FIXME: implement return "typedef char* __builtin_va_list;"; } - + virtual bool setCPU(const std::string &Name) { + return Name == "sm_10" || Name == "sm_13" || Name == "sm_20"; + } virtual bool setFeatureEnabled(llvm::StringMap &Features, StringRef Name, bool Enabled) const; }; - const Builtin::Info PTXTargetInfo::BuiltinInfo[] = { + 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/BuiltinsPTX.def" +#include "clang/Basic/BuiltinsNVPTX.def" }; - const char * const PTXTargetInfo::GCCRegNames[] = { + const char * const NVPTXTargetInfo::GCCRegNames[] = { "r0" }; - void PTXTargetInfo::getGCCRegNames(const char * const *&Names, + void NVPTXTargetInfo::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 { + bool NVPTXTargetInfo::setFeatureEnabled(llvm::StringMap &Features, + StringRef Name, + bool Enabled) const { if(std::binary_search(AvailableFeatures.begin(), AvailableFeatures.end(), Name)) { Features[Name] = Enabled; @@ -1055,117 +1039,28 @@ namespace { } } - 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 - 3, // opencl_local - 4, // opencl_constant - 1, // cuda_device - 4, // cuda_constant - 3, // cuda_shared - }; - class NVPTXTargetInfo : public TargetInfo { - static const char * const GCCRegNames[]; - public: - NVPTXTargetInfo(const std::string& triple) : TargetInfo(triple) { - BigEndian = false; - TLSSupported = false; - LongWidth = LongAlign = 64; - AddrSpaceMap = &NVPTXAddrSpaceMap; - } - virtual void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const { - Builder.defineMacro("__PTX__"); - } - virtual void getTargetBuiltins(const Builtin::Info *&Records, - unsigned &NumRecords) const { - // FIXME: implement. - Records = 0; - NumRecords = 0; - } - virtual bool hasFeature(StringRef Feature) const { - return Feature == "nvptx"; - } - - 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 setCPU(const std::string &Name) { - return Name == "sm_10"; - } - }; - - const char * const NVPTXTargetInfo::GCCRegNames[] = { - "r0" - }; - - void NVPTXTargetInfo::getGCCRegNames(const char * const *&Names, - unsigned &NumNames) const { - Names = GCCRegNames; - NumNames = llvm::array_lengthof(GCCRegNames); - } - 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 97ca238..357b3fe 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 6dd1018..4a094bb 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 2212d42..1c0d943 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 2da61ec..61d4d6b 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 ecca851..f0bf295 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 6f33640..34a21c6 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 4d6fa10..1d7e497 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() { } -- 2.7.4