11 files changed, 1057 insertions, 27 deletions

diff --git a/test/CodeGen/NVPTX/branch-fold.ll b/test/CodeGen/NVPTX/branch-fold.ll
new file mode 100644
index 000000000000..2b9cd0a35d92
--- /dev/null
+++ b/test/CodeGen/NVPTX/branch-fold.ll
@@ -0,0 +1,40 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -disable-cgp | FileCheck %s
+; Disable CGP which also folds branches, so that only BranchFolding is under
+; the spotlight.
+
+target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+define void @foo(i32 %x, float* %output) {
+; CHECK-LABEL: .visible .func foo(
+; CHECK-NOT: bra.uni
+; CHECK-NOT: LBB0_
+  %1 = icmp eq i32 %x, 1
+  br i1 %1, label %then, label %else
+
+then:
+  br label %merge
+
+else:
+  br label %merge
+
+merge:
+  store float 2.0, float* %output
+  ret void
+}
+
+; PR24299. no crash
+define ptx_kernel void @hoge() #0 {
+; CHECK-LABEL: .visible .entry hoge(
+bb:
+  br i1 undef, label %bb1, label %bb4
+
+bb1:                                              ; preds = %bb1, %bb
+  %tmp = phi i64 [ %tmp2, %bb1 ], [ undef, %bb ]
+  %tmp2 = add nsw i64 %tmp, 1
+  %tmp3 = icmp sle i64 %tmp, 0
+  br i1 %tmp3, label %bb1, label %bb4
+
+bb4:                                              ; preds = %bb4, %bb1, %bb
+  br label %bb4
+}
diff --git a/test/CodeGen/NVPTX/bypass-div.ll b/test/CodeGen/NVPTX/bypass-div.ll
new file mode 100644
index 000000000000..bd98c9a5b0b1
--- /dev/null
+++ b/test/CodeGen/NVPTX/bypass-div.ll
@@ -0,0 +1,80 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+
+; 64-bit divides and rems should be split into a fast and slow path where
+; the fast path uses a 32-bit operation.
+
+define void @sdiv64(i64 %a, i64 %b, i64* %retptr) {
+; CHECK-LABEL: sdiv64(
+; CHECK:        div.s64
+; CHECK:        div.u32
+; CHECK:        ret
+  %d = sdiv i64 %a, %b
+  store i64 %d, i64* %retptr
+  ret void
+}
+
+define void @udiv64(i64 %a, i64 %b, i64* %retptr) {
+; CHECK-LABEL: udiv64(
+; CHECK:        div.u64
+; CHECK:        div.u32
+; CHECK:        ret
+  %d = udiv i64 %a, %b
+  store i64 %d, i64* %retptr
+  ret void
+}
+
+define void @srem64(i64 %a, i64 %b, i64* %retptr) {
+; CHECK-LABEL: srem64(
+; CHECK:        rem.s64
+; CHECK:        rem.u32
+; CHECK:        ret
+  %d = srem i64 %a, %b
+  store i64 %d, i64* %retptr
+  ret void
+}
+
+define void @urem64(i64 %a, i64 %b, i64* %retptr) {
+; CHECK-LABEL: urem64(
+; CHECK:        rem.u64
+; CHECK:        rem.u32
+; CHECK:        ret
+  %d = urem i64 %a, %b
+  store i64 %d, i64* %retptr
+  ret void
+}
+
+define void @sdiv32(i32 %a, i32 %b, i32* %retptr) {
+; CHECK-LABEL: sdiv32(
+; CHECK: div.s32
+; CHECK-NOT: div.
+  %d = sdiv i32 %a, %b
+  store i32 %d, i32* %retptr
+  ret void
+}
+
+define void @udiv32(i32 %a, i32 %b, i32* %retptr) {
+; CHECK-LABEL: udiv32(
+; CHECK: div.u32
+; CHECK-NOT: div.
+  %d = udiv i32 %a, %b
+  store i32 %d, i32* %retptr
+  ret void
+}
+
+define void @srem32(i32 %a, i32 %b, i32* %retptr) {
+; CHECK-LABEL: srem32(
+; CHECK: rem.s32
+; CHECK-NOT: rem.
+  %d = srem i32 %a, %b
+  store i32 %d, i32* %retptr
+  ret void
+}
+
+define void @urem32(i32 %a, i32 %b, i32* %retptr) {
+; CHECK-LABEL: urem32(
+; CHECK: rem.u32
+; CHECK-NOT: rem.
+  %d = urem i32 %a, %b
+  store i32 %d, i32* %retptr
+  ret void
+}
diff --git a/test/CodeGen/NVPTX/combine-min-max.ll b/test/CodeGen/NVPTX/combine-min-max.ll
new file mode 100644
index 000000000000..64bb7a37ffd2
--- /dev/null
+++ b/test/CodeGen/NVPTX/combine-min-max.ll
@@ -0,0 +1,307 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -O2 | FileCheck %s
+
+; *************************************
+; * Cases with no min/max
+
+define i32 @ab_eq_i32(i32 %a, i32 %b) {
+; LABEL: @ab_slt_i32
+; CHECK-NOT: min
+; CHECK-NOT: max
+  %cmp = icmp eq i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+define i64 @ba_ne_i64(i64 %a, i64 %b) {
+; LABEL: @ab_ne_i64
+; CHECK-NOT: min
+; CHECK-NOT: max
+  %cmp = icmp ne i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+; PTX does have e.g. max.s16, but at least as of Kepler (sm_3x) that
+; gets compiled to SASS that converts the 16 bit parameters to 32 bit
+; before using a 32 bit instruction. That is probably not a win and
+; NVCC 7.5 does not emit 16 bit min/max either, presumably for that
+; reason.
+define i16 @ab_ugt_i16(i16 %a, i16 %b) {
+; LABEL: @ab_ugt_i16
+; CHECK-NOT: min
+; CHECK-NOT: max
+  %cmp = icmp ugt i16 %a, %b
+  %sel = select i1 %cmp, i16 %a, i16 %b
+  ret i16 %sel
+}
+
+
+; *************************************
+; * All variations with i32
+
+; *** ab, unsigned, i32
+define i32 @ab_ugt_i32(i32 %a, i32 %b) {
+; LABEL: @ab_ugt_i32
+; CHECK: max.u32
+  %cmp = icmp ugt i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+define i32 @ab_uge_i32(i32 %a, i32 %b) {
+; LABEL: @ab_uge_i32
+; CHECK: max.u32
+  %cmp = icmp uge i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+define i32 @ab_ult_i32(i32 %a, i32 %b) {
+; LABEL: @ab_ult_i32
+; CHECK: min.u32
+  %cmp = icmp ult i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+define i32 @ab_ule_i32(i32 %a, i32 %b) {
+; LABEL: @ab_ule_i32
+; CHECK: min.u32
+  %cmp = icmp ule i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+; *** ab, signed, i32
+define i32 @ab_sgt_i32(i32 %a, i32 %b) {
+; LABEL: @ab_ugt_i32
+; CHECK: max.s32
+  %cmp = icmp sgt i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+define i32 @ab_sge_i32(i32 %a, i32 %b) {
+; LABEL: @ab_sge_i32
+; CHECK: max.s32
+  %cmp = icmp sge i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+define i32 @ab_slt_i32(i32 %a, i32 %b) {
+; LABEL: @ab_slt_i32
+; CHECK: min.s32
+  %cmp = icmp slt i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+define i32 @ab_sle_i32(i32 %a, i32 %b) {
+; LABEL: @ab_sle_i32
+; CHECK: min.s32
+  %cmp = icmp sle i32 %a, %b
+  %sel = select i1 %cmp, i32 %a, i32 %b
+  ret i32 %sel
+}
+
+; *** ba, unsigned, i32
+define i32 @ba_ugt_i32(i32 %a, i32 %b) {
+; LABEL: @ba_ugt_i32
+; CHECK: min.u32
+  %cmp = icmp ugt i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+define i32 @ba_uge_i32(i32 %a, i32 %b) {
+; LABEL: @ba_uge_i32
+; CHECK: min.u32
+  %cmp = icmp uge i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+define i32 @ba_ult_i32(i32 %a, i32 %b) {
+; LABEL: @ba_ult_i32
+; CHECK: max.u32
+  %cmp = icmp ult i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+define i32 @ba_ule_i32(i32 %a, i32 %b) {
+; LABEL: @ba_ule_i32
+; CHECK: max.u32
+  %cmp = icmp ule i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+; *** ba, signed, i32
+define i32 @ba_sgt_i32(i32 %a, i32 %b) {
+; LBAEL: @ba_ugt_i32
+; CHECK: min.s32
+  %cmp = icmp sgt i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+define i32 @ba_sge_i32(i32 %a, i32 %b) {
+; LABEL: @ba_sge_i32
+; CHECK: min.s32
+  %cmp = icmp sge i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+define i32 @ba_slt_i32(i32 %a, i32 %b) {
+; LABEL: @ba_slt_i32
+; CHECK: max.s32
+  %cmp = icmp slt i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+define i32 @ba_sle_i32(i32 %a, i32 %b) {
+; LABEL: @ba_sle_i32
+; CHECK: max.s32
+  %cmp = icmp sle i32 %a, %b
+  %sel = select i1 %cmp, i32 %b, i32 %a
+  ret i32 %sel
+}
+
+; *************************************
+; * All variations with i64
+
+; *** ab, unsigned, i64
+define i64 @ab_ugt_i64(i64 %a, i64 %b) {
+; LABEL: @ab_ugt_i64
+; CHECK: max.u64
+  %cmp = icmp ugt i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+define i64 @ab_uge_i64(i64 %a, i64 %b) {
+; LABEL: @ab_uge_i64
+; CHECK: max.u64
+  %cmp = icmp uge i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+define i64 @ab_ult_i64(i64 %a, i64 %b) {
+; LABEL: @ab_ult_i64
+; CHECK: min.u64
+  %cmp = icmp ult i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+define i64 @ab_ule_i64(i64 %a, i64 %b) {
+; LABEL: @ab_ule_i64
+; CHECK: min.u64
+  %cmp = icmp ule i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+; *** ab, signed, i64
+define i64 @ab_sgt_i64(i64 %a, i64 %b) {
+; LABEL: @ab_ugt_i64
+; CHECK: max.s64
+  %cmp = icmp sgt i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+define i64 @ab_sge_i64(i64 %a, i64 %b) {
+; LABEL: @ab_sge_i64
+; CHECK: max.s64
+  %cmp = icmp sge i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+define i64 @ab_slt_i64(i64 %a, i64 %b) {
+; LABEL: @ab_slt_i64
+; CHECK: min.s64
+  %cmp = icmp slt i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+define i64 @ab_sle_i64(i64 %a, i64 %b) {
+; LABEL: @ab_sle_i64
+; CHECK: min.s64
+  %cmp = icmp sle i64 %a, %b
+  %sel = select i1 %cmp, i64 %a, i64 %b
+  ret i64 %sel
+}
+
+; *** ba, unsigned, i64
+define i64 @ba_ugt_i64(i64 %a, i64 %b) {
+; LABEL: @ba_ugt_i64
+; CHECK: min.u64
+  %cmp = icmp ugt i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+define i64 @ba_uge_i64(i64 %a, i64 %b) {
+; LABEL: @ba_uge_i64
+; CHECK: min.u64
+  %cmp = icmp uge i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+define i64 @ba_ult_i64(i64 %a, i64 %b) {
+; LABEL: @ba_ult_i64
+; CHECK: max.u64
+  %cmp = icmp ult i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+define i64 @ba_ule_i64(i64 %a, i64 %b) {
+; LABEL: @ba_ule_i64
+; CHECK: max.u64
+  %cmp = icmp ule i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+; *** ba, signed, i64
+define i64 @ba_sgt_i64(i64 %a, i64 %b) {
+; LBAEL: @ba_ugt_i64
+; CHECK: min.s64
+  %cmp = icmp sgt i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+define i64 @ba_sge_i64(i64 %a, i64 %b) {
+; LABEL: @ba_sge_i64
+; CHECK: min.s64
+  %cmp = icmp sge i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+define i64 @ba_slt_i64(i64 %a, i64 %b) {
+; LABEL: @ba_slt_i64
+; CHECK: max.s64
+  %cmp = icmp slt i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
+
+define i64 @ba_sle_i64(i64 %a, i64 %b) {
+; LABEL: @ba_sle_i64
+; CHECK: max.s64
+  %cmp = icmp sle i64 %a, %b
+  %sel = select i1 %cmp, i64 %b, i64 %a
+  ret i64 %sel
+}
diff --git a/test/CodeGen/NVPTX/fma-assoc.ll b/test/CodeGen/NVPTX/fma-assoc.ll
index fc04c61dd691..80a08a86316c 100644
--- a/test/CodeGen/NVPTX/fma-assoc.ll
+++ b/test/CodeGen/NVPTX/fma-assoc.ll
@@ -23,3 +23,16 @@ define ptx_device double @t1_f64(double %x, double %y, double %z,
   %d = fadd double %c, %z
   ret double %d
 }
+
+define double @two_choices(double %val1, double %val2) {
+; CHECK-LABEL: two_choices(
+; CHECK: mul.f64
+; CHECK-NOT: mul.f64
+; CHECK: fma.rn.f64
+  %1 = fmul double %val1, %val2
+  %2 = fmul double %1, %1
+  %3 = fadd double %1, %2
+
+  ret double %3
+}
+
diff --git a/test/CodeGen/NVPTX/global-addrspace.ll b/test/CodeGen/NVPTX/global-addrspace.ll
new file mode 100644
index 000000000000..4da14c7ff4fe
--- /dev/null
+++ b/test/CodeGen/NVPTX/global-addrspace.ll
@@ -0,0 +1,12 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
+
+; PTX32: .visible .global .align 4 .u32 i;
+; PTX32: .visible .const .align 4 .u32 j;
+; PTX32: .visible .shared .align 4 .u32 k;
+; PTX64: .visible .global .align 4 .u32 i;
+; PTX64: .visible .const .align 4 .u32 j;
+; PTX64: .visible .shared .align 4 .u32 k;
+@i = addrspace(1) externally_initialized global i32 0, align 4
+@j = addrspace(4) externally_initialized global i32 0, align 4
+@k = addrspace(3) global i32 undef, align 4
diff --git a/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll b/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
new file mode 100644
index 000000000000..d93499b47f59
--- /dev/null
+++ b/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
@@ -0,0 +1,264 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck -check-prefix=SM20 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck -check-prefix=SM35 %s
+
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
+target triple = "nvptx64-unknown-unknown"
+
+; SM20-LABEL: .visible .entry foo1(
+; SM20: ld.global.f32
+; SM35-LABEL: .visible .entry foo1(
+; SM35: ld.global.nc.f32
+define void @foo1(float * noalias readonly %from, float * %to) {
+  %1 = load float, float * %from
+  store float %1, float * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo2(
+; SM20: ld.global.f64
+; SM35-LABEL: .visible .entry foo2(
+; SM35: ld.global.nc.f64
+define void @foo2(double * noalias readonly %from, double * %to) {
+  %1 = load double, double * %from
+  store double %1, double * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo3(
+; SM20: ld.global.u16
+; SM35-LABEL: .visible .entry foo3(
+; SM35: ld.global.nc.u16
+define void @foo3(i16 * noalias readonly %from, i16 * %to) {
+  %1 = load i16, i16 * %from
+  store i16 %1, i16 * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo4(
+; SM20: ld.global.u32
+; SM35-LABEL: .visible .entry foo4(
+; SM35: ld.global.nc.u32
+define void @foo4(i32 * noalias readonly %from, i32 * %to) {
+  %1 = load i32, i32 * %from
+  store i32 %1, i32 * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo5(
+; SM20: ld.global.u64
+; SM35-LABEL: .visible .entry foo5(
+; SM35: ld.global.nc.u64
+define void @foo5(i64 * noalias readonly %from, i64 * %to) {
+  %1 = load i64, i64 * %from
+  store i64 %1, i64 * %to
+  ret void
+}
+
+; i128 is non standard integer in nvptx64
+; SM20-LABEL: .visible .entry foo6(
+; SM20: ld.global.u64
+; SM20: ld.global.u64
+; SM35-LABEL: .visible .entry foo6(
+; SM35: ld.global.nc.u64
+; SM35: ld.global.nc.u64
+define void @foo6(i128 * noalias readonly %from, i128 * %to) {
+  %1 = load i128, i128 * %from
+  store i128 %1, i128 * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo7(
+; SM20: ld.global.v2.u8
+; SM35-LABEL: .visible .entry foo7(
+; SM35: ld.global.nc.v2.u8
+define void @foo7(<2 x i8> * noalias readonly %from, <2 x i8> * %to) {
+  %1 = load <2 x i8>, <2 x i8> * %from
+  store <2 x i8> %1, <2 x i8> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo8(
+; SM20: ld.global.v2.u16
+; SM35-LABEL: .visible .entry foo8(
+; SM35: ld.global.nc.v2.u16
+define void @foo8(<2 x i16> * noalias readonly %from, <2 x i16> * %to) {
+  %1 = load <2 x i16>, <2 x i16> * %from
+  store <2 x i16> %1, <2 x i16> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo9(
+; SM20: ld.global.v2.u32
+; SM35-LABEL: .visible .entry foo9(
+; SM35: ld.global.nc.v2.u32
+define void @foo9(<2 x i32> * noalias readonly %from, <2 x i32> * %to) {
+  %1 = load <2 x i32>, <2 x i32> * %from
+  store <2 x i32> %1, <2 x i32> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo10(
+; SM20: ld.global.v2.u64
+; SM35-LABEL: .visible .entry foo10(
+; SM35: ld.global.nc.v2.u64
+define void @foo10(<2 x i64> * noalias readonly %from, <2 x i64> * %to) {
+  %1 = load <2 x i64>, <2 x i64> * %from
+  store <2 x i64> %1, <2 x i64> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo11(
+; SM20: ld.global.v2.f32
+; SM35-LABEL: .visible .entry foo11(
+; SM35: ld.global.nc.v2.f32
+define void @foo11(<2 x float> * noalias readonly %from, <2 x float> * %to) {
+  %1 = load <2 x float>, <2 x float> * %from
+  store <2 x float> %1, <2 x float> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo12(
+; SM20: ld.global.v2.f64
+; SM35-LABEL: .visible .entry foo12(
+; SM35: ld.global.nc.v2.f64
+define void @foo12(<2 x double> * noalias readonly %from, <2 x double> * %to) {
+  %1 = load <2 x double>, <2 x double> * %from
+  store <2 x double> %1, <2 x double> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo13(
+; SM20: ld.global.v4.u8
+; SM35-LABEL: .visible .entry foo13(
+; SM35: ld.global.nc.v4.u8
+define void @foo13(<4 x i8> * noalias readonly %from, <4 x i8> * %to) {
+  %1 = load <4 x i8>, <4 x i8> * %from
+  store <4 x i8> %1, <4 x i8> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo14(
+; SM20: ld.global.v4.u16
+; SM35-LABEL: .visible .entry foo14(
+; SM35: ld.global.nc.v4.u16
+define void @foo14(<4 x i16> * noalias readonly %from, <4 x i16> * %to) {
+  %1 = load <4 x i16>, <4 x i16> * %from
+  store <4 x i16> %1, <4 x i16> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo15(
+; SM20: ld.global.v4.u32
+; SM35-LABEL: .visible .entry foo15(
+; SM35: ld.global.nc.v4.u32
+define void @foo15(<4 x i32> * noalias readonly %from, <4 x i32> * %to) {
+  %1 = load <4 x i32>, <4 x i32> * %from
+  store <4 x i32> %1, <4 x i32> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo16(
+; SM20: ld.global.v4.f32
+; SM35-LABEL: .visible .entry foo16(
+; SM35: ld.global.nc.v4.f32
+define void @foo16(<4 x float> * noalias readonly %from, <4 x float> * %to) {
+  %1 = load <4 x float>, <4 x float> * %from
+  store <4 x float> %1, <4 x float> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo17(
+; SM20: ld.global.v2.f64
+; SM20: ld.global.v2.f64
+; SM35-LABEL: .visible .entry foo17(
+; SM35: ld.global.nc.v2.f64
+; SM35: ld.global.nc.v2.f64
+define void @foo17(<4 x double> * noalias readonly %from, <4 x double> * %to) {
+  %1 = load <4 x double>, <4 x double> * %from
+  store <4 x double> %1, <4 x double> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo18(
+; SM20: ld.global.u64
+; SM35-LABEL: .visible .entry foo18(
+; SM35: ld.global.nc.u64
+define void @foo18(float ** noalias readonly %from, float ** %to) {
+  %1 = load float *, float ** %from
+  store float * %1, float ** %to
+  ret void
+}
+
+; Test that we can infer a cached load for a pointer induction variable.
+; SM20-LABEL: .visible .entry foo19(
+; SM20: ld.global.f32
+; SM35-LABEL: .visible .entry foo19(
+; SM35: ld.global.nc.f32
+define void @foo19(float * noalias readonly %from, float * %to, i32 %n) {
+entry:
+  br label %loop
+
+loop:
+  %i = phi i32 [ 0, %entry ], [ %nexti, %loop ]
+  %sum = phi float [ 0.0, %entry ], [ %nextsum, %loop ]
+  %ptr = getelementptr inbounds float, float * %from, i32 %i
+  %value = load float, float * %ptr, align 4
+  %nextsum = fadd float %value, %sum
+  %nexti = add nsw i32 %i, 1
+  %exitcond = icmp eq i32 %nexti, %n
+  br i1 %exitcond, label %exit, label %loop
+
+exit:
+  store float %nextsum, float * %to
+  ret void
+}
+
+; This test captures the case of a non-kernel function. In a
+; non-kernel function, without interprocedural analysis, we do not
+; know that the parameter is global. We also do not know that the
+; pointed-to memory is never written to (for the duration of the
+; kernel). For both reasons, we cannot use a cached load here.
+; SM20-LABEL: notkernel(
+; SM20: ld.f32
+; SM35-LABEL: notkernel(
+; SM35: ld.f32
+define void @notkernel(float * noalias readonly %from, float * %to) {
+  %1 = load float, float * %from
+  store float %1, float * %to
+  ret void
+}
+
+; As @notkernel, but with the parameter explicitly marked as global. We still
+; do not know that the parameter is never written to (for the duration of the
+; kernel). This case does not currently come up normally since we do not infer
+; that pointers are global interprocedurally as of 2015-08-05.
+; SM20-LABEL: notkernel2(
+; SM20: ld.global.f32
+; SM35-LABEL: notkernel2(
+; SM35: ld.global.f32
+define void @notkernel2(float addrspace(1) * noalias readonly %from, float * %to) {
+  %1 = load float, float addrspace(1) * %from
+  store float %1, float * %to
+  ret void
+}
+
+!nvvm.annotations = !{!1 ,!2 ,!3 ,!4 ,!5 ,!6, !7 ,!8 ,!9 ,!10 ,!11 ,!12, !13, !14, !15, !16, !17, !18, !19}
+!1 = !{void (float *, float *)* @foo1, !"kernel", i32 1}
+!2 = !{void (double *, double *)* @foo2, !"kernel", i32 1}
+!3 = !{void (i16 *, i16 *)* @foo3, !"kernel", i32 1}
+!4 = !{void (i32 *, i32 *)* @foo4, !"kernel", i32 1}
+!5 = !{void (i64 *, i64 *)* @foo5, !"kernel", i32 1}
+!6 = !{void (i128 *, i128 *)* @foo6, !"kernel", i32 1}
+!7 = !{void (<2 x i8> *, <2 x i8> *)* @foo7, !"kernel", i32 1}
+!8 = !{void (<2 x i16> *, <2 x i16> *)* @foo8, !"kernel", i32 1}
+!9 = !{void (<2 x i32> *, <2 x i32> *)* @foo9, !"kernel", i32 1}
+!10 = !{void (<2 x i64> *, <2 x i64> *)* @foo10, !"kernel", i32 1}
+!11 = !{void (<2 x float> *, <2 x float> *)* @foo11, !"kernel", i32 1}
+!12 = !{void (<2 x double> *, <2 x double> *)* @foo12, !"kernel", i32 1}
+!13 = !{void (<4 x i8> *, <4 x i8> *)* @foo13, !"kernel", i32 1}
+!14 = !{void (<4 x i16> *, <4 x i16> *)* @foo14, !"kernel", i32 1}
+!15 = !{void (<4 x i32> *, <4 x i32> *)* @foo15, !"kernel", i32 1}
+!16 = !{void (<4 x float> *, <4 x float> *)* @foo16, !"kernel", i32 1}
+!17 = !{void (<4 x double> *, <4 x double> *)* @foo17, !"kernel", i32 1}
+!18 = !{void (float **, float **)* @foo18, !"kernel", i32 1}
+!19 = !{void (float *, float *, i32)* @foo19, !"kernel", i32 1}
diff --git a/test/CodeGen/NVPTX/lower-aggr-copies.ll b/test/CodeGen/NVPTX/lower-aggr-copies.ll
index c3adfc4646cf..ef570982b808 100644
--- a/test/CodeGen/NVPTX/lower-aggr-copies.ll
+++ b/test/CodeGen/NVPTX/lower-aggr-copies.ll
@@ -1,35 +1,68 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix PTX
+; RUN: opt < %s -S -nvptx-lower-aggr-copies | FileCheck %s --check-prefix IR
 
 ; Verify that the NVPTXLowerAggrCopies pass works as expected - calls to
 ; llvm.mem* intrinsics get lowered to loops.
 
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target triple = "nvptx64-unknown-unknown"
+
 declare void @llvm.memcpy.p0i8.p0i8.i64(i8* nocapture, i8* nocapture readonly, i64, i32, i1) #1
+declare void @llvm.memmove.p0i8.p0i8.i64(i8* nocapture, i8* nocapture readonly, i64, i32, i1) #1
 declare void @llvm.memset.p0i8.i64(i8* nocapture, i8, i64, i32, i1) #1
 
 define i8* @memcpy_caller(i8* %dst, i8* %src, i64 %n) #0 {
 entry:
   tail call void @llvm.memcpy.p0i8.p0i8.i64(i8* %dst, i8* %src, i64 %n, i32 1, i1 false)
   ret i8* %dst
-; CHECK-LABEL: .visible .func (.param .b32 func_retval0) memcpy_caller
-; CHECK: LBB[[LABEL:[_0-9]+]]:
-; CHECK:      ld.u8 %rs[[REG:[0-9]+]]
-; CHECK:      st.u8 [%r{{[0-9]+}}], %rs[[REG]]
-; CHECK:      add.s64 %rd[[COUNTER:[0-9]+]], %rd[[COUNTER]], 1
-; CHECK-NEXT: setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
-; CHECK-NEXT: @%p[[PRED]] bra LBB[[LABEL]]
+
+; IR-LABEL:   @memcpy_caller
+; IR:         loadstoreloop:
+; IR:         [[LOADPTR:%[0-9]+]] = getelementptr inbounds i8, i8* %src, i64
+; IR-NEXT:    [[VAL:%[0-9]+]] = load i8, i8* [[LOADPTR]]
+; IR-NEXT:    [[STOREPTR:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64
+; IR-NEXT:    store i8 [[VAL]], i8* [[STOREPTR]]
+
+; PTX-LABEL:  .visible .func (.param .b64 func_retval0) memcpy_caller
+; PTX:        LBB[[LABEL:[_0-9]+]]:
+; PTX:        ld.u8 %rs[[REG:[0-9]+]]
+; PTX:        st.u8 [%rd{{[0-9]+}}], %rs[[REG]]
+; PTX:        add.s64 %rd[[COUNTER:[0-9]+]], %rd[[COUNTER]], 1
+; PTX-NEXT:   setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
+; PTX-NEXT:   @%p[[PRED]] bra LBB[[LABEL]]
 }
 
 define i8* @memcpy_volatile_caller(i8* %dst, i8* %src, i64 %n) #0 {
 entry:
   tail call void @llvm.memcpy.p0i8.p0i8.i64(i8* %dst, i8* %src, i64 %n, i32 1, i1 true)
   ret i8* %dst
-; CHECK-LABEL: .visible .func (.param .b32 func_retval0) memcpy_volatile_caller
-; CHECK: LBB[[LABEL:[_0-9]+]]:
-; CHECK:      ld.volatile.u8 %rs[[REG:[0-9]+]]
-; CHECK:      st.volatile.u8 [%r{{[0-9]+}}], %rs[[REG]]
-; CHECK:      add.s64 %rd[[COUNTER:[0-9]+]], %rd[[COUNTER]], 1
-; CHECK-NEXT: setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
-; CHECK-NEXT: @%p[[PRED]] bra LBB[[LABEL]]
+
+; IR-LABEL:   @memcpy_volatile_caller
+; IR:         load volatile
+; IR:         store volatile
+
+; PTX-LABEL:  .visible .func (.param .b64 func_retval0) memcpy_volatile_caller
+; PTX:        LBB[[LABEL:[_0-9]+]]:
+; PTX:        ld.volatile.u8 %rs[[REG:[0-9]+]]
+; PTX:        st.volatile.u8 [%rd{{[0-9]+}}], %rs[[REG]]
+; PTX:        add.s64 %rd[[COUNTER:[0-9]+]], %rd[[COUNTER]], 1
+; PTX-NEXT:   setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
+; PTX-NEXT:   @%p[[PRED]] bra LBB[[LABEL]]
+}
+
+define i8* @memcpy_casting_caller(i32* %dst, i32* %src, i64 %n) #0 {
+entry:
+  %0 = bitcast i32* %dst to i8*
+  %1 = bitcast i32* %src to i8*
+  tail call void @llvm.memcpy.p0i8.p0i8.i64(i8* %0, i8* %1, i64 %n, i32 1, i1 false)
+  ret i8* %0
+
+; Check that casts in calls to memcpy are handled properly
+; IR-LABEL:   @memcpy_casting_caller
+; IR:         [[DSTCAST:%[0-9]+]] = bitcast i32* %dst to i8*
+; IR:         [[SRCCAST:%[0-9]+]] = bitcast i32* %src to i8*
+; IR:         getelementptr inbounds i8, i8* [[SRCCAST]]
+; IR:         getelementptr inbounds i8, i8* [[DSTCAST]]
 }
 
 define i8* @memset_caller(i8* %dst, i32 %c, i64 %n) #0 {
@@ -37,11 +70,52 @@ entry:
   %0 = trunc i32 %c to i8
   tail call void @llvm.memset.p0i8.i64(i8* %dst, i8 %0, i64 %n, i32 1, i1 false)
   ret i8* %dst
-; CHECK-LABEL: .visible .func (.param .b32 func_retval0) memset_caller(
-; CHECK:      ld.param.u8 %rs[[REG:[0-9]+]]
-; CHECK:      LBB[[LABEL:[_0-9]+]]:
-; CHECK:      st.u8 [%r{{[0-9]+}}], %rs[[REG]]
-; CHECK:      add.s64 %rd[[COUNTER:[0-9]+]], %rd[[COUNTER]], 1
-; CHECK-NEXT: setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
-; CHECK-NEXT: @%p[[PRED]] bra LBB[[LABEL]]
+
+; IR-LABEL:   @memset_caller
+; IR:         [[VAL:%[0-9]+]] = trunc i32 %c to i8
+; IR:         loadstoreloop:
+; IR:         [[STOREPTR:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64
+; IR-NEXT:    store i8 [[VAL]], i8* [[STOREPTR]]
+
+; PTX-LABEL:  .visible .func (.param .b64 func_retval0) memset_caller(
+; PTX:        ld.param.u8 %rs[[REG:[0-9]+]]
+; PTX:        LBB[[LABEL:[_0-9]+]]:
+; PTX:        st.u8 [%rd{{[0-9]+}}], %rs[[REG]]
+; PTX:        add.s64 %rd[[COUNTER:[0-9]+]], %rd[[COUNTER]], 1
+; PTX-NEXT:   setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
+; PTX-NEXT:   @%p[[PRED]] bra LBB[[LABEL]]
+}
+
+define i8* @memmove_caller(i8* %dst, i8* %src, i64 %n) #0 {
+entry:
+  tail call void @llvm.memmove.p0i8.p0i8.i64(i8* %dst, i8* %src, i64 %n, i32 1, i1 false)
+  ret i8* %dst
+
+; IR-LABEL:   @memmove_caller
+; IR:         icmp ult i8* %src, %dst
+; IR:         [[PHIVAL:%[0-9a-zA-Z_]+]] = phi i64
+; IR-NEXT:    %index_ptr = sub i64 [[PHIVAL]], 1
+; IR:         [[FWDPHIVAL:%[0-9a-zA-Z_]+]] = phi i64
+; IR:         {{%[0-9a-zA-Z_]+}} = add i64 [[FWDPHIVAL]], 1
+
+; PTX-LABEL:  .visible .func (.param .b64 func_retval0) memmove_caller(
+; PTX:        ld.param.u64 %rd[[N:[0-9]+]]
+; PTX:        setp.eq.s64 %p[[NEQ0:[0-9]+]], %rd[[N]], 0
+; PTX:        setp.ge.u64 %p[[SRC_GT_THAN_DST:[0-9]+]], %rd{{[0-9]+}}, %rd{{[0-9]+}}
+; PTX-NEXT:   @%p[[SRC_GT_THAN_DST]] bra LBB[[FORWARD_BB:[0-9_]+]]
+; -- this is the backwards copying BB
+; PTX:        @%p[[NEQ0]] bra LBB[[EXIT:[0-9_]+]]
+; PTX:        add.s64 %rd[[N]], %rd[[N]], -1
+; PTX:        ld.u8 %rs[[ELEMENT:[0-9]+]]
+; PTX:        st.u8 [%rd{{[0-9]+}}], %rs[[ELEMENT]]
+; -- this is the forwards copying BB
+; PTX:        LBB[[FORWARD_BB]]:
+; PTX:        @%p[[NEQ0]] bra LBB[[EXIT]]
+; PTX:        ld.u8 %rs[[ELEMENT2:[0-9]+]]
+; PTX:        st.u8 [%rd{{[0-9]+}}], %rs[[ELEMENT2]]
+; PTX:        add.s64 %rd[[INDEX:[0-9]+]], %rd[[INDEX]], 1
+; -- exit block
+; PTX:        LBB[[EXIT]]:
+; PTX-NEXT:   st.param.b64 [func_retval0
+; PTX-NEXT:   ret
 }
diff --git a/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll b/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll
index 0de72c4a1aed..2fffa3eeac15 100644
--- a/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll
+++ b/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll
@@ -1,7 +1,7 @@
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
 
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
-target triple = "nvptx64-unknown-unknown"
+target triple = "nvptx64-nvidia-cuda"
 
 ; Verify that both %input and %output are converted to global pointers and then
 ; addrspacecast'ed back to the original type.
@@ -26,6 +26,22 @@ define void @kernel2(float addrspace(1)* %input, float addrspace(1)* %output) {
   ret void
 }
 
-!nvvm.annotations = !{!0, !1}
+%struct.S = type { i32*, i32* }
+
+define void @ptr_in_byval(%struct.S* byval %input, i32* %output) {
+; CHECK-LABEL: .visible .entry ptr_in_byval(
+; CHECK: cvta.to.global.u64
+; CHECK: cvta.to.global.u64
+  %b_ptr = getelementptr inbounds %struct.S, %struct.S* %input, i64 0, i32 1
+  %b = load i32*, i32** %b_ptr, align 4
+  %v = load i32, i32* %b, align 4
+; CHECK: ld.global.u32
+  store i32 %v, i32* %output, align 4
+; CHECK: st.global.u32
+  ret void
+}
+
+!nvvm.annotations = !{!0, !1, !2}
 !0 = !{void (float*, float*)* @kernel, !"kernel", i32 1}
 !1 = !{void (float addrspace(1)*, float addrspace(1)*)* @kernel2, !"kernel", i32 1}
+!2 = !{void (%struct.S*, i32*)* @ptr_in_byval, !"kernel", i32 1}
diff --git a/test/CodeGen/NVPTX/reg-copy.ll b/test/CodeGen/NVPTX/reg-copy.ll
new file mode 100644
index 000000000000..98ee49d39023
--- /dev/null
+++ b/test/CodeGen/NVPTX/reg-copy.ll
@@ -0,0 +1,224 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-unknown-unknown"
+
+define void @PR24303(float* %f) {
+; CHECK-LABEL: .visible .entry PR24303(
+; Do not use mov.f or mov.u to convert between float and int.
+; CHECK-NOT: mov.{{f|u}}{{32|64}} %f{{[0-9]+}}, %r{{[0-9]+}}
+; CHECK-NOT: mov.{{f|u}}{{32|64}} %r{{[0-9]+}}, %f{{[0-9]+}}
+entry:
+  %arrayidx1 = getelementptr inbounds float, float* %f, i64 1
+  %0 = load float, float* %f, align 4
+  %1 = load float, float* %arrayidx1, align 4
+  %arrayidx2 = getelementptr inbounds float, float* %f, i64 2
+  %arrayidx3 = getelementptr inbounds float, float* %f, i64 3
+  %2 = load float, float* %arrayidx2, align 4
+  %3 = load float, float* %arrayidx3, align 4
+  %mul.i = fmul float %0, %2
+  %mul4.i = fmul float %1, %3
+  %mul5.i = fmul float %0, %3
+  %mul6.i = fmul float %1, %2
+  %sub.i = fsub float %mul.i, %mul4.i
+  %4 = bitcast float %sub.i to i32
+  %add.i = fadd float %mul6.i, %mul5.i
+  %5 = bitcast float %add.i to i32
+  %6 = tail call float @llvm.nvvm.fabs.f(float %sub.i) #2
+  %7 = fcmp ugt float %6, 0x7FF0000000000000
+  br i1 %7, label %land.lhs.true.i, label %_ZN12cuda_builtinmlIfEENS_7complexIT_EERKS3_S5_.exit
+
+land.lhs.true.i:                                  ; preds = %entry
+  %8 = tail call float @llvm.nvvm.fabs.f(float %add.i) #2
+  %9 = fcmp ugt float %8, 0x7FF0000000000000
+  br i1 %9, label %if.then.i, label %_ZN12cuda_builtinmlIfEENS_7complexIT_EERKS3_S5_.exit
+
+if.then.i:                                        ; preds = %land.lhs.true.i
+  %10 = tail call float @llvm.nvvm.fabs.f(float %0) #2
+  %11 = fcmp oeq float %10, 0x7FF0000000000000
+  %.pre.i = tail call float @llvm.nvvm.fabs.f(float %1) #2
+  %12 = fcmp oeq float %.pre.i, 0x7FF0000000000000
+  %or.cond.i = or i1 %11, %12
+  br i1 %or.cond.i, label %if.then.14.i, label %if.end.31.i
+
+if.then.14.i:                                     ; preds = %if.then.i
+  %13 = bitcast float %0 to i32
+  %14 = and i32 %13, -2147483648
+  %15 = select i1 %11, i32 1065353216, i32 0
+  %16 = or i32 %15, %14
+  %17 = bitcast i32 %16 to float
+  %18 = bitcast float %1 to i32
+  %19 = and i32 %18, -2147483648
+  %20 = select i1 %12, i32 1065353216, i32 0
+  %21 = or i32 %20, %19
+  %22 = bitcast i32 %21 to float
+  %23 = tail call float @llvm.nvvm.fabs.f(float %2) #2
+  %24 = fcmp ugt float %23, 0x7FF0000000000000
+  br i1 %24, label %if.then.24.i, label %if.end.i
+
+if.then.24.i:                                     ; preds = %if.then.14.i
+  %25 = bitcast float %2 to i32
+  %26 = and i32 %25, -2147483648
+  %27 = bitcast i32 %26 to float
+  br label %if.end.i
+
+if.end.i:                                         ; preds = %if.then.24.i, %if.then.14.i
+  %__c.0.i = phi float [ %27, %if.then.24.i ], [ %2, %if.then.14.i ]
+  %28 = tail call float @llvm.nvvm.fabs.f(float %3) #2
+  %29 = fcmp ugt float %28, 0x7FF0000000000000
+  br i1 %29, label %if.then.28.i, label %if.end.31.i
+
+if.then.28.i:                                     ; preds = %if.end.i
+  %30 = bitcast float %3 to i32
+  %31 = and i32 %30, -2147483648
+  %32 = bitcast i32 %31 to float
+  br label %if.end.31.i
+
+if.end.31.i:                                      ; preds = %if.then.28.i, %if.end.i, %if.then.i
+  %__d.1.i = phi float [ %32, %if.then.28.i ], [ %3, %if.end.i ], [ %3, %if.then.i ]
+  %__c.1.i = phi float [ %__c.0.i, %if.then.28.i ], [ %__c.0.i, %if.end.i ], [ %2, %if.then.i ]
+  %__b.0.i = phi float [ %22, %if.then.28.i ], [ %22, %if.end.i ], [ %1, %if.then.i ]
+  %__a.0.i = phi float [ %17, %if.then.28.i ], [ %17, %if.end.i ], [ %0, %if.then.i ]
+  %__recalc.0.off0.i = phi i1 [ true, %if.then.28.i ], [ true, %if.end.i ], [ false, %if.then.i ]
+  %33 = tail call float @llvm.nvvm.fabs.f(float %__c.1.i) #2
+  %34 = fcmp oeq float %33, 0x7FF0000000000000
+  %.pre6.i = tail call float @llvm.nvvm.fabs.f(float %__d.1.i) #2
+  %35 = fcmp oeq float %.pre6.i, 0x7FF0000000000000
+  %or.cond8.i = or i1 %34, %35
+  br i1 %or.cond8.i, label %if.then.37.i, label %if.end.56.i
+
+if.then.37.i:                                     ; preds = %if.end.31.i
+  %36 = bitcast float %__c.1.i to i32
+  %37 = and i32 %36, -2147483648
+  %38 = select i1 %34, i32 1065353216, i32 0
+  %39 = or i32 %38, %37
+  %40 = bitcast i32 %39 to float
+  %41 = bitcast float %__d.1.i to i32
+  %42 = and i32 %41, -2147483648
+  %43 = select i1 %35, i32 1065353216, i32 0
+  %44 = or i32 %43, %42
+  %45 = bitcast i32 %44 to float
+  %46 = tail call float @llvm.nvvm.fabs.f(float %__a.0.i) #2
+  %47 = fcmp ugt float %46, 0x7FF0000000000000
+  br i1 %47, label %if.then.48.i, label %if.end.50.i
+
+if.then.48.i:                                     ; preds = %if.then.37.i
+  %48 = bitcast float %__a.0.i to i32
+  %49 = and i32 %48, -2147483648
+  %50 = bitcast i32 %49 to float
+  br label %if.end.50.i
+
+if.end.50.i:                                      ; preds = %if.then.48.i, %if.then.37.i
+  %__a.1.i = phi float [ %50, %if.then.48.i ], [ %__a.0.i, %if.then.37.i ]
+  %51 = tail call float @llvm.nvvm.fabs.f(float %__b.0.i) #2
+  %52 = fcmp ugt float %51, 0x7FF0000000000000
+  br i1 %52, label %if.then.53.i, label %if.then.93.i
+
+if.then.53.i:                                     ; preds = %if.end.50.i
+  %53 = bitcast float %__b.0.i to i32
+  %54 = and i32 %53, -2147483648
+  %55 = bitcast i32 %54 to float
+  br label %if.then.93.i
+
+if.end.56.i:                                      ; preds = %if.end.31.i
+  br i1 %__recalc.0.off0.i, label %if.then.93.i, label %land.lhs.true.58.i
+
+land.lhs.true.58.i:                               ; preds = %if.end.56.i
+  %56 = tail call float @llvm.nvvm.fabs.f(float %mul.i) #2
+  %57 = fcmp oeq float %56, 0x7FF0000000000000
+  br i1 %57, label %if.then.70.i, label %lor.lhs.false.61.i
+
+lor.lhs.false.61.i:                               ; preds = %land.lhs.true.58.i
+  %58 = tail call float @llvm.nvvm.fabs.f(float %mul4.i) #2
+  %59 = fcmp oeq float %58, 0x7FF0000000000000
+  br i1 %59, label %if.then.70.i, label %lor.lhs.false.64.i
+
+lor.lhs.false.64.i:                               ; preds = %lor.lhs.false.61.i
+  %60 = tail call float @llvm.nvvm.fabs.f(float %mul5.i) #2
+  %61 = fcmp oeq float %60, 0x7FF0000000000000
+  br i1 %61, label %if.then.70.i, label %lor.lhs.false.67.i
+
+lor.lhs.false.67.i:                               ; preds = %lor.lhs.false.64.i
+  %62 = tail call float @llvm.nvvm.fabs.f(float %mul6.i) #2
+  %63 = fcmp oeq float %62, 0x7FF0000000000000
+  br i1 %63, label %if.then.70.i, label %_ZN12cuda_builtinmlIfEENS_7complexIT_EERKS3_S5_.exit
+
+if.then.70.i:                                     ; preds = %lor.lhs.false.67.i, %lor.lhs.false.64.i, %lor.lhs.false.61.i, %land.lhs.true.58.i
+  %64 = tail call float @llvm.nvvm.fabs.f(float %__a.0.i) #2
+  %65 = fcmp ugt float %64, 0x7FF0000000000000
+  br i1 %65, label %if.then.73.i, label %if.end.75.i
+
+if.then.73.i:                                     ; preds = %if.then.70.i
+  %66 = bitcast float %__a.0.i to i32
+  %67 = and i32 %66, -2147483648
+  %68 = bitcast i32 %67 to float
+  br label %if.end.75.i
+
+if.end.75.i:                                      ; preds = %if.then.73.i, %if.then.70.i
+  %__a.3.i = phi float [ %68, %if.then.73.i ], [ %__a.0.i, %if.then.70.i ]
+  %69 = tail call float @llvm.nvvm.fabs.f(float %__b.0.i) #2
+  %70 = fcmp ugt float %69, 0x7FF0000000000000
+  br i1 %70, label %if.then.78.i, label %if.end.80.i
+
+if.then.78.i:                                     ; preds = %if.end.75.i
+  %71 = bitcast float %__b.0.i to i32
+  %72 = and i32 %71, -2147483648
+  %73 = bitcast i32 %72 to float
+  br label %if.end.80.i
+
+if.end.80.i:                                      ; preds = %if.then.78.i, %if.end.75.i
+  %__b.3.i = phi float [ %73, %if.then.78.i ], [ %__b.0.i, %if.end.75.i ]
+  %74 = fcmp ugt float %33, 0x7FF0000000000000
+  br i1 %74, label %if.then.83.i, label %if.end.85.i
+
+if.then.83.i:                                     ; preds = %if.end.80.i
+  %75 = bitcast float %__c.1.i to i32
+  %76 = and i32 %75, -2147483648
+  %77 = bitcast i32 %76 to float
+  br label %if.end.85.i
+
+if.end.85.i:                                      ; preds = %if.then.83.i, %if.end.80.i
+  %__c.3.i = phi float [ %77, %if.then.83.i ], [ %__c.1.i, %if.end.80.i ]
+  %78 = fcmp ugt float %.pre6.i, 0x7FF0000000000000
+  br i1 %78, label %if.then.88.i, label %if.then.93.i
+
+if.then.88.i:                                     ; preds = %if.end.85.i
+  %79 = bitcast float %__d.1.i to i32
+  %80 = and i32 %79, -2147483648
+  %81 = bitcast i32 %80 to float
+  br label %if.then.93.i
+
+if.then.93.i:                                     ; preds = %if.then.88.i, %if.end.85.i, %if.end.56.i, %if.then.53.i, %if.end.50.i
+  %__d.4.ph.i = phi float [ %__d.1.i, %if.end.85.i ], [ %81, %if.then.88.i ], [ %__d.1.i, %if.end.56.i ], [ %45, %if.end.50.i ], [ %45, %if.then.53.i ]
+  %__c.4.ph.i = phi float [ %__c.3.i, %if.end.85.i ], [ %__c.3.i, %if.then.88.i ], [ %__c.1.i, %if.end.56.i ], [ %40, %if.end.50.i ], [ %40, %if.then.53.i ]
+  %__b.4.ph.i = phi float [ %__b.3.i, %if.end.85.i ], [ %__b.3.i, %if.then.88.i ], [ %__b.0.i, %if.end.56.i ], [ %__b.0.i, %if.end.50.i ], [ %55, %if.then.53.i ]
+  %__a.4.ph.i = phi float [ %__a.3.i, %if.end.85.i ], [ %__a.3.i, %if.then.88.i ], [ %__a.0.i, %if.end.56.i ], [ %__a.1.i, %if.end.50.i ], [ %__a.1.i, %if.then.53.i ]
+  %mul95.i = fmul float %__c.4.ph.i, %__a.4.ph.i
+  %mul96.i = fmul float %__d.4.ph.i, %__b.4.ph.i
+  %sub97.i = fsub float %mul95.i, %mul96.i
+  %mul98.i = fmul float %sub97.i, 0x7FF0000000000000
+  %82 = bitcast float %mul98.i to i32
+  %mul100.i = fmul float %__d.4.ph.i, %__a.4.ph.i
+  %mul101.i = fmul float %__c.4.ph.i, %__b.4.ph.i
+  %add102.i = fadd float %mul101.i, %mul100.i
+  %mul103.i = fmul float %add102.i, 0x7FF0000000000000
+  %83 = bitcast float %mul103.i to i32
+  br label %_ZN12cuda_builtinmlIfEENS_7complexIT_EERKS3_S5_.exit
+
+_ZN12cuda_builtinmlIfEENS_7complexIT_EERKS3_S5_.exit: ; preds = %if.then.93.i, %lor.lhs.false.67.i, %land.lhs.true.i, %entry
+  %84 = phi i32 [ %4, %land.lhs.true.i ], [ %4, %entry ], [ %82, %if.then.93.i ], [ %4, %lor.lhs.false.67.i ]
+  %85 = phi i32 [ %5, %land.lhs.true.i ], [ %5, %entry ], [ %83, %if.then.93.i ], [ %5, %lor.lhs.false.67.i ]
+  %arrayidx5 = getelementptr inbounds float, float* %f, i64 5
+  %86 = bitcast float* %arrayidx5 to i32*
+  store i32 %84, i32* %86, align 4
+  %arrayidx7 = getelementptr inbounds float, float* %f, i64 6
+  %87 = bitcast float* %arrayidx7 to i32*
+  store i32 %85, i32* %87, align 4
+  ret void
+}
+
+declare float @llvm.nvvm.fabs.f(float)
+
+!nvvm.annotations = !{!0}
+
+!0 = !{void (float*)* @PR24303, !"kernel", i32 1}
diff --git a/test/CodeGen/NVPTX/symbol-naming.ll b/test/CodeGen/NVPTX/symbol-naming.ll
index 0f176934ca39..7a3e6310ffdf 100644
--- a/test/CodeGen/NVPTX/symbol-naming.ll
+++ b/test/CodeGen/NVPTX/symbol-naming.ll
@@ -7,10 +7,10 @@
 ; PTX32-NOT: .str
 ; PTX64-NOT: .str
 
-; PTX32-DAG: _$_str1
+; PTX32-DAG: _$_str.1
 ; PTX32-DAG: _$_str
 
-; PTX64-DAG: _$_str1
+; PTX64-DAG: _$_str.1
 ; PTX64-DAG: _$_str
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
diff --git a/test/CodeGen/NVPTX/vector-call.ll b/test/CodeGen/NVPTX/vector-call.ll
index a03d7fd41914..968d1d4a5f51 100644
--- a/test/CodeGen/NVPTX/vector-call.ll
+++ b/test/CodeGen/NVPTX/vector-call.ll
@@ -4,7 +4,7 @@ target triple = "nvptx-unknown-cuda"
 
 declare void @bar(<4 x i32>)
 
-; CHECK-LABEL @foo
+; CHECK-LABEL: @foo
 define void @foo(<4 x i32> %a) {
 ; CHECK: st.param.v4.b32
   tail call void @bar(<4 x i32> %a)