Vendor import of llvm trunk r154661:

http://llvm.org/svn/llvm-project/llvm/trunk@r154661

9 files changed, 83 insertions, 49 deletions

diff --git a/test/CodeGen/PTX/cvt.ll b/test/CodeGen/PTX/cvt.ll
index a643d25..f55070a 100644
--- a/test/CodeGen/PTX/cvt.ll
+++ b/test/CodeGen/PTX/cvt.ll
@@ -172,9 +172,9 @@ define ptx_device i64 @cvt_i64_f64(double %x) {
 ; f32
 
 define ptx_device float @cvt_f32_preds(i1 %x) {
-; CHECK: mov.b32 %f0, 1065353216;
-; CHECK: mov.b32 %f1, 0;
-; CHECK: selp.f32 %ret{{[0-9]+}}, %f0, %f1, %p{{[0-9]+}};
+; CHECK: mov.b32 %f0, 0;
+; CHECK: mov.b32 %f1, 1065353216;
+; CHECK: selp.f32 %ret{{[0-9]+}}, %f1, %f0, %p{{[0-9]+}};
 ; CHECK: ret;
 	%a = uitofp i1 %x to float
 	ret float %a
@@ -232,9 +232,9 @@ define ptx_device float @cvt_f32_s64(i64 %x) {
 ; f64
 
 define ptx_device double @cvt_f64_preds(i1 %x) {
-; CHECK: mov.b64 %fd0, 4575657221408423936;
-; CHECK: mov.b64 %fd1, 0;
-; CHECK: selp.f64 %ret{{[0-9]+}}, %fd0, %fd1, %p{{[0-9]+}};
+; CHECK: mov.b64 %fd0, 0;
+; CHECK: mov.b64 %fd1, 4575657221408423936;
+; CHECK: selp.f64 %ret{{[0-9]+}}, %fd1, %fd0, %p{{[0-9]+}};
 ; CHECK: ret;
 	%a = uitofp i1 %x to double
 	ret double %a
diff --git a/test/CodeGen/PTX/dg.exp b/test/CodeGen/PTX/dg.exp
deleted file mode 100644
index 2c304b5..0000000
--- a/test/CodeGen/PTX/dg.exp
+++ /dev/null
@@ -1,5 +0,0 @@
-load_lib llvm.exp
-
-if { [llvm_supports_target PTX] } {
-  RunLLVMTests [lsort [glob -nocomplain $srcdir/$subdir/*.{ll,c,cpp}]]
-}
diff --git a/test/CodeGen/PTX/ld.ll b/test/CodeGen/PTX/ld.ll
index 81fd33a..e55820d 100644
--- a/test/CodeGen/PTX/ld.ll
+++ b/test/CodeGen/PTX/ld.ll
@@ -1,48 +1,48 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
-;CHECK: .extern .global .b8 array_i16[20];
+;CHECK: .extern .global .b16 array_i16[10];
 @array_i16 = external global [10 x i16]
 
-;CHECK: .extern .const .b8 array_constant_i16[20];
+;CHECK: .extern .const .b16 array_constant_i16[10];
 @array_constant_i16 = external addrspace(1) constant [10 x i16]
 
-;CHECK: .extern .shared .b8 array_shared_i16[20];
+;CHECK: .extern .shared .b16 array_shared_i16[10];
 @array_shared_i16 = external addrspace(4) global [10 x i16]
 
-;CHECK: .extern .global .b8 array_i32[40];
+;CHECK: .extern .global .b32 array_i32[10];
 @array_i32 = external global [10 x i32]
 
-;CHECK: .extern .const .b8 array_constant_i32[40];
+;CHECK: .extern .const .b32 array_constant_i32[10];
 @array_constant_i32 = external addrspace(1) constant [10 x i32]
 
-;CHECK: .extern .shared .b8 array_shared_i32[40];
+;CHECK: .extern .shared .b32 array_shared_i32[10];
 @array_shared_i32 = external addrspace(4) global [10 x i32]
 
-;CHECK: .extern .global .b8 array_i64[80];
+;CHECK: .extern .global .b64 array_i64[10];
 @array_i64 = external global [10 x i64]
 
-;CHECK: .extern .const .b8 array_constant_i64[80];
+;CHECK: .extern .const .b64 array_constant_i64[10];
 @array_constant_i64 = external addrspace(1) constant [10 x i64]
 
-;CHECK: .extern .shared .b8 array_shared_i64[80];
+;CHECK: .extern .shared .b64 array_shared_i64[10];
 @array_shared_i64 = external addrspace(4) global [10 x i64]
 
-;CHECK: .extern .global .b8 array_float[40];
+;CHECK: .extern .global .b32 array_float[10];
 @array_float = external global [10 x float]
 
-;CHECK: .extern .const .b8 array_constant_float[40];
+;CHECK: .extern .const .b32 array_constant_float[10];
 @array_constant_float = external addrspace(1) constant [10 x float]
 
-;CHECK: .extern .shared .b8 array_shared_float[40];
+;CHECK: .extern .shared .b32 array_shared_float[10];
 @array_shared_float = external addrspace(4) global [10 x float]
 
-;CHECK: .extern .global .b8 array_double[80];
+;CHECK: .extern .global .b64 array_double[10];
 @array_double = external global [10 x double]
 
-;CHECK: .extern .const .b8 array_constant_double[80];
+;CHECK: .extern .const .b64 array_constant_double[10];
 @array_constant_double = external addrspace(1) constant [10 x double]
 
-;CHECK: .extern .shared .b8 array_shared_double[80];
+;CHECK: .extern .shared .b64 array_shared_double[10];
 @array_shared_double = external addrspace(4) global [10 x double]
 
 
diff --git a/test/CodeGen/PTX/lit.local.cfg b/test/CodeGen/PTX/lit.local.cfg
new file mode 100644
index 0000000..e748f7f
--- /dev/null
+++ b/test/CodeGen/PTX/lit.local.cfg
@@ -0,0 +1,6 @@
+config.suffixes = ['.ll', '.c', '.cpp']
+
+targets = set(config.root.targets_to_build.split())
+if not 'PTX' in targets:
+    config.unsupported = True
+
diff --git a/test/CodeGen/PTX/mad-disabling.ll b/test/CodeGen/PTX/mad-disabling.ll
index ad7b341..603c3ba 100644
--- a/test/CodeGen/PTX/mad-disabling.ll
+++ b/test/CodeGen/PTX/mad-disabling.ll
@@ -1,8 +1,13 @@
-; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20 | grep "mad"
-; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20,+no-fma | grep -v "mad"
+; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20 | FileCheck %s -check-prefix=FMA
+; RUN: llc < %s -march=ptx32 -mattr=+ptx20,+sm20,+no-fma | FileCheck %s -check-prefix=MUL
+; RUN: llc < %s -march=ptx64 -mattr=+ptx20,+sm20 | FileCheck %s -check-prefix=FMA
+; RUN: llc < %s -march=ptx64 -mattr=+ptx20,+sm20,+no-fma | FileCheck %s -check-prefix=MUL
 
 define ptx_device float @test_mul_add_f(float %x, float %y, float %z) {
 entry:
+; FMA: mad.rn.f32
+; MUL: mul.rn.f32
+; MUL: add.rn.f32
   %a = fmul float %x, %y
   %b = fadd float %a, %z
   ret float %b
@@ -10,6 +15,9 @@ entry:
 
 define ptx_device double @test_mul_add_d(double %x, double %y, double %z) {
 entry:
+; FMA: mad.rn.f64
+; MUL: mul.rn.f64
+; MUL: add.rn.f64
   %a = fmul double %x, %y
   %b = fadd double %a, %z
   ret double %b
diff --git a/test/CodeGen/PTX/mov.ll b/test/CodeGen/PTX/mov.ll
index 75555a7..9e501be 100644
--- a/test/CodeGen/PTX/mov.ll
+++ b/test/CodeGen/PTX/mov.ll
@@ -31,31 +31,31 @@ define ptx_device double @t1_f64() {
 }
 
 define ptx_device i16 @t2_u16(i16 %x) {
-; CHECK: mov.b16 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.b16 %ret{{[0-9]+}}, %arg{{[0-9]+}};
 ; CHECK: ret;
 	ret i16 %x
 }
 
 define ptx_device i32 @t2_u32(i32 %x) {
-; CHECK: mov.b32 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.b32 %ret{{[0-9]+}}, %arg{{[0-9]+}};
 ; CHECK: ret;
 	ret i32 %x
 }
 
 define ptx_device i64 @t2_u64(i64 %x) {
-; CHECK: mov.b64 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.b64 %ret{{[0-9]+}}, %arg{{[0-9]+}};
 ; CHECK: ret;
 	ret i64 %x
 }
 
 define ptx_device float @t3_f32(float %x) {
-; CHECK: mov.f32 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.f32 %ret{{[0-9]+}}, %arg{{[0-9]+}};
 ; CHECK: ret;
 	ret float %x
 }
 
 define ptx_device double @t3_f64(double %x) {
-; CHECK: mov.f64 %ret{{[0-9]+}}, %param{{[0-9]+}};
+; CHECK: mov.f64 %ret{{[0-9]+}}, %arg{{[0-9]+}};
 ; CHECK: ret;
 	ret double %x
 }
diff --git a/test/CodeGen/PTX/parameter-order.ll b/test/CodeGen/PTX/parameter-order.ll
index 09015da..377f173 100644
--- a/test/CodeGen/PTX/parameter-order.ll
+++ b/test/CodeGen/PTX/parameter-order.ll
@@ -1,6 +1,6 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
-; CHECK: .func (.reg .b32 %ret{{[0-9]+}}) test_parameter_order (.reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}}, .reg .b32 %param{{[0-9]+}})
+; CHECK: .func (.reg .b32 %ret{{[0-9]+}}) test_parameter_order (.reg .f32 %arg{{[0-9]+}}, .reg .b32 %arg{{[0-9]+}}, .reg .b32 %arg{{[0-9]+}}, .reg .f32 %arg{{[0-9]+}})
 define ptx_device i32 @test_parameter_order(float %a, i32 %b, i32 %c, float %d) {
 ; CHECK: sub.u32 %ret{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}
 	%result = sub i32 %b, %c
diff --git a/test/CodeGen/PTX/printf.ll b/test/CodeGen/PTX/printf.ll
new file mode 100644
index 0000000..f901b20
--- /dev/null
+++ b/test/CodeGen/PTX/printf.ll
@@ -0,0 +1,25 @@
+; RUN: llc < %s -march=ptx64 -mattr=+ptx20,+sm20 | FileCheck %s
+
+declare i32 @printf(i8*, ...)
+
+@str = private unnamed_addr constant [6 x i8] c"test\0A\00"
+
+define ptx_device void @t1_printf() {
+; CHECK: mov.u64 %rd{{[0-9]+}}, $L__str;
+; CHECK: call.uni	(__localparam_{{[0-9]+}}), vprintf, (__localparam_{{[0-9]+}}, __localparam_{{[0-9]+}});
+; CHECK: ret;
+    %1 = call i32 (i8*, ...)* @printf(i8* getelementptr inbounds ([6 x i8]* @str, i64 0, i64 0))
+	ret void
+}
+
+@str2 = private unnamed_addr constant [11 x i8] c"test = %f\0A\00"
+
+define ptx_device void @t2_printf() {
+; CHECK: .local .align 8 .b8 __local{{[0-9]+}}[{{[0-9]+}}];
+; CHECK: mov.u64 %rd{{[0-9]+}}, $L__str2;
+; CHECK: cvta.local.u64  %rd{{[0-9]+}}, __local{{[0-9+]}};
+; CHECK: call.uni	(__localparam_{{[0-9]+}}), vprintf, (__localparam_{{[0-9]+}}, __localparam_{{[0-9]+}});
+; CHECK: ret;
+  %1 = call i32 (i8*, ...)* @printf(i8* getelementptr inbounds ([11 x i8]* @str2, i64 0, i64 0), double 0x3FF3333340000000)
+  ret void
+}
diff --git a/test/CodeGen/PTX/st.ll b/test/CodeGen/PTX/st.ll
index 63ef58c..c794363 100644
--- a/test/CodeGen/PTX/st.ll
+++ b/test/CodeGen/PTX/st.ll
@@ -1,48 +1,48 @@
 ; RUN: llc < %s -march=ptx32 | FileCheck %s
 
-;CHECK: .extern .global .b8 array_i16[20];
+;CHECK: .extern .global .b16 array_i16[10];
 @array_i16 = external global [10 x i16]
 
-;CHECK: .extern .const .b8 array_constant_i16[20];
+;CHECK: .extern .const .b16 array_constant_i16[10];
 @array_constant_i16 = external addrspace(1) constant [10 x i16]
 
-;CHECK: .extern .shared .b8 array_shared_i16[20];
+;CHECK: .extern .shared .b16 array_shared_i16[10];
 @array_shared_i16 = external addrspace(4) global [10 x i16]
 
-;CHECK: .extern .global .b8 array_i32[40];
+;CHECK: .extern .global .b32 array_i32[10];
 @array_i32 = external global [10 x i32]
 
-;CHECK: .extern .const .b8 array_constant_i32[40];
+;CHECK: .extern .const .b32 array_constant_i32[10];
 @array_constant_i32 = external addrspace(1) constant [10 x i32]
 
-;CHECK: .extern .shared .b8 array_shared_i32[40];
+;CHECK: .extern .shared .b32 array_shared_i32[10];
 @array_shared_i32 = external addrspace(4) global [10 x i32]
 
-;CHECK: .extern .global .b8 array_i64[80];
+;CHECK: .extern .global .b64 array_i64[10];
 @array_i64 = external global [10 x i64]
 
-;CHECK: .extern .const .b8 array_constant_i64[80];
+;CHECK: .extern .const .b64 array_constant_i64[10];
 @array_constant_i64 = external addrspace(1) constant [10 x i64]
 
-;CHECK: .extern .shared .b8 array_shared_i64[80];
+;CHECK: .extern .shared .b64 array_shared_i64[10];
 @array_shared_i64 = external addrspace(4) global [10 x i64]
 
-;CHECK: .extern .global .b8 array_float[40];
+;CHECK: .extern .global .b32 array_float[10];
 @array_float = external global [10 x float]
 
-;CHECK: .extern .const .b8 array_constant_float[40];
+;CHECK: .extern .const .b32 array_constant_float[10];
 @array_constant_float = external addrspace(1) constant [10 x float]
 
-;CHECK: .extern .shared .b8 array_shared_float[40];
+;CHECK: .extern .shared .b32 array_shared_float[10];
 @array_shared_float = external addrspace(4) global [10 x float]
 
-;CHECK: .extern .global .b8 array_double[80];
+;CHECK: .extern .global .b64 array_double[10];
 @array_double = external global [10 x double]
 
-;CHECK: .extern .const .b8 array_constant_double[80];
+;CHECK: .extern .const .b64 array_constant_double[10];
 @array_constant_double = external addrspace(1) constant [10 x double]
 
-;CHECK: .extern .shared .b8 array_shared_double[80];
+;CHECK: .extern .shared .b64 array_shared_double[10];
 @array_shared_double = external addrspace(4) global [10 x double]