vendor/llvm/llvm-trunk-r321017

16 files changed, 840 insertions, 73 deletions

diff --git a/test/CodeGen/NVPTX/atomics-sm60.ll b/test/CodeGen/NVPTX/atomics-sm60.ll
new file mode 100644
index 0000000000000..0b5bafb780c59
--- /dev/null
+++ b/test/CodeGen/NVPTX/atomics-sm60.ll
@@ -0,0 +1,19 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_60 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 | FileCheck %s
+
+; CHECK-LABEL .func test(
+define void @test(double* %dp0, double addrspace(1)* %dp1, double addrspace(3)* %dp3, double %d) {
+; CHECK: atom.add.f64
+  %r1 = call double @llvm.nvvm.atomic.load.add.f64.p0f64(double* %dp0, double %d)
+; CHECK: atom.global.add.f64
+  %r2 = call double @llvm.nvvm.atomic.load.add.f64.p1f64(double addrspace(1)* %dp1, double %d)
+; CHECK: atom.shared.add.f64
+  %ret = call double @llvm.nvvm.atomic.load.add.f64.p3f64(double addrspace(3)* %dp3, double %d)
+  ret void
+}
+
+declare double @llvm.nvvm.atomic.load.add.f64.p0f64(double* nocapture, double) #1
+declare double @llvm.nvvm.atomic.load.add.f64.p1f64(double addrspace(1)* nocapture, double) #1
+declare double @llvm.nvvm.atomic.load.add.f64.p3f64(double addrspace(3)* nocapture, double) #1
+
+attributes #1 = { argmemonly nounwind }
diff --git a/test/CodeGen/NVPTX/barrier.ll b/test/CodeGen/NVPTX/barrier.ll
new file mode 100644
index 0000000000000..0bca04e6ddad6
--- /dev/null
+++ b/test/CodeGen/NVPTX/barrier.ll
@@ -0,0 +1,32 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
+
+declare void @llvm.nvvm.bar.warp.sync(i32)
+declare void @llvm.nvvm.barrier.sync(i32)
+declare void @llvm.nvvm.barrier.sync.cnt(i32, i32)
+
+; CHECK-LABEL: .func{{.*}}barrier.sync
+define void @barrier.sync(i32 %id, i32 %cnt) {
+  ; CHECK: ld.param.u32 	[[ID:%r[0-9]+]], [barrier.sync_param_0];
+  ; CHECK: ld.param.u32 	[[CNT:%r[0-9]+]], [barrier.sync_param_1];
+
+  ; CHECK:  barrier.sync [[ID]], [[CNT]];
+  call void @llvm.nvvm.barrier.sync.cnt(i32 %id, i32 %cnt)
+  ; CHECK:  barrier.sync [[ID]], 2;
+  call void @llvm.nvvm.barrier.sync.cnt(i32 %id, i32 2)
+  ; CHECK:  barrier.sync 3, [[CNT]];
+  call void @llvm.nvvm.barrier.sync.cnt(i32 3, i32 %cnt)
+  ; CHECK:  barrier.sync 4, 5;
+  call void @llvm.nvvm.barrier.sync.cnt(i32 4, i32 5)
+
+  ; CHECK: barrier.sync [[ID]];
+  call void @llvm.nvvm.barrier.sync(i32 %id)
+  ; CHECK: barrier.sync 1;
+  call void @llvm.nvvm.barrier.sync(i32 1)
+
+  ; CHECK: bar.warp.sync [[ID]];
+  call void @llvm.nvvm.bar.warp.sync(i32 %id)
+  ; CHECK: bar.warp.sync 6;
+  call void @llvm.nvvm.bar.warp.sync(i32 6)
+  ret void;
+}
+
diff --git a/test/CodeGen/NVPTX/fns.ll b/test/CodeGen/NVPTX/fns.ll
new file mode 100644
index 0000000000000..7673e43449c30
--- /dev/null
+++ b/test/CodeGen/NVPTX/fns.ll
@@ -0,0 +1,36 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
+
+declare i32 @llvm.nvvm.fns(i32, i32, i32)
+
+; CHECK-LABEL: .func{{.*}}fns
+define i32 @fns(i32 %mask, i32 %base, i32 %offset) {
+  ; CHECK: ld.param.u32 	[[MASK:%r[0-9]+]], [fns_param_0];
+  ; CHECK: ld.param.u32 	[[BASE:%r[0-9]+]], [fns_param_1];
+  ; CHECK: ld.param.u32 	[[OFFSET:%r[0-9]+]], [fns_param_2];
+
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, [[MASK]], [[BASE]], [[OFFSET]];
+  %r0 = call i32 @llvm.nvvm.fns(i32 %mask, i32 %base, i32 %offset);
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, [[MASK]], [[BASE]], 0;
+  %r1 = call i32 @llvm.nvvm.fns(i32 %mask, i32 %base, i32 0);
+  %r01 = add i32 %r0, %r1;
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, [[MASK]], 1, [[OFFSET]];
+  %r2 = call i32 @llvm.nvvm.fns(i32 %mask, i32 1, i32 %offset);
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, [[MASK]], 1, 0;
+  %r3 = call i32 @llvm.nvvm.fns(i32 %mask, i32 1, i32 0);
+  %r23 = add i32 %r2, %r3;
+  %r0123 = add i32 %r01, %r23;
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, 2, [[BASE]], [[OFFSET]];
+  %r4 = call i32 @llvm.nvvm.fns(i32 2, i32 %base, i32 %offset);
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, 2, [[BASE]], 0;
+  %r5 = call i32 @llvm.nvvm.fns(i32 2, i32 %base, i32 0);
+  %r45 = add i32 %r4, %r5;
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, 2, 1, [[OFFSET]];
+  %r6 = call i32 @llvm.nvvm.fns(i32 2, i32 1, i32 %offset);
+  ; CHECK:  fns.b32 	{{%r[0-9]+}}, 2, 1, 0;
+  %r7 = call i32 @llvm.nvvm.fns(i32 2, i32 1, i32 0);
+  %r67 = add i32 %r6, %r7;
+  %r4567 = add i32 %r45, %r67;
+  %r = add i32 %r0123, %r4567;
+  ret i32 %r;
+}
+
diff --git a/test/CodeGen/NVPTX/generic-to-nvvm-ir.ll b/test/CodeGen/NVPTX/generic-to-nvvm-ir.ll
index 2bd946d01eb7b..5df5183dc2fba 100644
--- a/test/CodeGen/NVPTX/generic-to-nvvm-ir.ll
+++ b/test/CodeGen/NVPTX/generic-to-nvvm-ir.ll
@@ -1,6 +1,6 @@
 ; Verify functionality of NVPTXGenericToNVVM.cpp pass.
 ;
-; RUN: opt < %s -march nvptx64 -S -generic-to-nvvm -verify-debug-info | FileCheck %s
+; RUN: opt < %s -march nvptx64 -S -generic-to-nvvm | FileCheck %s
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
@@ -42,9 +42,9 @@ declare void @extfunc(i8 signext)
 !3 = !{!4}
 ; Find list of global variables and make sure it's the one used by DICompileUnit
 ; CHECK: [[GLOBALSNODE]] = !{[[GVNODE:![0-9]+]]}
-!4 = !DIGlobalVariableExpression(var: !DIGlobalVariable(name: "static_var", scope: !0, file: !1, line: 2, type: !5, isLocal: false, isDefinition: true))
+!4 = !DIGlobalVariableExpression(var: !DIGlobalVariable(name: "static_var", scope: !0, file: !1, line: 2, type: !5, isLocal: false, isDefinition: true), expr: !DIExpression())
 ; Debug info must also be updated to reflect new address space.
-; CHECK: [[GVNODE]] = !DIGlobalVariableExpression(var: [[GVVAR:.*]])
+; CHECK: [[GVNODE]] = !DIGlobalVariableExpression(var: [[GVVAR:.*]], expr: !DIExpression())
 ; CHECK: [[GVVAR]] = !DIGlobalVariable(name: "static_var"
 ; CHECK-SAME: scope: [[CUNODE]]
 ; CHECK-SAME: type: [[TYPENODE:![0-9]+]]
diff --git a/test/CodeGen/NVPTX/i128-global.ll b/test/CodeGen/NVPTX/i128-global.ll
new file mode 100644
index 0000000000000..f53575d4ddb36
--- /dev/null
+++ b/test/CodeGen/NVPTX/i128-global.ll
@@ -0,0 +1,7 @@
+; RUN: llc < %s -O0 -march=nvptx64 -mcpu=sm_20 | FileCheck %s
+
+; CHECK: .visible .global .align 16 .b8 G1[16] = {1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+@G1 = global i128 1
+
+; CHECK: .visible .global .align 16 .b8 G2[16];
+@G2 = global i128 0
diff --git a/test/CodeGen/NVPTX/i128-param.ll b/test/CodeGen/NVPTX/i128-param.ll
new file mode 100644
index 0000000000000..7cb603546aedf
--- /dev/null
+++ b/test/CodeGen/NVPTX/i128-param.ll
@@ -0,0 +1,58 @@
+; RUN: llc < %s -O0 -march=nvptx -mcpu=sm_20 | FileCheck %s
+
+; CHECK-LABEL: .visible .func callee(
+; CHECK-NEXT: .param .align 16 .b8 callee_param_0[16],
+; CHECK-NEXT: .param .align 16 .b8 callee_param_1[16],
+define void @callee(i128, i128, i128*) {
+  ; CHECK-DAG: ld.param.v2.u64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [callee_param_0];
+  ; CHECK-DAG: ld.param.v2.u64 {%[[REG2:rd[0-9]+]], %[[REG3:rd[0-9]+]]}, [callee_param_1];
+
+  ; CHECK:      mul.lo.s64 %[[REG4:rd[0-9]+]], %[[REG0]], %[[REG3]];
+	; CHECK-NEXT: mul.hi.u64 %[[REG5:rd[0-9]+]], %[[REG0]], %[[REG2]];
+	; CHECK-NEXT: add.s64    %[[REG6:rd[0-9]+]], %[[REG5]], %[[REG4]];
+	; CHECK-NEXT: mul.lo.s64 %[[REG7:rd[0-9]+]], %[[REG1]], %[[REG2]];
+	; CHECK-NEXT: add.s64    %[[REG8:rd[0-9]+]], %[[REG6]], %[[REG7]];
+	; CHECK-NEXT: mul.lo.s64 %[[REG9:rd[0-9]+]], %[[REG0]], %[[REG2]];
+  %a = mul i128 %0, %1
+
+  store i128 %a, i128* %2
+  ret void
+}
+
+; CHECK-LABEL: .visible .entry caller_kernel(
+; CHECK-NEXT: .param .align 16 .b8 caller_kernel_param_0[16],
+; CHECK-NEXT: .param .align 16 .b8 caller_kernel_param_1[16],
+define ptx_kernel void @caller_kernel(i128, i128, i128*) {
+start:
+  ; CHECK-DAG: ld.param.v2.u64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [caller_kernel_param_0];
+  ; CHECK-DAG: ld.param.v2.u64 {%[[REG2:rd[0-9]+]], %[[REG3:rd[0-9]+]]}, [caller_kernel_param_1];
+
+  ; CHECK:      { // callseq [[CALLSEQ_ID:[0-9]]], 0
+	; CHECK:      .param .align 16 .b8 param0[16];
+	; CHECK-NEXT: st.param.v2.b64 	[param0+0], {%[[REG0]], %[[REG1]]}
+	; CHECK:      .param .align 16 .b8 param1[16];
+	; CHECK-NEXT: st.param.v2.b64 	[param1+0], {%[[REG2]], %[[REG3]]}
+	; CHECK:      } // callseq [[CALLSEQ_ID]]
+  call void @callee(i128 %0, i128 %1, i128* %2)
+
+  ret void
+}
+
+; CHECK-LABEL: .visible .func caller_func(
+; CHECK-NEXT: .param .align 16 .b8 caller_func_param_0[16],
+; CHECK-NEXT: .param .align 16 .b8 caller_func_param_1[16],
+define void @caller_func(i128, i128, i128*) {
+start:
+  ; CHECK-DAG: ld.param.v2.u64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [caller_func_param_0]
+  ; CHECK-DAG: ld.param.v2.u64 {%[[REG2:rd[0-9]+]], %[[REG3:rd[0-9]+]]}, [caller_func_param_1]
+
+  ; CHECK: { // callseq [[CALLSEQ_ID:[0-9]]], 0
+	; CHECK: .param .align 16 .b8 param0[16];
+	; CHECK: st.param.v2.b64 	[param0+0], {%[[REG0]], %[[REG1]]}
+	; CHECK: .param .align 16 .b8 param1[16];
+  ; CHECK: st.param.v2.b64 	[param1+0], {%[[REG2]], %[[REG3]]}
+	; CHECK: } // callseq [[CALLSEQ_ID]]
+  call void @callee(i128 %0, i128 %1, i128* %2)
+
+  ret void
+}
diff --git a/test/CodeGen/NVPTX/i128-retval.ll b/test/CodeGen/NVPTX/i128-retval.ll
new file mode 100644
index 0000000000000..015b0199d835e
--- /dev/null
+++ b/test/CodeGen/NVPTX/i128-retval.ll
@@ -0,0 +1,28 @@
+; RUN: llc < %s -O0 -march=nvptx64 -mcpu=sm_20 | FileCheck %s
+
+; CHECK-LABEL: .visible .func (.param .align 16 .b8 func_retval0[16]) callee(
+define i128 @callee(i128) {
+  ; CHECK: ld.param.v2.u64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [callee_param_0];
+  ; CHECK: st.param.v2.b64 [func_retval0+0], {%[[REG0]], %[[REG1]]}
+  ret i128 %0
+}
+
+; CHECK-LABEL: .visible .func caller(
+define void @caller(i128, i128*) {
+start:
+  ; CHECK-DAG: ld.param.v2.u64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [caller_param_0];
+  ; CHECK-DAG: ld.param.u64 %[[OUT:rd[0-9]+]],  [caller_param_1];
+
+  ; CHECK: { // callseq 0, 0
+	; CHECK: .param .align 16 .b8 retval0[16];
+	; CHECK: call.uni (retval0),
+  ; CHECK: ld.param.v2.b64 {%[[REG2:rd[0-9]+]], %[[REG3:rd[0-9]+]]}, [retval0+0];
+	; CHECK: } // callseq 0
+  %a = call i128 @callee(i128 %0)
+
+	; CHECK-DAG: st.u64 [%[[OUT]]], %[[REG2]];
+	; CHECK-DAG: st.u64 [%[[OUT]]+8], %[[REG3]];
+  store i128 %a, i128* %1
+
+  ret void
+}
diff --git a/test/CodeGen/NVPTX/ld-st-addrrspace.py b/test/CodeGen/NVPTX/ld-st-addrrspace.py
new file mode 100644
index 0000000000000..c9440666d99b5
--- /dev/null
+++ b/test/CodeGen/NVPTX/ld-st-addrrspace.py
@@ -0,0 +1,97 @@
+# This test generates all variants of load/store instructions and verifies that
+# LLVM generates correct PTX for them.
+
+# RUN: python %s > %t.ll
+# RUN: llc < %t.ll -march=nvptx64 -mcpu=sm_30 | FileCheck -check-prefixes=CHECK,CHECK_P64 %t.ll
+# RUN: llc < %t.ll -march=nvptx -mcpu=sm_30 | FileCheck -check-prefixes=CHECK,CHECK_P32 %t.ll
+
+from itertools import product
+from string import Template
+
+llvm_type_to_ptx_type = {
+    "i8": "u8",
+    "i16": "u16",
+    "i32": "u32",
+    "i64": "u64",
+    "half": "b16",
+    "<2 x half>": "b32",
+    "float": "f32",
+    "double": "f64"
+}
+
+llvm_type_to_ptx_reg = {
+    "i8": "r",
+    "i16": "r",
+    "i32": "r",
+    "i64": "rd",
+    "half": "h",
+    "<2 x half>": "hh",
+    "float": "f",
+    "double": "fd"
+}
+
+addrspace_id = {
+    "": 0,
+    ".global": 1,
+    ".shared": 3,
+    ".const": 4,
+    ".local": 5,
+    ".param": 101
+}
+
+
+def gen_load_tests():
+  load_template = """
+define ${type} @ld${_volatile}${_space}.${ptx_type}(${type} addrspace(${asid})* %ptr) {
+; CHECK_P32: ld${_volatile}${_volatile_as}.${ptx_type} %${ptx_reg}{{[0-9]+}}, [%r{{[0-9]+}}]
+; CHECK_P64: ld${_volatile}${_volatile_as}.${ptx_type} %${ptx_reg}{{[0-9]+}}, [%rd{{[0-9]+}}]
+; CHECK: ret
+  %p = ${generic_ptr}
+  %a = load ${volatile} ${type}, ${type}* %p
+  ret ${type} %a
+}
+"""
+  for op_type, volatile, space in product(
+      ["i8", "i16", "i32", "i64", "half", "float", "double", "<2 x half>"],
+      [True, False],  # volatile
+      ["", ".shared", ".global", ".const", ".local", ".param"]):
+
+    # Volatile is only supported for global, shared and generic.
+    if volatile and not space in ["", ".global", ".shared"]:
+      continue
+
+    # Volatile is only supported for global, shared and generic.
+    # All other volatile accesses are done in generic AS.
+    if volatile and not space in ["", ".global", ".shared"]:
+      volatile_as = ""
+    else:
+      volatile_as = space
+
+    params = {
+        "type": op_type,
+        "volatile": "volatile" if volatile else "",
+        "_volatile": ".volatile" if volatile else "",
+        "_volatile_as": volatile_as,
+        "_space": space,
+        "ptx_reg": llvm_type_to_ptx_reg[op_type],
+        "ptx_type": llvm_type_to_ptx_type[op_type],
+        "asid": addrspace_id[space],
+    }
+
+    # LLVM does not accept "addrspacecast Type* addrspace(0) to Type*", so we
+    # need to avoid it for generic pointer tests.
+    if space:
+      generic_ptr_template = ("addrspacecast ${type} addrspace(${asid})* %ptr "
+                              "to ${type}*")
+    else:
+      generic_ptr_template = "select i1 true, ${type}* %ptr, ${type}* %ptr"
+    params["generic_ptr"] = Template(generic_ptr_template).substitute(params)
+
+    print(Template(load_template).substitute(params))
+
+
+def main():
+  gen_load_tests()
+
+
+main()
diff --git a/test/CodeGen/NVPTX/lower-aggr-copies.ll b/test/CodeGen/NVPTX/lower-aggr-copies.ll
index 4298442157e23..1da1af6594705 100644
--- a/test/CodeGen/NVPTX/lower-aggr-copies.ll
+++ b/test/CodeGen/NVPTX/lower-aggr-copies.ll
@@ -1,6 +1,5 @@
 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -O0 | FileCheck %s --check-prefix PTX
 ; RUN: opt < %s -S -nvptx-lower-aggr-copies | FileCheck %s --check-prefix IR
-; RUN: opt < %s -S -nvptx-lower-aggr-copies -use-wide-memcpy-loop-lowering=true | FileCheck %s --check-prefix WIR
 
 ; Verify that the NVPTXLowerAggrCopies pass works as expected - calls to
 ; llvm.mem* intrinsics get lowered to loops.
@@ -18,13 +17,22 @@ entry:
   ret i8* %dst
 
 ; IR-LABEL:   @memcpy_caller
-; IR:         [[CMPREG:%[0-9]+]] = icmp eq i64 0, %n
-; IR:         br i1 [[CMPREG]], label %split, label %loadstoreloop
-; 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]]
+; IR:         entry:
+; IR:         [[Cond:%[0-9]+]] = icmp ne i64 %n, 0
+; IR:         br i1 [[Cond]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
+
+; IR:         loop-memcpy-expansion:
+; IR:         %loop-index = phi i64 [ 0, %entry ], [ [[IndexInc:%[0-9]+]], %loop-memcpy-expansion ]
+; IR:         [[SrcGep:%[0-9]+]] = getelementptr inbounds i8, i8* %src, i64 %loop-index
+; IR:         [[Load:%[0-9]+]] = load i8, i8* [[SrcGep]]
+; IR:         [[DstGep:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64 %loop-index
+; IR:         store i8 [[Load]], i8* [[DstGep]]
+; IR:         [[IndexInc]] = add i64 %loop-index, 1
+; IR:         [[Cond2:%[0-9]+]] = icmp ult i64 [[IndexInc]], %n
+; IR:         br i1 [[Cond2]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
+
+; IR-LABEL:   post-loop-memcpy-expansion:
+; IR:         ret i8* %dst
 
 ; PTX-LABEL:  .visible .func (.param .b64 func_retval0) memcpy_caller
 ; PTX:        LBB[[LABEL:[_0-9]+]]:
@@ -34,22 +42,6 @@ entry:
 ; PTX:        setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
 ; PTX:        @%p[[PRED]] bra LBB[[LABEL]]
 
-; WIR-LABEL:   @memcpy_caller
-; WIR:         entry:
-; WIR:         [[LoopCount:%[0-9]+]] = udiv i64 %n, 1
-; WIR:         [[ResidualSize:%[0-9]+]] = urem i64 %n, 1
-; WIR:         [[Cond:%[0-9]+]] = icmp ne i64 [[LoopCount]], 0
-; WIR:         br i1 [[Cond]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
-
-; WIR:         loop-memcpy-expansion:
-; WIR:         %loop-index = phi i64 [ 0, %entry ], [ [[IndexInc:%[0-9]+]], %loop-memcpy-expansion ]
-; WIR:         [[SrcGep:%[0-9]+]] = getelementptr inbounds i8, i8* %src, i64 %loop-index
-; WIR:         [[Load:%[0-9]+]] = load i8, i8* [[SrcGep]]
-; WIR:         [[DstGep:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64 %loop-index
-; WIR:         store i8 [[Load]], i8* [[DstGep]]
-; WIR:         [[IndexInc]] = add i64 %loop-index, 1
-; WIR:         [[Cond2:%[0-9]+]] = icmp ult i64 [[IndexInc]], [[LoopCount]]
-; WIR:         br i1 [[Cond2]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
 }
 
 define i8* @memcpy_volatile_caller(i8* %dst, i8* %src, i64 %n) #0 {
@@ -58,8 +50,23 @@ entry:
   ret i8* %dst
 
 ; IR-LABEL:   @memcpy_volatile_caller
-; IR:         load volatile
-; IR:         store volatile
+; IR:         entry:
+; IR:         [[Cond:%[0-9]+]] = icmp ne i64 %n, 0
+; IR:         br i1 [[Cond]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
+
+; IR:         loop-memcpy-expansion:
+; IR:         %loop-index = phi i64 [ 0, %entry ], [ [[IndexInc:%[0-9]+]], %loop-memcpy-expansion ]
+; IR:         [[SrcGep:%[0-9]+]] = getelementptr inbounds i8, i8* %src, i64 %loop-index
+; IR:         [[Load:%[0-9]+]] = load volatile i8, i8* [[SrcGep]]
+; IR:         [[DstGep:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64 %loop-index
+; IR:         store volatile i8 [[Load]], i8* [[DstGep]]
+; IR:         [[IndexInc]] = add i64 %loop-index, 1
+; IR:         [[Cond2:%[0-9]+]] = icmp ult i64 [[IndexInc]], %n
+; IR:         br i1 [[Cond2]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
+
+; IR-LABEL:   post-loop-memcpy-expansion:
+; IR:         ret i8* %dst
+
 
 ; PTX-LABEL:  .visible .func (.param .b64 func_retval0) memcpy_volatile_caller
 ; PTX:        LBB[[LABEL:[_0-9]+]]:
@@ -68,23 +75,6 @@ entry:
 ; PTX:        add.s64 %rd[[COUNTER:[0-9]+]], %rd{{[0-9]+}}, 1
 ; PTX:        setp.lt.u64 %p[[PRED:[0-9]+]], %rd[[COUNTER]], %rd
 ; PTX:        @%p[[PRED]] bra LBB[[LABEL]]
-
-; WIR-LABEL:   @memcpy_volatile_caller
-; WIR:         entry:
-; WIR:         [[LoopCount:%[0-9]+]] = udiv i64 %n, 1
-; WIR:         [[ResidualSize:%[0-9]+]] = urem i64 %n, 1
-; WIR:         [[Cond:%[0-9]+]] = icmp ne i64 [[LoopCount]], 0
-; WIR:         br i1 [[Cond]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
-
-; WIR:         loop-memcpy-expansion:
-; WIR:         %loop-index = phi i64 [ 0, %entry ], [ [[IndexInc:%[0-9]+]], %loop-memcpy-expansion ]
-; WIR:         [[SrcGep:%[0-9]+]] = getelementptr inbounds i8, i8* %src, i64 %loop-index
-; WIR:         [[Load:%[0-9]+]] = load volatile i8, i8* [[SrcGep]]
-; WIR:         [[DstGep:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64 %loop-index
-; WIR:         store volatile i8 [[Load]], i8* [[DstGep]]
-; WIR:         [[IndexInc]] = add i64 %loop-index, 1
-; WIR:         [[Cond2:%[0-9]+]] = icmp ult i64 [[IndexInc]], [[LoopCount]]
-; WIR:         br i1 [[Cond2]], label %loop-memcpy-expansion, label %post-loop-memcpy-expansion
 }
 
 define i8* @memcpy_casting_caller(i32* %dst, i32* %src, i64 %n) #0 {
@@ -100,12 +90,6 @@ entry:
 ; IR:         [[SRCCAST:%[0-9]+]] = bitcast i32* %src to i8*
 ; IR:         getelementptr inbounds i8, i8* [[SRCCAST]]
 ; IR:         getelementptr inbounds i8, i8* [[DSTCAST]]
-
-; WIR-LABEL:   @memcpy_casting_caller
-; WIR:         [[DSTCAST:%[0-9]+]] = bitcast i32* %dst to i8*
-; WIR:         [[SRCCAST:%[0-9]+]] = bitcast i32* %src to i8*
-; WIR:         getelementptr inbounds i8, i8* [[SRCCAST]]
-; WIR:         getelementptr inbounds i8, i8* [[DSTCAST]]
 }
 
 define i8* @memcpy_known_size(i8* %dst, i8* %src) {
@@ -114,18 +98,18 @@ entry:
   ret i8* %dst
 
 ; Check that calls with compile-time constant size are handled correctly
-; WIR-LABEL:    @memcpy_known_size
-; WIR:          entry:
-; WIR:          br label %load-store-loop
-; WIR:          load-store-loop:
-; WIR:          %loop-index = phi i64 [ 0, %entry ], [ [[IndexInc:%[0-9]+]], %load-store-loop ]
-; WIR:          [[SrcGep:%[0-9]+]] = getelementptr inbounds i8, i8* %src, i64 %loop-index
-; WIR:          [[Load:%[0-9]+]] = load i8, i8* [[SrcGep]]
-; WIR:          [[DstGep:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64 %loop-index
-; WIR:          store i8 [[Load]], i8* [[DstGep]]
-; WIR:          [[IndexInc]] = add i64 %loop-index, 1
-; WIR:          [[Cond:%[0-9]+]] = icmp ult i64 %3, 144
-; WIR:          br i1 [[Cond]], label %load-store-loop, label %memcpy-split
+; IR-LABEL:    @memcpy_known_size
+; IR:          entry:
+; IR:          br label %load-store-loop
+; IR:          load-store-loop:
+; IR:          %loop-index = phi i64 [ 0, %entry ], [ [[IndexInc:%[0-9]+]], %load-store-loop ]
+; IR:          [[SrcGep:%[0-9]+]] = getelementptr inbounds i8, i8* %src, i64 %loop-index
+; IR:          [[Load:%[0-9]+]] = load i8, i8* [[SrcGep]]
+; IR:          [[DstGep:%[0-9]+]] = getelementptr inbounds i8, i8* %dst, i64 %loop-index
+; IR:          store i8 [[Load]], i8* [[DstGep]]
+; IR:          [[IndexInc]] = add i64 %loop-index, 1
+; IR:          [[Cond:%[0-9]+]] = icmp ult i64 %3, 144
+; IR:          br i1 [[Cond]], label %load-store-loop, label %memcpy-split
 }
 
 define i8* @memset_caller(i8* %dst, i32 %c, i64 %n) #0 {
diff --git a/test/CodeGen/NVPTX/match.ll b/test/CodeGen/NVPTX/match.ll
new file mode 100644
index 0000000000000..98e1791b8dcaa
--- /dev/null
+++ b/test/CodeGen/NVPTX/match.ll
@@ -0,0 +1,117 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx60 | FileCheck %s
+
+declare i32 @llvm.nvvm.match.any.sync.i32(i32, i32)
+declare i64 @llvm.nvvm.match.any.sync.i64(i32, i64)
+
+; CHECK-LABEL: .func{{.*}}match.any.sync.i32
+define i32 @match.any.sync.i32(i32 %mask, i32 %value) {
+  ; CHECK: ld.param.u32 	[[MASK:%r[0-9]+]], [match.any.sync.i32_param_0];
+  ; CHECK: ld.param.u32 	[[VALUE:%r[0-9]+]], [match.any.sync.i32_param_1];
+
+  ; CHECK:  match.any.sync.b32  [[V0:%r[0-9]+]], [[VALUE]], [[MASK]];
+  %v0 = call i32 @llvm.nvvm.match.any.sync.i32(i32 %mask, i32 %value)
+  ; CHECK:  match.any.sync.b32  [[V1:%r[0-9]+]], [[VALUE]], 1;
+  %v1 = call i32 @llvm.nvvm.match.any.sync.i32(i32 1, i32 %value)
+  ; CHECK:  match.any.sync.b32  [[V2:%r[0-9]+]], 2, [[MASK]];
+  %v2 = call i32 @llvm.nvvm.match.any.sync.i32(i32 %mask, i32 2)
+  ; CHECK:  match.any.sync.b32  [[V3:%r[0-9]+]], 4, 3;
+  %v3 = call i32 @llvm.nvvm.match.any.sync.i32(i32 3, i32 4)
+  %sum1 = add i32 %v0, %v1
+  %sum2 = add i32 %v2, %v3
+  %sum3 = add i32 %sum1, %sum2
+  ret i32 %sum3;
+}
+
+; CHECK-LABEL: .func{{.*}}match.any.sync.i64
+define i64 @match.any.sync.i64(i32 %mask, i64 %value) {
+  ; CHECK: ld.param.u32 	[[MASK:%r[0-9]+]], [match.any.sync.i64_param_0];
+  ; CHECK: ld.param.u64 	[[VALUE:%rd[0-9]+]], [match.any.sync.i64_param_1];
+
+  ; CHECK:  match.any.sync.b64  [[V0:%rd[0-9]+]], [[VALUE]], [[MASK]];
+  %v0 = call i64 @llvm.nvvm.match.any.sync.i64(i32 %mask, i64 %value)
+  ; CHECK:  match.any.sync.b64  [[V1:%rd[0-9]+]], [[VALUE]], 1;
+  %v1 = call i64 @llvm.nvvm.match.any.sync.i64(i32 1, i64 %value)
+  ; CHECK:  match.any.sync.b64  [[V2:%rd[0-9]+]], 2, [[MASK]];
+  %v2 = call i64 @llvm.nvvm.match.any.sync.i64(i32 %mask, i64 2)
+  ; CHECK:  match.any.sync.b64  [[V3:%rd[0-9]+]], 4, 3;
+  %v3 = call i64 @llvm.nvvm.match.any.sync.i64(i32 3, i64 4)
+  %sum1 = add i64 %v0, %v1
+  %sum2 = add i64 %v2, %v3
+  %sum3 = add i64 %sum1, %sum2
+  ret i64 %sum3;
+}
+
+declare {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32, i32)
+declare {i64, i1} @llvm.nvvm.match.all.sync.i64p(i32, i64)
+
+; CHECK-LABEL: .func{{.*}}match.all.sync.i32p(
+define {i32,i1} @match.all.sync.i32p(i32 %mask, i32 %value) {
+  ; CHECK: ld.param.u32 	[[MASK:%r[0-9]+]], [match.all.sync.i32p_param_0];
+  ; CHECK: ld.param.u32 	[[VALUE:%r[0-9]+]], [match.all.sync.i32p_param_1];
+
+  ; CHECK:  match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], [[MASK]];
+  %r1 = call {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32 %mask, i32 %value)
+  %v1 = extractvalue {i32, i1} %r1, 0
+  %p1 = extractvalue {i32, i1} %r1, 1
+
+  ; CHECK:  match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, 1, [[MASK]];
+  %r2 = call {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32 %mask, i32 1)
+  %v2 = extractvalue {i32, i1} %r2, 0
+  %p2 = extractvalue {i32, i1} %r2, 1
+
+  ; CHECK:  match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], 2;
+  %r3 = call {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32 2, i32 %value)
+  %v3 = extractvalue {i32, i1} %r3, 0
+  %p3 = extractvalue {i32, i1} %r3, 1
+
+  ; CHECK:  match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, 4, 3;
+  %r4 = call {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32 3, i32 4)
+  %v4 = extractvalue {i32, i1} %r4, 0
+  %p4 = extractvalue {i32, i1} %r4, 1
+
+  %vsum1 = add i32 %v1, %v2
+  %vsum2 = add i32 %v3, %v4
+  %vsum3 = add i32 %vsum1, %vsum2
+  %psum1 = add i1 %p1, %p2
+  %psum2 = add i1 %p3, %p4
+  %psum3 = add i1 %psum1, %psum2
+  %ret0 = insertvalue {i32, i1} undef, i32 %vsum3, 0
+  %ret1 = insertvalue {i32, i1} %ret0, i1 %psum3, 1
+  ret {i32, i1} %ret1;
+}
+
+; CHECK-LABEL: .func{{.*}}match.all.sync.i64p(
+define {i64,i1} @match.all.sync.i64p(i32 %mask, i64 %value) {
+  ; CHECK: ld.param.u32 	[[MASK:%r[0-9]+]], [match.all.sync.i64p_param_0];
+  ; CHECK: ld.param.u64 	[[VALUE:%rd[0-9]+]], [match.all.sync.i64p_param_1];
+
+  ; CHECK:  match.all.sync.b64 {{%rd[0-9]+\|%p[0-9]+}}, [[VALUE]], [[MASK]];
+  %r1 = call {i64, i1} @llvm.nvvm.match.all.sync.i64p(i32 %mask, i64 %value)
+  %v1 = extractvalue {i64, i1} %r1, 0
+  %p1 = extractvalue {i64, i1} %r1, 1
+
+  ; CHECK:  match.all.sync.b64 {{%rd[0-9]+\|%p[0-9]+}}, 1, [[MASK]];
+  %r2 = call {i64, i1} @llvm.nvvm.match.all.sync.i64p(i32 %mask, i64 1)
+  %v2 = extractvalue {i64, i1} %r2, 0
+  %p2 = extractvalue {i64, i1} %r2, 1
+
+  ; CHECK:  match.all.sync.b64 {{%rd[0-9]+\|%p[0-9]+}}, [[VALUE]], 2;
+  %r3 = call {i64, i1} @llvm.nvvm.match.all.sync.i64p(i32 2, i64 %value)
+  %v3 = extractvalue {i64, i1} %r3, 0
+  %p3 = extractvalue {i64, i1} %r3, 1
+
+  ; CHECK:  match.all.sync.b64 {{%rd[0-9]+\|%p[0-9]+}}, 4, 3;
+  %r4 = call {i64, i1} @llvm.nvvm.match.all.sync.i64p(i32 3, i64 4)
+  %v4 = extractvalue {i64, i1} %r4, 0
+  %p4 = extractvalue {i64, i1} %r4, 1
+
+  %vsum1 = add i64 %v1, %v2
+  %vsum2 = add i64 %v3, %v4
+  %vsum3 = add i64 %vsum1, %vsum2
+  %psum1 = add i1 %p1, %p2
+  %psum2 = add i1 %p3, %p4
+  %psum3 = add i1 %psum1, %psum2
+  %ret0 = insertvalue {i64, i1} undef, i64 %vsum3, 0
+  %ret1 = insertvalue {i64, i1} %ret0, i1 %psum3, 1
+  ret {i64, i1} %ret1;
+}
diff --git a/test/CodeGen/NVPTX/minmax-negative.ll b/test/CodeGen/NVPTX/minmax-negative.ll
new file mode 100644
index 0000000000000..73bd9e049e7a7
--- /dev/null
+++ b/test/CodeGen/NVPTX/minmax-negative.ll
@@ -0,0 +1,9 @@
+; RUN: llc < %s -march=nvptx -O0 | FileCheck %s
+
+define i16 @test1(i16* %sur1) {
+; CHECK-NOT: mov.u16 %rs{{[0-9]+}}, 32767
+  %_tmp21.i = icmp sle i16 0, 0
+  %_tmp22.i = select i1 %_tmp21.i, i16 0, i16 32767
+  store i16 %_tmp22.i, i16* %sur1
+  ret i16 0
+}
diff --git a/test/CodeGen/NVPTX/shfl-sync.ll b/test/CodeGen/NVPTX/shfl-sync.ll
new file mode 100644
index 0000000000000..521e045b65457
--- /dev/null
+++ b/test/CodeGen/NVPTX/shfl-sync.ll
@@ -0,0 +1,94 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 | FileCheck %s
+
+declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
+declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
+declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
+declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
+declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
+declare float @llvm.nvvm.shfl.sync.idx.f32(float, i32, i32, i32)
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.rrr
+define i32 @shfl.sync.rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
+  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 %c)
+  ret i32 %val
+}
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.irr
+define i32 @shfl.sync.irr(i32 %a, i32 %b, i32 %c) {
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], 1;
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 %c)
+  ret i32 %val
+}
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.rri
+define i32 @shfl.sync.rri(i32 %mask, i32 %a, i32 %b) {
+  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 1, [[MASK]];
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 1)
+  ret i32 %val
+}
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.iri
+define i32 @shfl.sync.iri(i32 %a, i32 %b) {
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 2)
+  ret i32 %val
+}
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.rir
+define i32 @shfl.sync.rir(i32 %mask, i32 %a, i32 %c) {
+  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, [[C]], [[MASK]];
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 %c)
+  ret i32 %val
+}
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.iir
+define i32 @shfl.sync.iir(i32 %a, i32 %c) {
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 %c)
+  ret i32 %val
+}
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.rii
+define i32 @shfl.sync.rii(i32 %mask, i32 %a) {
+  ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, 2, [[MASK]];
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 2)
+  ret i32 %val
+}
+
+; CHECK-LABEL: .func{{.*}}shfl.sync.iii
+define i32 @shfl.sync.iii(i32 %a, i32 %b) {
+  ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
+  ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, 3, 1;
+  ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
+  %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 3)
+  ret i32 %val
+}
diff --git a/test/CodeGen/NVPTX/sm-version-70.ll b/test/CodeGen/NVPTX/sm-version-70.ll
new file mode 100644
index 0000000000000..8b72d50747ae5
--- /dev/null
+++ b/test/CodeGen/NVPTX/sm-version-70.ll
@@ -0,0 +1,5 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 | FileCheck %s
+
+; CHECK: .version 6.0
+; CHECK: .target sm_70
diff --git a/test/CodeGen/NVPTX/symbol-naming.ll b/test/CodeGen/NVPTX/symbol-naming.ll
index 7a3e6310ffdf9..3f1caf927dc81 100644
--- a/test/CodeGen/NVPTX/symbol-naming.ll
+++ b/test/CodeGen/NVPTX/symbol-naming.ll
@@ -1,17 +1,17 @@
-; 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
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
 
 ; Verify that the NVPTX target removes invalid symbol names prior to emitting
 ; PTX.
 
-; PTX32-NOT: .str
-; PTX64-NOT: .str
+; CHECK-NOT: .str
+; CHECK-NOT: .function.
 
-; PTX32-DAG: _$_str.1
-; PTX32-DAG: _$_str
+; CHECK-DAG: _$_str
+; CHECK-DAG: _$_str1
 
-; PTX64-DAG: _$_str.1
-; PTX64-DAG: _$_str
+; CHECK-DAG: _$_function_$_
+; CHECK-DAG: _$_function_$_2
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-unknown-unknown"
@@ -22,10 +22,25 @@ target triple = "nvptx64-unknown-unknown"
 
 
 ; Function Attrs: nounwind
-define void @foo(i32 %a, float %b, i8 signext %c, i32 %e) {
+define internal void @.function.() {
 entry:
   %call = call i32 (i8*, ...) @printf(i8* getelementptr inbounds ([13 x i8], [13 x i8]* @.str, i32 0, i32 0))
   ret void
 }
 
+; Function Attrs: nounwind
+define internal void @_$_function_$_() {
+entry:
+  %call = call i32 (i8*, ...) @printf(i8* getelementptr inbounds ([13 x i8], [13 x i8]* @_$_str, i32 0, i32 0))
+  ret void
+}
+
+; Function Attrs: nounwind
+define void @global_function() {
+entry:
+  call void @.function.()
+  call void @_$_function_$_()
+  ret void
+}
+
 declare i32 @printf(i8*, ...)
diff --git a/test/CodeGen/NVPTX/vote.ll b/test/CodeGen/NVPTX/vote.ll
new file mode 100644
index 0000000000000..c28e795fe5733
--- /dev/null
+++ b/test/CodeGen/NVPTX/vote.ll
@@ -0,0 +1,65 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
+
+declare i1 @llvm.nvvm.vote.all(i1)
+; CHECK-LABEL: .func{{.*}}vote.all
+define i1 @vote.all(i1 %pred) {
+  ; CHECK: vote.all.pred
+  %val = call i1 @llvm.nvvm.vote.all(i1 %pred)
+  ret i1 %val
+}
+
+declare i1 @llvm.nvvm.vote.any(i1)
+; CHECK-LABEL: .func{{.*}}vote.any
+define i1 @vote.any(i1 %pred) {
+  ; CHECK: vote.any.pred
+  %val = call i1 @llvm.nvvm.vote.any(i1 %pred)
+  ret i1 %val
+}
+
+declare i1 @llvm.nvvm.vote.uni(i1)
+; CHECK-LABEL: .func{{.*}}vote.uni
+define i1 @vote.uni(i1 %pred) {
+  ; CHECK: vote.uni.pred
+  %val = call i1 @llvm.nvvm.vote.uni(i1 %pred)
+  ret i1 %val
+}
+
+declare i32 @llvm.nvvm.vote.ballot(i1)
+; CHECK-LABEL: .func{{.*}}vote.ballot
+define i32 @vote.ballot(i1 %pred) {
+  ; CHECK: vote.ballot.b32
+  %val = call i32 @llvm.nvvm.vote.ballot(i1 %pred)
+  ret i32 %val
+}
+
+declare i1 @llvm.nvvm.vote.all.sync(i32, i1)
+; CHECK-LABEL: .func{{.*}}vote.sync.all
+define i1 @vote.sync.all(i32 %mask, i1 %pred) {
+  ; CHECK: vote.sync.all.pred
+  %val = call i1 @llvm.nvvm.vote.all.sync(i32 %mask, i1 %pred)
+  ret i1 %val
+}
+
+declare i1 @llvm.nvvm.vote.any.sync(i32, i1)
+; CHECK-LABEL: .func{{.*}}vote.sync.any
+define i1 @vote.sync.any(i32 %mask, i1 %pred) {
+  ; CHECK: vote.sync.any.pred
+  %val = call i1 @llvm.nvvm.vote.any.sync(i32 %mask, i1 %pred)
+  ret i1 %val
+}
+
+declare i1 @llvm.nvvm.vote.uni.sync(i32, i1)
+; CHECK-LABEL: .func{{.*}}vote.sync.uni
+define i1 @vote.sync.uni(i32 %mask, i1 %pred) {
+  ; CHECK: vote.sync.uni.pred
+  %val = call i1 @llvm.nvvm.vote.uni.sync(i32 %mask, i1 %pred)
+  ret i1 %val
+}
+
+declare i32 @llvm.nvvm.vote.ballot.sync(i32, i1)
+; CHECK-LABEL: .func{{.*}}vote.sync.ballot
+define i32 @vote.sync.ballot(i32 %mask, i1 %pred) {
+  ; CHECK: vote.sync.ballot.b32
+  %val = call i32 @llvm.nvvm.vote.ballot.sync(i32 %mask, i1 %pred)
+  ret i32 %val
+}
diff --git a/test/CodeGen/NVPTX/wmma.py b/test/CodeGen/NVPTX/wmma.py
new file mode 100644
index 0000000000000..ad62b84f41773
--- /dev/null
+++ b/test/CodeGen/NVPTX/wmma.py
@@ -0,0 +1,201 @@
+# This test generates all variants of wmma intrinsics and verifies that LLVM
+# generates correct instructions for them.
+
+# RUN: python %s > %t.ll
+# RUN: llc < %t.ll -march=nvptx64 -mcpu=sm_70 -mattr=+ptx60 | FileCheck %t.ll
+
+from itertools import product
+from string import Template
+
+def make_wmma_slice_ty(abcd, itype):
+  elt_ty = "<2 x half>" if itype == "f16" else "float"
+  num_elts = 4 if abcd in "cd" and itype == "f16" else 8;
+  return [elt_ty] * num_elts
+
+def make_wmma_ld_ret_ty(abc, itype):
+  return "{%s}" % ", ".join(make_wmma_slice_ty(abc, itype))
+
+# Convenient test patterns.
+check_f16_8 = "{{%s}}" % ", *".join(["%hh[0-9]+"] * 8)
+check_f16_4 = "{{%s}}" % ", *".join(["%hh[0-9]+"] * 4)
+check_f32_8 = "{{%s}}" % ", *".join(["%f[0-9]+"] * 8)
+
+def gen_wmma_load_tests():
+  load_template = """
+declare ${ret_ty} @llvm.nvvm.wmma.load.$intrinsic_suffix(i8* %src ${extra_args});
+
+; CHECK-LABEL: .func {{.*}}test_wmma_load_${function_suffix}(
+define ${ret_ty} @test_wmma_load_${function_suffix}(i8* %src ${extra_args}) {
+; CHECK wmma.load.${intrinsic_suffix}
+; CHECK: {${check_result}}
+; CHECK: [%rd{{[0-9]+}}]${stride_pattern}
+  %v0 = call ${ret_ty} @llvm.nvvm.wmma.load.${intrinsic_suffix}(i8* %src ${extra_args});
+  ret ${ret_ty} %v0;
+}
+
+; CHECK-LABEL: .func{{.*}}test_wmma_load_${function_suffix}_o(
+define ${ret_ty} @test_wmma_load_${function_suffix}_o(i8* %src ${extra_args}) {
+; CHECK wmma.load.${intrinsic_suffix}
+; CHECK: {${check_result}}
+; CHECK: [%rd{{[0-9]+}}+128]${stride_pattern}
+  %src1 = getelementptr i8, i8* %src, i32 128;
+  %v0 = call ${ret_ty} @llvm.nvvm.wmma.load.${intrinsic_suffix}(i8* %src1 ${extra_args});
+  ret ${ret_ty} %v0;
+}
+"""
+  suffix_template = "${abc}.sync.${layout}.m16n16k16${space}${stride}.${itype}"
+  instruction_template = "${abc}.sync.${layout}.m16n16k16${space}.${itype}"
+
+  for abc, layout, space, stride, itype in product(
+      "abc",
+      ["row","col"],
+      ["",".shared",".global"],
+      ["", ".stride"],
+      ["f16", "f32"]):
+
+    params = {
+        "abc" : abc,
+        "layout" : layout,
+        "space" : space,
+        "stride" : stride,
+        "itype" : itype
+    }
+
+    if itype == "f32" and abc != "c":
+      continue
+
+    test_params = params
+    test_params["intrinsic_suffix"] = Template(suffix_template).substitute(params)
+    test_params["function_suffix"] = test_params["intrinsic_suffix"].replace(".","_")
+    test_params["instruction_suffix"] = Template(instruction_template).substitute(params)
+    test_params["ret_ty"] = make_wmma_ld_ret_ty(abc, itype)
+    if abc == "c" :
+      test_params["check_result"] = check_f16_4 if itype == "f16" else check_f32_8
+    else:
+      test_params["check_result"] = check_f16_8
+
+    if stride:
+      test_params["extra_args"] = ", i32 %stride";
+      test_params["stride_pattern"] = ", %r{{[0-9]+}}"
+    else:
+      test_params["extra_args"] = ""
+      test_params["stride_pattern"] = ""
+
+    print(Template(load_template).substitute(test_params))
+
+def make_wmma_slice_args(itype, abcd, prefix="v"):
+  return ", ".join(["%s %%%s%d" % (t, prefix, i) for i,t
+                  in enumerate(make_wmma_slice_ty(abcd, itype))])
+
+def gen_wmma_store_tests():
+  store_template = """
+declare void @llvm.nvvm.wmma.store.$intrinsic_suffix(i8* %src, ${args}${extra_args});
+
+; CHECK-LABEL: .func {{.*}}test_wmma_store_${function_suffix}(
+define void @test_wmma_store_${function_suffix}(i8* %src, ${args}${extra_args}) {
+; CHECK wmma.store.${intrinsic_suffix} {{.*}}[%rd{{[0-9+]}}
+; CHECK: {${check_args}}
+; CHECK: ${stride_pattern}
+  call void @llvm.nvvm.wmma.store.${intrinsic_suffix}(i8* %src, ${args} ${extra_args});
+  ret void
+}
+
+; CHECK-LABEL: .func{{.*}}test_wmma_store_${function_suffix}_o(
+define void @test_wmma_store_${function_suffix}_o(i8* %src, ${args}${extra_args}) {
+; CHECK wmma.store.${intrinsic_suffix} {{.*}}[%rd{{[0-9+]}}+128]
+; CHECK: ${check_args}
+; CHECK: ${stride_pattern}
+  %src1 = getelementptr i8, i8* %src, i32 128;
+  call void @llvm.nvvm.wmma.store.${intrinsic_suffix}(i8* %src1, ${args}${extra_args});
+  ret void
+}
+"""
+  suffix_template = "${abc}.sync.${layout}.m16n16k16${space}${stride}.${itype}"
+  instruction_template = "${abc}.sync.${layout}.m16n16k16${space}.${itype}"
+
+  for abc, layout, space, stride, itype in product(
+      "d",
+      ["row","col"],
+      ["",".shared",".global"],
+      ["", ".stride"],
+      ["f16", "f32"]):
+
+    params = {
+        "abc" : abc,
+        "layout" : layout,
+        "space" : space,
+        "stride" : stride,
+        "itype" : itype
+    }
+
+    test_params = params
+    test_params["intrinsic_suffix"] = Template(suffix_template).substitute(params)
+    test_params["function_suffix"] = test_params["intrinsic_suffix"].replace(".","_")
+    test_params["instruction_suffix"] = Template(instruction_template).substitute(params)
+    test_params["ret_ty"] = make_wmma_ld_ret_ty(abc, itype)
+    test_params["check_args"] = check_f16_4 if itype == "f16" else check_f32_8
+    if stride:
+      test_params["extra_args"] = ", i32 %stride";
+      test_params["stride_pattern"] = ", %r{{[0-9]+}};"
+    else:
+      test_params["extra_args"] = ""
+      test_params["stride_pattern"] = ";"
+    test_params["args"] = make_wmma_slice_args(itype, "d");
+
+    print(Template(store_template).substitute(test_params))
+
+def gen_wmma_mma_tests():
+  mma_template = """
+declare ${ret_ty} @llvm.nvvm.wmma.mma.sync.$intrinsic_suffix(
+        ${args});
+
+; CHECK-LABEL: .func {{.*}}test_wmma_mma_${function_suffix}(
+define ${ret_ty} @test_wmma_mma_${function_suffix}(
+        ${args}) {
+; CHECK wmma.mma.${intrinsic_suffix} {{.*}}[%rd{{[0-9+]}}
+; CHECK ${check_d}
+; CHECK ${check_ab}
+; CHECK ${check_ab}
+; CHECK ${check_c}
+  %r = call ${ret_ty} @llvm.nvvm.wmma.mma.sync.${intrinsic_suffix}(
+        ${args});
+  ret ${ret_ty} %r;
+}
+"""
+  suffix_template = "${alayout}.${blayout}.m16n16k16.${dtype}.${ctype}${satf}"
+
+  for alayout, blayout, ctype, dtype, satf in product(
+      ["row","col"],
+      ["row","col"],
+      ["f16", "f32"],
+      ["f16", "f32"],
+      [".satfinite", ""]):
+
+    params = {
+        "alayout" : alayout,
+        "blayout" : blayout,
+        "ctype" : ctype,
+        "dtype" : dtype,
+        "satf"  : satf
+    }
+
+    test_params = params
+    test_params["intrinsic_suffix"] = Template(suffix_template).substitute(params)
+    test_params["function_suffix"] = test_params["intrinsic_suffix"].replace(".", "_")
+    test_params["ret_ty"] = make_wmma_ld_ret_ty("d", dtype)
+    test_params["check_ab"] = check_f16_8
+    test_params["check_c"] = check_f16_4 if ctype == "f16" else check_f32_8
+    test_params["check_d"] = check_f16_4 if dtype == "f16" else check_f32_8
+    args = ",\n        ".join(make_wmma_slice_args(t, abcd, prefix=abcd)
+                              for abcd, t in (("a", "f16"),
+                                              ("b", "f16"),
+                                              ("c", ctype)))
+    test_params["args"] = args
+    print(Template(mma_template).substitute(test_params))
+
+def main():
+  gen_wmma_load_tests()
+  gen_wmma_store_tests()
+  gen_wmma_mma_tests()
+
+main()