vendor/clang/clang-release_39-r276489

18 files changed, 631 insertions, 223 deletions

diff --git a/test/CodeGenCUDA/Inputs/cuda-initializers.h b/test/CodeGenCUDA/Inputs/cuda-initializers.h
new file mode 100644
index 000000000000..186b16027651
--- /dev/null
+++ b/test/CodeGenCUDA/Inputs/cuda-initializers.h
@@ -0,0 +1,145 @@
+// CUDA struct types with interesting initialization properties.
+// Keep in sync with ../SemaCUDA/Inputs/cuda-initializers.h.
+
+// Base classes with different initializer variants.
+
+// trivial constructor -- allowed
+struct T {
+  int t;
+};
+
+// empty constructor
+struct EC {
+  int ec;
+  __device__ EC() {}     // -- allowed
+  __device__ EC(int) {}  // -- not allowed
+};
+
+// empty destructor
+struct ED {
+  __device__ ~ED() {}     // -- allowed
+};
+
+struct ECD {
+  __device__ ECD() {}     // -- allowed
+  __device__ ~ECD() {}    // -- allowed
+};
+
+// empty templated constructor -- allowed with no arguments
+struct ETC {
+  template <typename... T> __device__ ETC(T...) {}
+};
+
+// undefined constructor -- not allowed
+struct UC {
+  int uc;
+  __device__ UC();
+};
+
+// undefined destructor -- not allowed
+struct UD {
+  int ud;
+  __device__ ~UD();
+};
+
+// empty constructor w/ initializer list -- not allowed
+struct ECI {
+  int eci;
+  __device__ ECI() : eci(1) {}
+};
+
+// non-empty constructor -- not allowed
+struct NEC {
+  int nec;
+  __device__ NEC() { nec = 1; }
+};
+
+// non-empty destructor -- not allowed
+struct NED {
+  int ned;
+  __device__ ~NED() { ned = 1; }
+};
+
+// no-constructor,  virtual method -- not allowed
+struct NCV {
+  int ncv;
+  __device__ virtual void vm() {}
+};
+
+// virtual destructor -- not allowed.
+struct VD {
+  __device__ virtual ~VD() {}
+};
+
+// dynamic in-class field initializer -- not allowed
+__device__ int f();
+struct NCF {
+  int ncf = f();
+};
+
+// static in-class field initializer.  NVCC does not allow it, but
+// clang generates static initializer for this, so we'll accept it.
+// We still can't use it on __shared__ vars as they don't allow *any*
+// initializers.
+struct NCFS {
+  int ncfs = 3;
+};
+
+// undefined templated constructor -- not allowed
+struct UTC {
+  template <typename... T> __device__ UTC(T...);
+};
+
+// non-empty templated constructor -- not allowed
+struct NETC {
+  int netc;
+  template <typename... T> __device__ NETC(T...) { netc = 1; }
+};
+
+// Regular base class -- allowed
+struct T_B_T : T {};
+
+// Incapsulated object of allowed class -- allowed
+struct T_F_T {
+  T t;
+};
+
+// array of allowed objects -- allowed
+struct T_FA_T {
+  T t[2];
+};
+
+
+// Calling empty base class initializer is OK
+struct EC_I_EC : EC {
+  __device__ EC_I_EC() : EC() {}
+};
+
+// .. though passing arguments is not allowed.
+struct EC_I_EC1 : EC {
+  __device__ EC_I_EC1() : EC(1) {}
+};
+
+// Virtual base class -- not allowed
+struct T_V_T : virtual T {};
+
+// Inherited from or incapsulated class with non-empty constructor --
+// not allowed
+struct T_B_NEC : NEC {};
+struct T_F_NEC {
+  NEC nec;
+};
+struct T_FA_NEC {
+  NEC nec[2];
+};
+
+
+// Inherited from or incapsulated class with non-empty desstructor --
+// not allowed
+struct T_B_NED : NED {};
+struct T_F_NED {
+  NED ned;
+};
+struct T_FA_NED {
+  NED ned[2];
+};
diff --git a/test/CodeGenCUDA/Inputs/cuda.h b/test/CodeGenCUDA/Inputs/cuda.h
index a9a4595a14a9..9b9f43a1aaa9 100644
--- a/test/CodeGenCUDA/Inputs/cuda.h
+++ b/test/CodeGenCUDA/Inputs/cuda.h
@@ -18,3 +18,5 @@ typedef struct cudaStream *cudaStream_t;
 
 int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
                       cudaStream_t stream = 0);
