dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[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
This commit is contained in:
Artem Belevich 2015-09-28 22:54:08 +00:00
parent bfa22743ff
commit 757da43210
5 changed files with 19 additions and 19 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
include/clang/Basic/BuiltinsNVPTX.def
View File

@ -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")

1
lib/Basic/Targets.cpp
View File

@ -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;

7
test/CodeGen/builtins-nvptx.c
View File

@ -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() {

8
test/CodeGen/nvptx-inlineasm-ptx.c
View File

@ -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));

20
test/Preprocessor/init.c
View File

@ -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__ _