[CUDA] 32-bit NVPTX: long should be 32bit

[tra]

Summary

Currently 32-bit NVPTX happens to have 64-bit long which does not match host's long in 32-bit mode.

cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def

 
 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_clock64, "LLi", "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")
 

cfe/trunk/lib/Basic/Targets.cpp

void NVPTXTargetInfo::getGCCRegNames(const char *const *&Names,
                                      unsigned &NumNames) const {
   Names = GCCRegNames;
   NumNames = llvm::array_lengthof(GCCRegNames);
 }
 
 class NVPTX32TargetInfo : public NVPTXTargetInfo {
 public:
   NVPTX32TargetInfo(const llvm::Triple &Triple) : NVPTXTargetInfo(Triple) {
+    LongWidth = LongAlign = 32;
     PointerWidth = PointerAlign = 32;
     SizeType = TargetInfo::UnsignedInt;
     PtrDiffType = TargetInfo::SignedInt;
     IntPtrType = TargetInfo::SignedInt;
     DataLayoutString = "e-p:32:32-i64:64-v16:16-v32:32-n16:32:64";
   }
 };
 

cfe/trunk/test/CodeGen/builtins-nvptx.c

// CHECK: call i32 @llvm.ptx.read.lanemask.gt()
   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;
 
 }
 
-__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() {
 
 // CHECK: call i32 @llvm.ptx.read.pm0()
 // CHECK: call i32 @llvm.ptx.read.pm1()
 // CHECK: call i32 @llvm.ptx.read.pm2()
 // CHECK: call i32 @llvm.ptx.read.pm3()

cfe/trunk/test/CodeGen/nvptx-inlineasm-ptx.c

 // RUN: %clang_cc1 -triple nvptx-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s
 // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s
 
 void constraints() {
   char           c;
   unsigned char  uc;
   short          s;
   unsigned short us;
   int            i;
   unsigned int   ui;
-  long           l;
-  unsigned long  ul;
+  long long      ll;
+  unsigned long long ull;
   float          f;
   double         d;
 
   // CHECK: i8 asm sideeffect "mov.b8 $0, $1;", "=c,c"
   asm volatile ("mov.b8 %0, %1;" : "=c"(c) : "c"(c));
   // CHECK: i8 asm sideeffect "mov.b8 $0, $1;", "=c,c"
   asm volatile ("mov.b8 %0, %1;" : "=c"(uc) : "c"(uc));
 
   // CHECK: i16 asm sideeffect "mov.b16 $0, $1;", "=h,h"
   asm volatile ("mov.b16 %0, %1;" : "=h"(s) : "h"(s));
   // CHECK: i16 asm sideeffect "mov.b16 $0, $1;", "=h,h"
   asm volatile ("mov.b16 %0, %1;" : "=h"(us) : "h"(us));
 
   // CHECK: i32 asm sideeffect "mov.b32 $0, $1;", "=r,r"
   asm volatile ("mov.b32 %0, %1;" : "=r"(i) : "r"(i));
   // CHECK: i32 asm sideeffect "mov.b32 $0, $1;", "=r,r"
   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));
   // CHECK: double asm sideeffect "mov.b64 $0, $1;", "=d,d"
   asm volatile ("mov.b64 %0, %1;" : "=d"(d) : "d"(d));
 }

cfe/trunk/test/Preprocessor/init.c

 // NVPTX32:#define __INT_FAST16_FMTd__ "hd"
 // NVPTX32:#define __INT_FAST16_FMTi__ "hi"
 // NVPTX32:#define __INT_FAST16_MAX__ 32767
 // NVPTX32:#define __INT_FAST16_TYPE__ short
 // NVPTX32:#define __INT_FAST32_FMTd__ "d"
 // 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
 // NVPTX32:#define __INT_FAST8_TYPE__ signed char
 // NVPTX32:#define __INT_LEAST16_FMTd__ "hd"
 // NVPTX32:#define __INT_LEAST16_FMTi__ "hi"
 // NVPTX32:#define __INT_LEAST16_MAX__ 32767
 // NVPTX32:#define __INT_LEAST16_TYPE__ short
 // NVPTX32:#define __INT_LEAST32_FMTd__ "d"
 // 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
 // NVPTX32:#define __INT_LEAST8_TYPE__ signed char
 // NVPTX32:#define __INT_MAX__ 2147483647
 // NVPTX32:#define __LDBL_DENORM_MIN__ 4.9406564584124654e-324L
 // NVPTX32:#define __LDBL_DIG__ 15
 // NVPTX32:#define __LDBL_EPSILON__ 2.2204460492503131e-16L
 // NVPTX32:#define __LDBL_HAS_DENORM__ 1
 // NVPTX32:#define __LDBL_HAS_INFINITY__ 1
 // NVPTX32:#define __LDBL_HAS_QUIET_NAN__ 1
 // NVPTX32:#define __LDBL_MANT_DIG__ 53
 // NVPTX32:#define __LDBL_MAX_10_EXP__ 308
 // NVPTX32:#define __LDBL_MAX_EXP__ 1024
 // NVPTX32:#define __LDBL_MAX__ 1.7976931348623157e+308L
 // NVPTX32:#define __LDBL_MIN_10_EXP__ (-307)
 // NVPTX32:#define __LDBL_MIN_EXP__ (-1021)
 // 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
 // NVPTX32:#define __PRAGMA_REDEFINE_EXTNAME 1
 // NVPTX32:#define __PTRDIFF_TYPE__ int
 // NVPTX32:#define __PTRDIFF_WIDTH__ 32
 // NVPTX32:#define __PTX__ 1
 // NVPTX32:#define __SCHAR_MAX__ 127
 // NVPTX32:#define __SHRT_MAX__ 32767
 // NVPTX32:#define __SIG_ATOMIC_MAX__ 2147483647
 // NVPTX32:#define __SIG_ATOMIC_WIDTH__ 32
 // NVPTX32:#define __SIZEOF_DOUBLE__ 8
 // NVPTX32:#define __SIZEOF_FLOAT__ 4
 // 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
 // NVPTX32:#define __SIZEOF_SIZE_T__ 4
 // NVPTX32:#define __SIZEOF_WCHAR_T__ 4
 // NVPTX32:#define __SIZEOF_WINT_T__ 4
 // NVPTX32:#define __SIZE_MAX__ 4294967295U
 // NVPTX32:#define __SIZE_TYPE__ unsigned int

 // NVPTX32:#define __UINTPTR_MAX__ 4294967295U
 // NVPTX32:#define __UINTPTR_TYPE__ unsigned int
 // NVPTX32:#define __UINTPTR_WIDTH__ 32
 // NVPTX32:#define __UINT_FAST16_MAX__ 65535
 // NVPTX32:#define __UINT_FAST16_TYPE__ unsigned short
 // 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
 // NVPTX32:#define __UINT_LEAST16_TYPE__ unsigned short
 // 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__ _
 // NVPTX32:#define __WCHAR_MAX__ 2147483647
 // NVPTX32:#define __WCHAR_TYPE__ int
 // NVPTX32:#define __WCHAR_WIDTH__ 32
 // NVPTX32:#define __WINT_TYPE__ int
 // NVPTX32:#define __WINT_WIDTH__ 32