+
+extern "C" __device__ int printf(const char*, ...);
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index 31cba958e154..449529bb24b4 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -25,8 +25,6 @@ struct MyStruct {
 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @b = addrspace(3) global float undef
-// CHECK: @c = addrspace(3) global %struct.c undef
-// CHECK  @d = addrspace(3) global %struct.d undef
 
 __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
@@ -38,14 +36,6 @@ __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
   k++;
 
-  static int li;
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
-  li++;
-
-  __constant__ int lj;
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
-  lj++;
-
   __shared__ int lk;
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
   lk++;
@@ -102,32 +92,3 @@ __device__ float *func5() {
 }
 // CHECK: define float* @_Z5func5v()
 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
-
-struct StructWithCtor {
-  __device__ StructWithCtor(): data(1) {}
-  __device__ StructWithCtor(const StructWithCtor &second): data(second.data) {}
-  __device__ int getData() { return data; }
-  int data;
-};
-
-__device__ int construct_shared_struct() {
-// CHECK-LABEL: define i32 @_Z23construct_shared_structv()
-  __shared__ StructWithCtor s;
-// CHECK: call void @_ZN14StructWithCtorC1Ev(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
-  __shared__ StructWithCtor t(s);
-// CHECK: call void @_ZN14StructWithCtorC1ERKS_(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*), %struct.StructWithCtor* dereferenceable(4) addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
-  return t.getData();
-// CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*))
-}
-
-// Make sure we allow __shared__ structures with default or empty constructors.
-struct c {
-  int i;
-};
-__shared__ struct c c;
-
-struct d {
-  int i;
-  d() {}
-};
-__shared__ struct d d;
diff --git a/test/CodeGenCUDA/alias.cu b/test/CodeGenCUDA/alias.cu
new file mode 100644
index 000000000000..6efff6b92aa8
--- /dev/null
+++ b/test/CodeGenCUDA/alias.cu
@@ -0,0 +1,17 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
+// RUN:   -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Check that we don't generate an alias from "foo" to the mangled name for
+// ns::foo() -- nvptx doesn't support aliases.
+
+namespace ns {
+extern "C" {
+// CHECK-NOT: @foo = internal alias
+__device__ __attribute__((used)) static int foo() { return 0; }
+}
+}
diff --git a/test/CodeGenCUDA/convergent.cu b/test/CodeGenCUDA/convergent.cu
new file mode 100644
index 000000000000..6827c57d29fb
--- /dev/null
+++ b/test/CodeGenCUDA/convergent.cu
@@ -0,0 +1,45 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
+// RUN:   -disable-llvm-passes -o - %s | FileCheck -check-prefix DEVICE %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN:   -disable-llvm-passes -o - %s | \
+// RUN:  FileCheck -check-prefix HOST %s
+
+#include "Inputs/cuda.h"
+
+// DEVICE: Function Attrs:
+// DEVICE-SAME: convergent
+// DEVICE-NEXT: define void @_Z3foov
+__device__ void foo() {}
+
+// HOST: Function Attrs:
+// HOST-NOT: convergent
+// HOST-NEXT: define void @_Z3barv
+// DEVICE: Function Attrs:
+// DEVICE-SAME: convergent
+// DEVICE-NEXT: define void @_Z3barv
+__host__ __device__ void baz();
+__host__ __device__ void bar() {
+  // DEVICE: call void @_Z3bazv() [[CALL_ATTR:#[0-9]+]]
+  baz();
+  // DEVICE: call i32 asm "trap;", "=l"() [[ASM_ATTR:#[0-9]+]]
+  int x;
+  asm ("trap;" : "=l"(x));
+  // DEVICE: call void asm sideeffect "trap;", ""() [[ASM_ATTR:#[0-9]+]]
+  asm volatile ("trap;");
+}
+
+// DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
+// DEVICE: attributes [[BAZ_ATTR]] = {
+// DEVICE-SAME: convergent
+// DEVICE-SAME: }
+// DEVICE: attributes [[CALL_ATTR]] = { convergent }
+// DEVICE: attributes [[ASM_ATTR]] = { convergent
+
+// HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
+// HOST: attributes [[BAZ_ATTR]] = {
+// HOST-NOT: convergent
+// NOST-SAME: }
diff --git a/test/CodeGenCUDA/cuda-builtin-vars.cu b/test/CodeGenCUDA/cuda-builtin-vars.cu
index 834e16d04d67..c2159f5af141 100644
--- a/test/CodeGenCUDA/cuda-builtin-vars.cu
+++ b/test/CodeGenCUDA/cuda-builtin-vars.cu
@@ -6,21 +6,21 @@
 __attribute__((global))
 void kernel(int *out) {
   int i = 0;
-  out[i++] = threadIdx.x; // CHECK: call i32 @llvm.ptx.read.tid.x()
-  out[i++] = threadIdx.y; // CHECK: call i32 @llvm.ptx.read.tid.y()
-  out[i++] = threadIdx.z; // CHECK: call i32 @llvm.ptx.read.tid.z()
+  out[i++] = threadIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+  out[i++] = threadIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+  out[i++] = threadIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
 
-  out[i++] = blockIdx.x; // CHECK: call i32 @llvm.ptx.read.ctaid.x()
-  out[i++] = blockIdx.y; // CHECK: call i32 @llvm.ptx.read.ctaid.y()
-  out[i++] = blockIdx.z; // CHECK: call i32 @llvm.ptx.read.ctaid.z()
+  out[i++] = blockIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+  out[i++] = blockIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
+  out[i++] = blockIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
 
-  out[i++] = blockDim.x; // CHECK: call i32 @llvm.ptx.read.ntid.x()
-  out[i++] = blockDim.y; // CHECK: call i32 @llvm.ptx.read.ntid.y()
-  out[i++] = blockDim.z; // CHECK: call i32 @llvm.ptx.read.ntid.z()
+  out[i++] = blockDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
+  out[i++] = blockDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
+  out[i++] = blockDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
 
-  out[i++] = gridDim.x; // CHECK: call i32 @llvm.ptx.read.nctaid.x()
-  out[i++] = gridDim.y; // CHECK: call i32 @llvm.ptx.read.nctaid.y()
-  out[i++] = gridDim.z; // CHECK: call i32 @llvm.ptx.read.nctaid.z()
+  out[i++] = gridDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
+  out[i++] = gridDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
+  out[i++] = gridDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
 
   out[i++] = warpSize; // CHECK: store i32 32,
 
diff --git a/test/CodeGenCUDA/device-stub.cu b/test/CodeGenCUDA/device-stub.cu
index 7f5e159151cf..5979ba3fce60 100644
--- a/test/CodeGenCUDA/device-stub.cu
+++ b/test/CodeGenCUDA/device-stub.cu
@@ -1,7 +1,46 @@
-// RUN: %clang_cc1 -emit-llvm %s -fcuda-include-gpubinary %s -o - | FileCheck %s
+// RUN: echo "GPU binary would be here" > %t
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o - | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -fcuda-include-gpubinary %t -o -  -DNOGLOBALS \
+// RUN:   | FileCheck %s -check-prefix=NOGLOBALS
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefix=NOGPUBIN
 
 #include "Inputs/cuda.h"
 
+#ifndef NOGLOBALS
+// CHECK-DAG: @device_var = internal global i32
+__device__ int device_var;
+
+// CHECK-DAG: @constant_var = internal global i32
+__constant__ int constant_var;
+
+// CHECK-DAG: @shared_var = internal global i32
+__shared__ int shared_var;
+
+// Make sure host globals don't get internalized...
+// CHECK-DAG: @host_var = global i32
+int host_var;
+// ... and that extern vars remain external.
+// CHECK-DAG: @ext_host_var = external global i32
+extern int ext_host_var;
+
+// Shadows for external device-side variables are *definitions* of
+// those variables.
+// CHECK-DAG: @ext_device_var = internal global i32
+extern __device__ int ext_device_var;
+// CHECK-DAG: @ext_device_var = internal global i32
+extern __constant__ int ext_constant_var;
+
+void use_pointers() {
+  int *p;
+  p = &device_var;
+  p = &constant_var;
+  p = &shared_var;
+  p = &host_var;
+  p = &ext_device_var;
+  p = &ext_constant_var;
+  p = &ext_host_var;
+}
+
 // Make sure that all parts of GPU code init/cleanup are there:
 // * constant unnamed string with the kernel name
 // CHECK: private unnamed_addr constant{{.*}}kernelfunc{{.*}}\00"
@@ -31,10 +70,16 @@ __global__ void kernelfunc(int i, int j, int k) {}
 // CHECK: call{{.*}}cudaConfigureCall
 // CHECK: call{{.*}}kernelfunc
 void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
+#endif
 
-// Test that we've built a function to register kernels
-// CHECK: define internal void @__cuda_register_kernels
+// Test that we've built a function to register kernels and global vars.
+// CHECK: define internal void @__cuda_register_globals
 // CHECK: call{{.*}}cudaRegisterFunction(i8** %0, {{.*}}kernelfunc
+// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}device_var{{.*}}i32 0, i32 4, i32 0, i32 0
+// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}constant_var{{.*}}i32 0, i32 4, i32 1, i32 0
+// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_device_var{{.*}}i32 1, i32 4, i32 0, i32 0
+// CHECK-DAG: call{{.*}}cudaRegisterVar(i8** %0, {{.*}}ext_constant_var{{.*}}i32 1, i32 4, i32 1, i32 0
+// CHECK: ret void
 
 // Test that we've built contructor..
 // CHECK: define internal void @__cuda_module_ctor
@@ -42,11 +87,26 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
 // CHECK: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
 //   .. stores return value in __cuda_gpubin_handle
 // CHECK-NEXT: store{{.*}}__cuda_gpubin_handle
-//   .. and then calls __cuda_register_kernels
-// CHECK-NEXT: call void @__cuda_register_kernels
+//   .. and then calls __cuda_register_globals
+// CHECK-NEXT: call void @__cuda_register_globals
 
 // Test that we've created destructor.
 // CHECK: define internal void @__cuda_module_dtor
 // CHECK: load{{.*}}__cuda_gpubin_handle
 // CHECK-NEXT: call void @__cudaUnregisterFatBinary
 
+// There should be no __cuda_register_globals if we have no
+// device-side globals, but we still need to register GPU binary.
+// Skip GPU binary string first.
+// NOGLOBALS: @0 = private unnamed_addr constant{{.*}}
+// NOGLOBALS-NOT: define internal void @__cuda_register_globals
+// NOGLOBALS: define internal void @__cuda_module_ctor
+// NOGLOBALS: call{{.*}}cudaRegisterFatBinary{{.*}}__cuda_fatbin_wrapper
+// NOGLOBALS-NOT: call void @__cuda_register_globals
+// NOGLOBALS: define internal void @__cuda_module_dtor
+// NOGLOBALS: call void @__cudaUnregisterFatBinary
+
+// There should be no constructors/destructors if we have no GPU binary.
+// NOGPUBIN-NOT: define internal void @__cuda_register_globals
+// NOGPUBIN-NOT: define internal void @__cuda_module_ctor
+// NOGPUBIN-NOT: define internal void @__cuda_module_dtor
diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu
new file mode 100644
index 000000000000..6f2d9294131f
--- /dev/null
+++ b/test/CodeGenCUDA/device-var-init.cu
@@ -0,0 +1,198 @@
+// REQUIRES: nvptx-registered-target
+
+// Make sure we don't allow dynamic initialization for device
+// variables, but accept empty constructors allowed by CUDA.
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \
+// RUN:     -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck %s
+
+#ifdef __clang__
+#include "Inputs/cuda.h"
+#endif
+
+// Use the types we share with Sema tests.
+#include "Inputs/cuda-initializers.h"
+
+__device__ int d_v;
+// CHECK: @d_v = addrspace(1) externally_initialized global i32 0,
+__shared__ int s_v;
+// CHECK: @s_v = addrspace(3) global i32 undef,
+__constant__ int c_v;
+// CHECK: addrspace(4) externally_initialized global i32 0,
+
+__device__ int d_v_i = 1;
+// CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1,
+
+// trivial constructor -- allowed
+__device__ T d_t;
+// CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
+__shared__ T s_t;
+// CHECK: @s_t = addrspace(3) global %struct.T undef,
+__constant__ T c_t;
+// CHECK: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer,
+
+__device__ T d_t_i = {2};
+// CHECK: @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 },
+__constant__ T c_t_i = {2};
+// CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 },
+
+// empty constructor
+__device__ EC d_ec;
+// CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
+__shared__ EC s_ec;
+// CHECK: @s_ec = addrspace(3) global %struct.EC undef,
+__constant__ EC c_ec;
+// CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
+
+// empty destructor
+__device__ ED d_ed;
+// CHECK: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
+__shared__ ED s_ed;
+// CHECK: @s_ed = addrspace(3) global %struct.ED undef,
+__constant__ ED c_ed;
+// CHECK: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
+
+__device__ ECD d_ecd;
+// CHECK: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
+__shared__ ECD s_ecd;
+// CHECK: @s_ecd = addrspace(3) global %struct.ECD undef,
+__constant__ ECD c_ecd;
+// CHECK: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
+
+// empty templated constructor -- allowed with no arguments
+__device__ ETC d_etc;
+// CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
+__shared__ ETC s_etc;
+// CHECK: @s_etc = addrspace(3) global %struct.ETC undef,
+__constant__ ETC c_etc;
+// CHECK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer,
+
+__device__ NCFS d_ncfs;
+// CHECK: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 }
+__constant__ NCFS c_ncfs;
+// CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 }
+
+// Regular base class -- allowed
+__device__ T_B_T d_t_b_t;
+// CHECK: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer,
+__shared__ T_B_T s_t_b_t;
+// CHECK: @s_t_b_t = addrspace(3) global %struct.T_B_T undef,
+__constant__ T_B_T c_t_b_t;
+// CHECK: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer,
+
+// Incapsulated object of allowed class -- allowed
+__device__ T_F_T d_t_f_t;
+// CHECK: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer,
+__shared__ T_F_T s_t_f_t;
+// CHECK: @s_t_f_t = addrspace(3) global %struct.T_F_T undef,
+__constant__ T_F_T c_t_f_t;
+// CHECK: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer,
+
+// array of allowed objects -- allowed
+__device__ T_FA_T d_t_fa_t;
+// CHECK: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer,
+__shared__ T_FA_T s_t_fa_t;
+// CHECK: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef,
+__constant__ T_FA_T c_t_fa_t;
+// CHECK: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer,
+
+
+// Calling empty base class initializer is OK
+__device__ EC_I_EC d_ec_i_ec;
+// CHECK: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer,
+__shared__ EC_I_EC s_ec_i_ec;
+// CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef,
+__constant__ EC_I_EC c_ec_i_ec;
+// CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
+
+// We should not emit global initializers for device-side variables.
+// CHECK-NOT: @__cxx_global_var_init
+
+// Make sure that initialization restrictions do not apply to local
+// variables.
+__device__ void df() {
+  T t;
+  // CHECK-NOT: call
+  EC ec;
+  // CHECK:   call void @_ZN2ECC1Ev(%struct.EC* %ec)
+  ED ed;
+  // CHECK-NOT: call
+  ECD ecd;
+  // CHECK:   call void @_ZN3ECDC1Ev(%struct.ECD* %ecd)
+  ETC etc;
+  // CHECK:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
+  UC uc;
+  // undefined constructor -- not allowed
+  // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
+  UD ud;
+  // undefined destructor -- not allowed
+  // CHECK-NOT: call
+  ECI eci;
+  // empty constructor w/ initializer list -- not allowed
+  // CHECK:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
+  NEC nec;
+  // non-empty constructor -- not allowed
+  // CHECK:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+  // non-empty destructor -- not allowed
+  NED ned;
+  // no-constructor,  virtual method -- not allowed
+  // CHECK:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
+  NCV ncv;
+  // CHECK-NOT: call
+  VD vd;
+  // CHECK:   call void @_ZN2VDC1Ev(%struct.VD* %vd)
+  NCF ncf;
+  // CHECK:   call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
+  NCFS ncfs;
+  // CHECK:   call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
+  UTC utc;
+  // CHECK:   call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
+  NETC netc;
+  // CHECK:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
+  T_B_T t_b_t;
+  // CHECK-NOT: call
+  T_F_T t_f_t;
+  // CHECK-NOT: call
+  T_FA_T t_fa_t;
+  // CHECK-NOT: call
+  EC_I_EC ec_i_ec;
+  // CHECK:   call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
+  EC_I_EC1 ec_i_ec1;
+  // CHECK:   call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
+  T_V_T t_v_t;
+  // CHECK:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
+  T_B_NEC t_b_nec;
+  // CHECK:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
+  T_F_NEC t_f_nec;
+  // CHECK:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
+  T_FA_NEC t_fa_nec;
+  // CHECK:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+  T_B_NED t_b_ned;
+  // CHECK-NOT: call
+  T_F_NED t_f_ned;
+  // CHECK-NOT: call
+  T_FA_NED t_fa_ned;
+  // CHECK-NOT: call
+  static __shared__ EC s_ec;
+  // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
+  static __shared__ ETC s_etc;
+  // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
+
+  // anchor point separating constructors and destructors
+  df(); // CHECK: call void @_Z2dfv()
+
+  // Verify that we only call non-empty destructors
+  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned) #6
+  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned) #6
+  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned) #6
+  // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
+  // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
+  // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
+  // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
+  // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
+
+  // CHECK-NEXT: ret void
+}
+
+// We should not emit global init function.
+// CHECK-NOT: @_GLOBAL__sub_I
diff --git a/test/CodeGenCUDA/filter-decl.cu b/test/CodeGenCUDA/filter-decl.cu
index 023ae61f3af8..bc744a07a330 100644
--- a/test/CodeGenCUDA/filter-decl.cu
+++ b/test/CodeGenCUDA/filter-decl.cu
@@ -9,15 +9,15 @@
 // CHECK-DEVICE-NOT: module asm "file scope asm is host only"
 __asm__("file scope asm is host only");
 
