[CUDA] 32-bit NVPTX should have 32-bit long type.

Currently it's 64-bit which will lead to mismatch between host and
device code if we compile for i386.

Differential Revision: http://reviews.llvm.org/D13181

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@248753 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsNVPTX.def b/include/clang/Basic/BuiltinsNVPTX.def
index f1cc961dd0e6e34ff080a34705615618abd4c07c..3ab6413bb0ec3cd2b01ab840573b6c6db78eac47 100644 (file)
--- a/include/clang/Basic/BuiltinsNVPTX.def
+++ b/include/clang/Basic/BuiltinsNVPTX.def
@@ -50,7 +50,7 @@ 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_clock64, "LLi", "n")
 
 BUILTIN(__builtin_ptx_read_pm0, "i", "n")
 BUILTIN(__builtin_ptx_read_pm1, "i", "n")

diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp
index da91694768ba5aef40b8b58743fc981298ad3674..a544573567dc9ee2b5cd924f06e2047bc75dcfeb 100644 (file)
--- a/lib/Basic/Targets.cpp
+++ b/lib/Basic/Targets.cpp
@@ -1716,6 +1716,7 @@ void NVPTXTargetInfo::getGCCRegNames(const char *const *&Names,
 class NVPTX32TargetInfo : public NVPTXTargetInfo {
 public:
   NVPTX32TargetInfo(const llvm::Triple &Triple) : NVPTXTargetInfo(Triple) {
+    LongWidth = LongAlign = 32;
     PointerWidth = PointerAlign = 32;
     SizeType = TargetInfo::UnsignedInt;
     PtrDiffType = TargetInfo::SignedInt;

diff --git a/test/CodeGen/builtins-nvptx.c b/test/CodeGen/builtins-nvptx.c
index 498da476d102da572f20cfa24a75c1c26e824f22..d1abb9c95090d310f9a966ef02ff36867e30fa77 100644 (file)
--- a/test/CodeGen/builtins-nvptx.c
+++ b/test/CodeGen/builtins-nvptx.c
@@ -109,16 +109,15 @@ __device__ int read_lanemasks() {
 
 }
 
-__device__ long read_clocks() {
+__device__ long long read_clocks() {
 
 // CHECK: call i32 @llvm.ptx.read.clock()
 // CHECK: call i64 @llvm.ptx.read.clock64()
 
   int a = __builtin_ptx_read_clock();
-  long b = __builtin_ptx_read_clock64();
-
-  return (long)a + b;
+  long long b = __builtin_ptx_read_clock64();
 
+  return a + b;
 }
 
 __device__ int read_pms() {

diff --git a/test/CodeGen/nvptx-inlineasm-ptx.c b/test/CodeGen/nvptx-inlineasm-ptx.c
index 0a19123ee07479007a131138f40649a5f551ee7f..e5345d93de06d7edc7d2599122a225ef630495f4 100644 (file)
--- a/test/CodeGen/nvptx-inlineasm-ptx.c
+++ b/test/CodeGen/nvptx-inlineasm-ptx.c
@@ -8,8 +8,8 @@ void constraints() {
   unsigned short us;
   int            i;
   unsigned int   ui;
-  long           l;
-  unsigned long  ul;
+  long long      ll;
+  unsigned long long ull;
   float          f;
   double         d;
 
@@ -29,9 +29,9 @@ void constraints() {
   asm volatile ("mov.b32 %0, %1;" : "=r"(ui) : "r"(ui));
 
   // CHECK: i64 asm sideeffect "mov.b64 $0, $1;", "=l,l"
-  asm volatile ("mov.b64 %0, %1;" : "=l"(l) : "l"(l));
+  asm volatile ("mov.b64 %0, %1;" : "=l"(ll) : "l"(ll));
   // CHECK: i64 asm sideeffect "mov.b64 $0, $1;", "=l,l"
-  asm volatile ("mov.b64 %0, %1;" : "=l"(ul) : "l"(ul));
+  asm volatile ("mov.b64 %0, %1;" : "=l"(ull) : "l"(ull));
 
   // CHECK: float asm sideeffect "mov.b32 $0, $1;", "=f,f"
   asm volatile ("mov.b32 %0, %1;" : "=f"(f) : "f"(f));

diff --git a/test/Preprocessor/init.c b/test/Preprocessor/init.c
index bb5de5484d9d5c6a6761ae1231d04aa7e337333c..a2f5bf72544a164d1188b047d857ed2b40e77cab 100644 (file)
--- a/test/Preprocessor/init.c
+++ b/test/Preprocessor/init.c
@@ -4737,10 +4737,10 @@
 // NVPTX32:#define __INT_FAST32_FMTi__ "i"
 // NVPTX32:#define __INT_FAST32_MAX__ 2147483647
 // NVPTX32:#define __INT_FAST32_TYPE__ int
-// NVPTX32:#define __INT_FAST64_FMTd__ "ld"
-// NVPTX32:#define __INT_FAST64_FMTi__ "li"
+// NVPTX32:#define __INT_FAST64_FMTd__ "lld"
+// NVPTX32:#define __INT_FAST64_FMTi__ "lli"
 // NVPTX32:#define __INT_FAST64_MAX__ 9223372036854775807L
-// NVPTX32:#define __INT_FAST64_TYPE__ long int
+// NVPTX32:#define __INT_FAST64_TYPE__ long long int
 // NVPTX32:#define __INT_FAST8_FMTd__ "hhd"
 // NVPTX32:#define __INT_FAST8_FMTi__ "hhi"
 // NVPTX32:#define __INT_FAST8_MAX__ 127
@@ -4753,10 +4753,10 @@
 // NVPTX32:#define __INT_LEAST32_FMTi__ "i"
 // NVPTX32:#define __INT_LEAST32_MAX__ 2147483647
 // NVPTX32:#define __INT_LEAST32_TYPE__ int
-// NVPTX32:#define __INT_LEAST64_FMTd__ "ld"
-// NVPTX32:#define __INT_LEAST64_FMTi__ "li"
+// NVPTX32:#define __INT_LEAST64_FMTd__ "lld"
+// NVPTX32:#define __INT_LEAST64_FMTi__ "lli"
 // NVPTX32:#define __INT_LEAST64_MAX__ 9223372036854775807L
-// NVPTX32:#define __INT_LEAST64_TYPE__ long int
+// NVPTX32:#define __INT_LEAST64_TYPE__ long long int
 // NVPTX32:#define __INT_LEAST8_FMTd__ "hhd"
 // NVPTX32:#define __INT_LEAST8_FMTi__ "hhi"
 // NVPTX32:#define __INT_LEAST8_MAX__ 127
@@ -4777,7 +4777,7 @@
 // NVPTX32:#define __LDBL_MIN__ 2.2250738585072014e-308L
 // NVPTX32:#define __LITTLE_ENDIAN__ 1
 // NVPTX32:#define __LONG_LONG_MAX__ 9223372036854775807LL
-// NVPTX32:#define __LONG_MAX__ 9223372036854775807L
+// NVPTX32:#define __LONG_MAX__ 2147483647L
 // NVPTX32-NOT:#define __LP64__
 // NVPTX32:#define __NVPTX__ 1
 // NVPTX32:#define __POINTER_WIDTH__ 32
@@ -4794,7 +4794,7 @@
 // NVPTX32:#define __SIZEOF_INT__ 4
 // NVPTX32:#define __SIZEOF_LONG_DOUBLE__ 8
 // NVPTX32:#define __SIZEOF_LONG_LONG__ 8
-// NVPTX32:#define __SIZEOF_LONG__ 8
+// NVPTX32:#define __SIZEOF_LONG__ 4
 // NVPTX32:#define __SIZEOF_POINTER__ 4
 // NVPTX32:#define __SIZEOF_PTRDIFF_T__ 4
 // NVPTX32:#define __SIZEOF_SHORT__ 2
@@ -4828,7 +4828,7 @@
 // NVPTX32:#define __UINT_FAST32_MAX__ 4294967295U
 // NVPTX32:#define __UINT_FAST32_TYPE__ unsigned int
 // NVPTX32:#define __UINT_FAST64_MAX__ 18446744073709551615UL
-// NVPTX32:#define __UINT_FAST64_TYPE__ long unsigned int
+// NVPTX32:#define __UINT_FAST64_TYPE__ long long unsigned int
 // NVPTX32:#define __UINT_FAST8_MAX__ 255
 // NVPTX32:#define __UINT_FAST8_TYPE__ unsigned char
 // NVPTX32:#define __UINT_LEAST16_MAX__ 65535
@@ -4836,7 +4836,7 @@
 // NVPTX32:#define __UINT_LEAST32_MAX__ 4294967295U
 // NVPTX32:#define __UINT_LEAST32_TYPE__ unsigned int
 // NVPTX32:#define __UINT_LEAST64_MAX__ 18446744073709551615UL
-// NVPTX32:#define __UINT_LEAST64_TYPE__ long unsigned int
+// NVPTX32:#define __UINT_LEAST64_TYPE__ long long unsigned int
 // NVPTX32:#define __UINT_LEAST8_MAX__ 255
 // NVPTX32:#define __UINT_LEAST8_TYPE__ unsigned char
 // NVPTX32:#define __USER_LABEL_PREFIX__ _