[cfe-commits] r157403 - in /cfe/trunk: include/clang/Basic/BuiltinsNVPTX.def include/clang/Basic/BuiltinsPTX.def include/clang/Basic/TargetBuiltins.h lib/Basic/Targets.cpp lib/CodeGen/TargetInfo.cpp test/CodeGen/builtins-nvptx.c test/CodeGen/builtins-ptx.c test/CodeGen/nvptx-cc.c test/CodeGen/ptx-cc.c test/CodeGenCUDA/address-spaces.cu test/CodeGenCUDA/ptx-kernels.cu test/CodeGenOpenCL/ptx-calls.cl test/CodeGenOpenCL/ptx-kernels.cl
Justin Holewinski
jholewinski at nvidia.com
Thu May 24 10:43:12 PDT 2012
Author: jholewinski
Date: Thu May 24 12:43:12 2012
New Revision: 157403
URL: http://llvm.org/viewvc/llvm-project?rev=157403&view=rev
Log:
Replace PTX back-end with NVPTX back-end in all places where Clang cares
NV_CONTRIB
Added:
cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def
- copied, changed from r157399, cfe/trunk/include/clang/Basic/BuiltinsPTX.def
cfe/trunk/test/CodeGen/builtins-nvptx.c
- copied, changed from r157399, cfe/trunk/test/CodeGen/builtins-ptx.c
cfe/trunk/test/CodeGen/nvptx-cc.c
- copied, changed from r157399, cfe/trunk/test/CodeGen/ptx-cc.c
Removed:
cfe/trunk/include/clang/Basic/BuiltinsPTX.def
cfe/trunk/test/CodeGen/builtins-ptx.c
cfe/trunk/test/CodeGen/ptx-cc.c
Modified:
cfe/trunk/include/clang/Basic/TargetBuiltins.h
cfe/trunk/lib/Basic/Targets.cpp
cfe/trunk/lib/CodeGen/TargetInfo.cpp
cfe/trunk/test/CodeGenCUDA/address-spaces.cu
cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl
cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl
Copied: cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def (from r157399, cfe/trunk/include/clang/Basic/BuiltinsPTX.def)
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def?p2=cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def&p1=cfe/trunk/include/clang/Basic/BuiltinsPTX.def&r1=157399&r2=157403&rev=157403&view=diff
==============================================================================
(empty)
Removed: cfe/trunk/include/clang/Basic/BuiltinsPTX.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsPTX.def?rev=157402&view=auto
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsPTX.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsPTX.def (removed)
@@ -1,62 +0,0 @@
-//===--- BuiltinsPTX.def - PTX Builtin function database ----*- C++ -*-===//
-//
-// The LLVM Compiler Infrastructure
-//
-// This file is distributed under the University of Illinois Open Source
-// License. See LICENSE.TXT for details.
-//
-//===----------------------------------------------------------------------===//
-//
-// This file defines the PTX-specific builtin function database. Users of
-// this file must define the BUILTIN macro to make use of this information.
-//
-//===----------------------------------------------------------------------===//
-
-// The format of this database matches clang/Basic/Builtins.def.
-
-BUILTIN(__builtin_ptx_read_tid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_tid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_tid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_tid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_ntid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_ntid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_ntid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_ntid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_ctaid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_ctaid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_ctaid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_ctaid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_nctaid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_nctaid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_nctaid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_nctaid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_laneid, "i", "nc")
-BUILTIN(__builtin_ptx_read_warpid, "i", "nc")
-BUILTIN(__builtin_ptx_read_nwarpid, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_smid, "i", "nc")
-BUILTIN(__builtin_ptx_read_nsmid, "i", "nc")
-BUILTIN(__builtin_ptx_read_gridid, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_lanemask_eq, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_le, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_lt, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_ge, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_gt, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_clock, "i", "n")
-BUILTIN(__builtin_ptx_read_clock64, "Li", "n")
-
-BUILTIN(__builtin_ptx_read_pm0, "i", "n")
-BUILTIN(__builtin_ptx_read_pm1, "i", "n")
-BUILTIN(__builtin_ptx_read_pm2, "i", "n")
-BUILTIN(__builtin_ptx_read_pm3, "i", "n")
-
-BUILTIN(__builtin_ptx_bar_sync, "vi", "n")
-
-
-#undef BUILTIN
Modified: cfe/trunk/include/clang/Basic/TargetBuiltins.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetBuiltins.h?rev=157403&r1=157402&r2=157403&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/TargetBuiltins.h (original)
+++ cfe/trunk/include/clang/Basic/TargetBuiltins.h Thu May 24 12:43:12 2012
@@ -35,12 +35,12 @@
};
}
- /// 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
};
}
Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=157403&r1=157402&r2=157403&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Thu May 24 12:43:12 2012
@@ -946,57 +946,39 @@
} // 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<llvm::StringRef> 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 @@
// 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<bool> &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<bool> &Features,
- StringRef Name,
- bool Enabled) const {
+ bool NVPTXTargetInfo::setFeatureEnabled(llvm::StringMap<bool> &Features,
+ StringRef Name,
+ bool Enabled) const {
if(std::binary_search(AvailableFeatures.begin(), AvailableFeatures.end(),
Name)) {
Features[Name] = Enabled;
@@ -1055,117 +1039,28 @@
}
}
- 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 @@
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:
Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=157403&r1=157402&r2=157403&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Thu May 24 12:43:12 2012
@@ -2891,14 +2891,14 @@
}
//===----------------------------------------------------------------------===//
-// 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 @@
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 @@
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 @@
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 @@
}
-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<FunctionDecl>(D);
if (!FD) return;
@@ -3704,11 +3706,9 @@
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));
Copied: cfe/trunk/test/CodeGen/builtins-nvptx.c (from r157399, cfe/trunk/test/CodeGen/builtins-ptx.c)
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-nvptx.c?p2=cfe/trunk/test/CodeGen/builtins-nvptx.c&p1=cfe/trunk/test/CodeGen/builtins-ptx.c&r1=157399&r2=157403&rev=157403&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/builtins-ptx.c (original)
+++ cfe/trunk/test/CodeGen/builtins-nvptx.c Thu May 24 12:43:12 2012
@@ -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() {
Removed: cfe/trunk/test/CodeGen/builtins-ptx.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-ptx.c?rev=157402&view=auto
==============================================================================
--- cfe/trunk/test/CodeGen/builtins-ptx.c (original)
+++ cfe/trunk/test/CodeGen/builtins-ptx.c (removed)
@@ -1,99 +0,0 @@
-// RUN: %clang_cc1 -triple ptx32-unknown-unknown -emit-llvm -o %t %s
-// RUN: %clang_cc1 -triple ptx64-unknown-unknown -emit-llvm -o %t %s
-
-
-int read_tid() {
-
- int x = __builtin_ptx_read_tid_x();
- int y = __builtin_ptx_read_tid_y();
- int z = __builtin_ptx_read_tid_z();
- int w = __builtin_ptx_read_tid_w();
-
- return x + y + z + w;
-
-}
-
-int read_ntid() {
-
- int x = __builtin_ptx_read_ntid_x();
- int y = __builtin_ptx_read_ntid_y();
- int z = __builtin_ptx_read_ntid_z();
- int w = __builtin_ptx_read_ntid_w();
-
- return x + y + z + w;
-
-}
-
-int read_ctaid() {
-
- int x = __builtin_ptx_read_ctaid_x();
- int y = __builtin_ptx_read_ctaid_y();
- int z = __builtin_ptx_read_ctaid_z();
- int w = __builtin_ptx_read_ctaid_w();
-
- return x + y + z + w;
-
-}
-
-int read_nctaid() {
-
- int x = __builtin_ptx_read_nctaid_x();
- int y = __builtin_ptx_read_nctaid_y();
- int z = __builtin_ptx_read_nctaid_z();
- int w = __builtin_ptx_read_nctaid_w();
-
- return x + y + z + w;
-
-}
-
-int read_ids() {
-
- int a = __builtin_ptx_read_laneid();
- int b = __builtin_ptx_read_warpid();
- int c = __builtin_ptx_read_nwarpid();
- int d = __builtin_ptx_read_smid();
- int e = __builtin_ptx_read_nsmid();
- int f = __builtin_ptx_read_gridid();
-
- return a + b + c + d + e + f;
-
-}
-
-int read_lanemasks() {
-
- int a = __builtin_ptx_read_lanemask_eq();
- int b = __builtin_ptx_read_lanemask_le();
- int c = __builtin_ptx_read_lanemask_lt();
- int d = __builtin_ptx_read_lanemask_ge();
- int e = __builtin_ptx_read_lanemask_gt();
-
- return a + b + c + d + e;
-
-}
-
-
-long read_clocks() {
-
- int a = __builtin_ptx_read_clock();
- long b = __builtin_ptx_read_clock64();
-
- return (long)a + b;
-
-}
-
-int read_pms() {
-
- int a = __builtin_ptx_read_pm0();
- int b = __builtin_ptx_read_pm1();
- int c = __builtin_ptx_read_pm2();
- int d = __builtin_ptx_read_pm3();
-
- return a + b + c + d;
-
-}
-
-void sync() {
-
- __builtin_ptx_bar_sync(0);
-
-}
Copied: cfe/trunk/test/CodeGen/nvptx-cc.c (from r157399, cfe/trunk/test/CodeGen/ptx-cc.c)
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/nvptx-cc.c?p2=cfe/trunk/test/CodeGen/nvptx-cc.c&p1=cfe/trunk/test/CodeGen/ptx-cc.c&r1=157399&r2=157403&rev=157403&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/ptx-cc.c (original)
+++ cfe/trunk/test/CodeGen/nvptx-cc.c Thu May 24 12:43:12 2012
@@ -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) {
Removed: cfe/trunk/test/CodeGen/ptx-cc.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/ptx-cc.c?rev=157402&view=auto
==============================================================================
--- cfe/trunk/test/CodeGen/ptx-cc.c (original)
+++ cfe/trunk/test/CodeGen/ptx-cc.c (removed)
@@ -1,9 +0,0 @@
-// 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
-
-// Just make sure Clang uses the proper calling convention for the PTX back-end.
-// If something is wrong, the back-end will fail.
-void foo(float* a,
- float* b) {
- a[0] = b[0];
-}
Modified: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=157403&r1=157402&r2=157403&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Thu May 24 12:43:12 2012
@@ -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++;
}
Modified: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu?rev=157403&r1=157402&r2=157403&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Thu May 24 12:43:12 2012
@@ -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"
Modified: cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl?rev=157403&r1=157402&r2=157403&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/ptx-calls.cl Thu May 24 12:43:12 2012
@@ -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() {
}
Modified: cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl?rev=157403&r1=157402&r2=157403&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/ptx-kernels.cl Thu May 24 12:43:12 2012
@@ -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() {
}
More information about the cfe-commits
mailing list