-// CHECK-HOST-NOT: constantdata = externally_initialized global
+// CHECK-HOST: constantdata = internal global
 // CHECK-DEVICE: constantdata = externally_initialized global
 __constant__ char constantdata[256];
 
-// CHECK-HOST-NOT: devicedata = externally_initialized global
+// CHECK-HOST: devicedata = internal global
 // CHECK-DEVICE: devicedata = externally_initialized global
 __device__ char devicedata[256];
 
-// CHECK-HOST-NOT: shareddata = global
+// CHECK-HOST: shareddata = internal global
 // CHECK-DEVICE: shareddata = global
 __shared__ char shareddata[256];
 
diff --git a/test/CodeGenCUDA/flush-denormals.cu b/test/CodeGenCUDA/flush-denormals.cu
new file mode 100644
index 000000000000..e528d7b102d4
--- /dev/null
+++ b/test/CodeGenCUDA/flush-denormals.cu
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -fcuda-is-device \
+// RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
+// RUN:   FileCheck %s -check-prefix CHECK -check-prefix NOFTZ
+// RUN: %clang_cc1 -fcuda-is-device -fcuda-flush-denormals-to-zero \
+// RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
+// RUN:   FileCheck %s -check-prefix CHECK -check-prefix FTZ
+
+#include "Inputs/cuda.h"
+
+// Checks that device function calls get emitted with the "ntpvx-f32ftz"
+// attribute set to "true" when we compile CUDA device code with
+// -fcuda-flush-denormals-to-zero.  Further, check that we reflect the presence
+// or absence of -fcuda-flush-denormals-to-zero in a module flag.
+
+// CHECK-LABEL: define void @foo() #0
+extern "C" __device__ void foo() {}
+
+// FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true"
+// NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz"
+
+// FTZ:!llvm.module.flags = !{[[MODFLAG:![0-9]+]]}
+// FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
+
+// NOFTZ:!llvm.module.flags = !{[[MODFLAG:![0-9]+]]}
+// NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}
diff --git a/test/CodeGenCUDA/fp-contract.cu b/test/CodeGenCUDA/fp-contract.cu
new file mode 100644
index 000000000000..070ebaea44ee
--- /dev/null
+++ b/test/CodeGenCUDA/fp-contract.cu
@@ -0,0 +1,32 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// By default we should fuse multiply/add into fma instruction.
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -disable-llvm-passes -o - %s | FileCheck -check-prefix ENABLED %s
+
+// Explicit -ffp-contract=fast
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -ffp-contract=fast -disable-llvm-passes -o - %s \
+// RUN:   | FileCheck -check-prefix ENABLED %s
+
+// Explicit -ffp-contract=on -- fusing by front-end (disabled).
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -ffp-contract=on -disable-llvm-passes -o - %s \
+// RUN:   | FileCheck -check-prefix DISABLED %s
+
+// Explicit -ffp-contract=off should disable instruction fusing.
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -ffp-contract=off -disable-llvm-passes -o - %s \
+// RUN:   | FileCheck -check-prefix DISABLED %s
+
+
+#include "Inputs/cuda.h"
+
+__host__ __device__ float func(float a, float b, float c) { return a + b * c; }
+// ENABLED:       fma.rn.f32
+// ENABLED-NEXT:  st.param.f32
+
+// DISABLED:      mul.rn.f32
+// DISABLED-NEXT: add.rn.f32
+// DISABLED-NEXT: st.param.f32
diff --git a/test/CodeGenCUDA/function-overload.cu b/test/CodeGenCUDA/function-overload.cu
index a12ef82773a2..380304af8222 100644
--- a/test/CodeGenCUDA/function-overload.cu
+++ b/test/CodeGenCUDA/function-overload.cu
@@ -1,168 +1,18 @@
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
-// Make sure we handle target overloads correctly.
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
-// RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
+// Make sure we handle target overloads correctly.  Most of this is checked in
+// sema, but special functions like constructors and destructors are here.
+//
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
-// RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
 
