r248753 - [CUDA] 32-bit NVPTX should have 32-bit long type.
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 28 15:54:08 PDT 2015
Author: tra
Date: Mon Sep 28 17:54:08 2015
New Revision: 248753
URL: http://llvm.org/viewvc/llvm-project?rev=248753&view=rev
Log:
[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
Modified:
cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def
cfe/trunk/lib/Basic/Targets.cpp
cfe/trunk/test/CodeGen/builtins-nvptx.c
cfe/trunk/test/CodeGen/nvptx-inlineasm-ptx.c
cfe/trunk/test/Preprocessor/init.c
Modified: cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def?rev=248753&r1=248752&r2=248753&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def Mon Sep 28 17:54:08 2015
@@ -50,7 +50,7 @@ BUILTIN(__builtin_ptx_read_lanemask_ge,
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")
Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=248753&r1=248752&r2=248753&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Mon Sep 28 17:54:08 2015
@@ -1716,6 +1716,7 @@ void NVPTXTargetInfo::getGCCRegNames(con
class NVPTX32TargetInfo : public NVPTXTargetInfo {
public:
NVPTX32TargetInfo(const llvm::Triple &Triple) : NVPTXTargetInfo(Triple) {
+ LongWidth = LongAlign = 32;
PointerWidth = PointerAlign = 32;
SizeType = TargetInfo::UnsignedInt;
PtrDiffType = TargetInfo::SignedInt;
Modified: cfe/trunk/test/CodeGen/builtins-nvptx.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-nvptx.c?rev=248753&r1=248752&r2=248753&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/builtins-nvptx.c (original)
+++ cfe/trunk/test/CodeGen/builtins-nvptx.c Mon Sep 28 17:54:08 2015
@@ -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() {
Modified: cfe/trunk/test/CodeGen/nvptx-inlineasm-ptx.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/nvptx-inlineasm-ptx.c?rev=248753&r1=248752&r2=248753&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/nvptx-inlineasm-ptx.c (original)
+++ cfe/trunk/test/CodeGen/nvptx-inlineasm-ptx.c Mon Sep 28 17:54:08 2015
@@ -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));
Modified: cfe/trunk/test/Preprocessor/init.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Preprocessor/init.c?rev=248753&r1=248752&r2=248753&view=diff
==============================================================================
--- cfe/trunk/test/Preprocessor/init.c (original)
+++ cfe/trunk/test/Preprocessor/init.c Mon Sep 28 17:54:08 2015
@@ -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__ _
More information about the cfe-commits
mailing list