-// Check target overloads handling with disabled call target checks.
-// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
-// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
-// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
-// RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
-// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
-// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
-// RUN:    -fcuda-is-device -o - %s \
-// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
-// RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s
-
 #include "Inputs/cuda.h"
 
-typedef int (*fp_t)(void);
-typedef void (*gp_t)(void);
-
-// CHECK-HOST: @hp = global i32 ()* @_Z1hv
-// CHECK-HOST: @chp = global i32 ()* @ch
-// CHECK-HOST: @dhp = global i32 ()* @_Z2dhv
-// CHECK-HOST: @cdhp = global i32 ()* @cdh
-// CHECK-HOST: @gp = global void ()* @_Z1gv
-
-// CHECK-BOTH-LABEL: define i32 @_Z2dhv()
-__device__ int dh(void) { return 1; }
-// CHECK-DEVICE: ret i32 1
-__host__ int dh(void) { return 2; }
-// CHECK-HOST:   ret i32 2
-
-// CHECK-BOTH-LABEL: define i32 @_Z2hdv()
-__host__ __device__ int hd(void) { return 3; }
-// CHECK-BOTH:   ret i32 3
-
-// CHECK-DEVICE-LABEL: define i32 @_Z1dv()
-__device__ int d(void) { return 8; }
-// CHECK-DEVICE:   ret i32 8
-
-// CHECK-HOST-LABEL: define i32 @_Z1hv()
-__host__ int h(void) { return 9; }
-// CHECK-HOST:   ret i32 9
-
-// CHECK-BOTH-LABEL: define void @_Z1gv()
-__global__ void g(void) {}
-// CHECK-BOTH:   ret void
-
-// mangled names of extern "C" __host__ __device__ functions clash
-// with those of their __host__/__device__ counterparts, so
-// overloading of extern "C" functions can only happen for __host__
-// and __device__ functions -- we never codegen them in the same
-// compilation and therefore mangled name conflict is not a problem.
-
-// CHECK-BOTH-LABEL: define i32 @cdh()
-extern "C" __device__ int cdh(void) {return 10;}
-// CHECK-DEVICE:   ret i32 10
-extern "C" __host__ int cdh(void) {return 11;}
-// CHECK-HOST:     ret i32 11
-
-// CHECK-DEVICE-LABEL: define i32 @cd()
-extern "C" __device__ int cd(void) {return 12;}
-// CHECK-DEVICE:   ret i32 12
-
-// CHECK-HOST-LABEL: define i32 @ch()
-extern "C" __host__ int ch(void) {return 13;}
-// CHECK-HOST:     ret i32 13
-
-// CHECK-BOTH-LABEL: define i32 @chd()
-extern "C" __host__ __device__ int chd(void) {return 14;}
-// CHECK-BOTH:     ret i32 14
-
-// CHECK-HOST-LABEL: define void @_Z5hostfv()
-__host__ void hostf(void) {
-#if defined (NOCHECKS)
-  fp_t dp = d;   // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
-  fp_t cdp = cd; // CHECK-HOST-NC: store {{.*}} @cd, {{.*}} %cdp,
-#endif
-  fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
-  fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp,
-  fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp,
-  fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp,
-  fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp,
-  fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
-  gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp,
-
-#if defined (NOCHECKS)
-  d();     // CHECK-HOST-NC: call i32 @_Z1dv()
-  cd();    // CHECK-HOST-NC: call i32 @cd()
-#endif
-  h();     // CHECK-HOST: call i32 @_Z1hv()
-  ch();    // CHECK-HOST: call i32 @ch()
-  dh();    // CHECK-HOST: call i32 @_Z2dhv()
-  cdh();   // CHECK-HOST: call i32 @cdh()
-  g<<<0,0>>>();  // CHECK-HOST: call void @_Z1gv()
-}
-
-// CHECK-DEVICE-LABEL: define void @_Z7devicefv()
-__device__ void devicef(void) {
-  fp_t dp = d;   // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
-  fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp,
-#if defined (NOCHECKS)
-  fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
-  fp_t chp = ch; // CHECK-DEVICE-NC: store {{.*}} @ch, {{.*}} %chp,
-#endif
-  fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp,
-  fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp,
-  fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp,
-  fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
-
-  d();     // CHECK-DEVICE: call i32 @_Z1dv()
-  cd();    // CHECK-DEVICE: call i32 @cd()
-#if defined (NOCHECKS)
-  h();     // CHECK-DEVICE-NC: call i32 @_Z1hv()
-  ch();    // CHECK-DEVICE-NC: call i32 @ch()
-#endif
-  dh();    // CHECK-DEVICE: call i32 @_Z2dhv()
-  cdh();   // CHECK-DEVICE: call i32 @cdh()
-}
-
-// CHECK-BOTH-LABEL: define void @_Z11hostdevicefv()
-__host__ __device__ void hostdevicef(void) {
-#if defined (NOCHECKS)
-  fp_t dp = d;   // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
-  fp_t cdp = cd; // CHECK-BOTH-NC: store {{.*}} @cd, {{.*}} %cdp,
-  fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
-  fp_t chp = ch; // CHECK-BOTH-NC: store {{.*}} @ch, {{.*}} %chp,
-#endif
-  fp_t dhp = dh; // CHECK-BOTH: store {{.*}} @_Z2dhv, {{.*}} %dhp,
-  fp_t cdhp = cdh; // CHECK-BOTH: store {{.*}} @cdh, {{.*}} %cdhp,
-  fp_t hdp = hd; // CHECK-BOTH: store {{.*}} @_Z2hdv, {{.*}} %hdp,
-  fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp,
-#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
-  gp_t gp = g; // CHECK-HOST-NC: store {{.*}} @_Z1gv, {{.*}} %gp,
-#endif
-
-#if defined (NOCHECKS)
-  d();     // CHECK-BOTH-NC: call i32 @_Z1dv()
-  cd();    // CHECK-BOTH-NC: call i32 @cd()
-  h();     // CHECK-BOTH-NC: call i32 @_Z1hv()
-  ch();    // CHECK-BOTH-NC: call i32 @ch()
-#endif
-  dh();    // CHECK-BOTH: call i32 @_Z2dhv()
-  cdh();   // CHECK-BOTH: call i32 @cdh()
-#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
-  g<<<0,0>>>();  // CHECK-HOST-NC: call void @_Z1gv()
-#endif
-}
-
-// Test for address of overloaded function resolution in the global context.
-fp_t hp = h;
-fp_t chp = ch;
-fp_t dhp = dh;
-fp_t cdhp = cdh;
-gp_t gp = g;
-
-int x;
 // Check constructors/destructors for D/H functions
+int x;
 struct s_cd_dh {
   __host__ s_cd_dh() { x = 11; }
   __device__ s_cd_dh() { x = 12; }
@@ -211,4 +61,3 @@ void wrapper() {
 // CHECK-HOST:   store i32 21,
 // CHECK-DEVICE: store i32 22,
 // CHECK-BOTH: ret void
-
diff --git a/test/CodeGenCUDA/host-device-calls-host.cu b/test/CodeGenCUDA/host-device-calls-host.cu
index 8140f619361b..94796a3c233c 100644
--- a/test/CodeGenCUDA/host-device-calls-host.cu
+++ b/test/CodeGenCUDA/host-device-calls-host.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-allow-host-calls-from-host-device -fcuda-is-device -Wno-cuda-compat -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -Wno-cuda-compat -emit-llvm -o - | FileCheck %s
 
 #include "Inputs/cuda.h"
 
diff --git a/test/CodeGenCUDA/launch-bounds.cu b/test/CodeGenCUDA/launch-bounds.cu
index ecbd0ad70580..6c369c6f3f0d 100644
--- a/test/CodeGenCUDA/launch-bounds.cu
+++ b/test/CodeGenCUDA/launch-bounds.cu
@@ -79,3 +79,8 @@ Kernel7()
 }
 // CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"maxntidx",
 // CHECK-NOT: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"minctasm",
+
+const char constchar = 12;
+__global__ void __launch_bounds__(constint, constchar) Kernel8() {}
+// CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel8{{.*}}, !"maxntidx", i32 100
+// CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel8{{.*}}, !"minctasm", i32 12
diff --git a/test/CodeGenCUDA/link-device-bitcode.cu b/test/CodeGenCUDA/link-device-bitcode.cu
index de3d39c20b49..869fcb1bc938 100644
--- a/test/CodeGenCUDA/link-device-bitcode.cu
+++ b/test/CodeGenCUDA/link-device-bitcode.cu
@@ -4,10 +4,10 @@
 // REQUIRES: nvptx-registered-target
 //
 // Prepare bitcode file to link with
-// RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc -o %t.bc \
-// RUN:    %S/Inputs/device-code.ll
-// RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc -o %t-2.bc \
-// RUN:    %S/Inputs/device-code-2.ll
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc \
+// RUN:    -disable-llvm-passes -o %t.bc %S/Inputs/device-code.ll
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -emit-llvm-bc \
+// RUN:    -disable-llvm-passes -o %t-2.bc %S/Inputs/device-code-2.ll
 //
 // Make sure function in device-code gets linked in and internalized.
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fcuda-is-device \
diff --git a/test/CodeGenCUDA/printf-aggregate.cu b/test/CodeGenCUDA/printf-aggregate.cu
new file mode 100644
index 000000000000..2e703b81d09b
--- /dev/null
+++ b/test/CodeGenCUDA/printf-aggregate.cu
@@ -0,0 +1,17 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: not %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \
+// RUN:   -o - %s 2>&1 | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Check that we don't crash when asked to printf a non-scalar arg.
+struct Struct {
+  int x;
+  int y;
+};
+__device__ void PrintfNonScalar() {
+  // CHECK: cannot compile this non-scalar arg to printf
+  printf("%d", Struct());
+}
diff --git a/test/CodeGenCUDA/printf.cu b/test/CodeGenCUDA/printf.cu
new file mode 100644
index 000000000000..dc3f4ea788f2
--- /dev/null
+++ b/test/CodeGenCUDA/printf.cu
@@ -0,0 +1,43 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \
+// RUN:   -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+extern "C" __device__ int vprintf(const char*, const char*);
+
+// Check a simple call to printf end-to-end.
+// CHECK: [[SIMPLE_PRINTF_TY:%[a-zA-Z0-9_]+]] = type { i32, i64, double }
+__device__ int CheckSimple() {
+  // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca [[SIMPLE_PRINTF_TY]]
+  // CHECK: [[FMT:%[0-9]+]] = load{{.*}}%fmt
+  const char* fmt = "%d %lld %f";
+  // CHECK: [[PTR0:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 0
+  // CHECK: store i32 1, i32* [[PTR0]], align 4
+  // CHECK: [[PTR1:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 1
+  // CHECK: store i64 2, i64* [[PTR1]], align 8
+  // CHECK: [[PTR2:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 2
+  // CHECK: store double 3.0{{[^,]*}}, double* [[PTR2]], align 8
+  // CHECK: [[BUF_CAST:%[0-9]+]] = bitcast [[SIMPLE_PRINTF_TY]]* [[BUF]] to i8*
+  // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(i8* [[FMT]], i8* [[BUF_CAST]])
+  // CHECK: ret i32 [[RET]]
+  return printf(fmt, 1, 2ll, 3.0);
+}
+
+__device__ void CheckNoArgs() {
+  // CHECK: call i32 @vprintf({{.*}}, i8* null){{$}}
+  printf("hello, world!");
+}
+
+// Check that printf's alloca happens in the entry block, not inside the if
+// statement.
+__device__ bool foo();
+__device__ void CheckAllocaIsInEntryBlock() {
+  // CHECK: alloca %printf_args
+  // CHECK: call {{.*}} @_Z3foov()
+  if (foo()) {
+    printf("%d", 42);
+  }
+}
diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu
index 6280e604f2ed..1d330bdf6a49 100644
--- a/test/CodeGenCUDA/ptx-kernels.cu
+++ b/test/CodeGenCUDA/ptx-kernels.cu
@@ -19,8 +19,17 @@ __global__ void global_function() {
 
 // Make sure host-instantiated kernels are preserved on device side.
 template <typename T> __global__ void templated_kernel(T param) {}
-// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_
-void host_function() { templated_kernel<<<0,0>>>(0); }
+// CHECK-DAG: define void @_Z16templated_kernelIiEvT_(
+
+namespace {
+__global__ void anonymous_ns_kernel() {}
+// CHECK-DAG: define void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(
+}
+
+void host_function() {
+  templated_kernel<<<0, 0>>>(0);
+  anonymous_ns_kernel<<<0,0>>>();
+}
 
 // CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
 // CHECK: !{{[0-9]+}} = !{void (i32)* @_Z16templated_kernelIiEvT_, !"kernel", i32 1}