vendor/clang/clang-release_350-r216957

7 files changed, 126 insertions, 10 deletions

diff --git a/test/CodeGenCUDA/Inputs/cuda.h b/test/CodeGenCUDA/Inputs/cuda.h
new file mode 100644
index 000000000000..a9a4595a14a9
--- /dev/null
+++ b/test/CodeGenCUDA/Inputs/cuda.h
@@ -0,0 +1,20 @@
+/* Minimal declarations for CUDA support.  Testing purposes only. */
+
+#include <stddef.h>
+
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+typedef struct cudaStream *cudaStream_t;
+
+int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                      cudaStream_t stream = 0);
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index 04344526f40e..b80820683f3e 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -1,6 +1,9 @@
 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+// Verifies Clang emits correct address spaces and addrspacecast instructions
+// for CUDA code.
+
+#include "Inputs/cuda.h"
 
 // CHECK: @i = addrspace(1) global
 __device__ int i;
@@ -11,6 +14,18 @@ __constant__ int j;
 // CHECK: @k = addrspace(3) global
 __shared__ int k;
 
+struct MyStruct {
+  int data1;
+  int data2;
+};
+
+// CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer
+// CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00
+// CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer
+// 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 0.000000e+00
+
 __device__ void foo() {
   // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*)
   i++;
@@ -22,15 +37,66 @@ __device__ void foo() {
   k++;
 
   static int li;
-  // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li
+  // CHECK: load i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
   li++;
 
   __constant__ int lj;
-  // CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj
+  // CHECK: load i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
   lj++;
 
   __shared__ int lk;
-  // CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk
+  // CHECK: load i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
   lk++;
 }
 
+__device__ void func0() {
+  __shared__ MyStruct a;
+  MyStruct *ap = &a; // composite type
+  ap->data1 = 1;
+  ap->data2 = 2;
+}
+// CHECK: define void @_Z5func0v()
+// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap
+
+__device__ void callee(float *ap) {
+  *ap = 1.0f;
+}
+
+__device__ void func1() {
+  __shared__ float a;
+  callee(&a); // implicit cast from parameters
+}
+// CHECK: define void @_Z5func1v()
+// CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*))
+
+__device__ void func2() {
+  __shared__ float a[256];
+  float *ap = &a[128]; // implicit cast from a decayed array
+  *ap = 1.0f;
+}
+// CHECK: define void @_Z5func2v()
+// CHECK: store float* getelementptr inbounds ([256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap
+
+__device__ void func3() {
+  __shared__ float a;
+  float *ap = reinterpret_cast<float *>(&a); // explicit cast
+  *ap = 1.0f;
+}
+// CHECK: define void @_Z5func3v()
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap
+
+__device__ void func4() {
+  __shared__ float a;
+  float *ap = (float *)&a; // explicit c-style cast
+  *ap = 1.0f;
+}
+// CHECK: define void @_Z5func4v()
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap
+
+__shared__ float b;
+
+__device__ float *func5() {
+  return &b; // implicit cast from a return value
+}
+// CHECK: define float* @_Z5func5v()
+// CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
diff --git a/test/CodeGenCUDA/device-stub.cu b/test/CodeGenCUDA/device-stub.cu
index af73ea993f45..ed94d108487c 100644
--- a/test/CodeGenCUDA/device-stub.cu
+++ b/test/CodeGenCUDA/device-stub.cu
@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 // Test that we build the correct number of calls to cudaSetupArgument followed
 // by a call to cudaLaunch.
diff --git a/test/CodeGenCUDA/filter-decl.cu b/test/CodeGenCUDA/filter-decl.cu
index b758632d12f7..faaeb69fe1c7 100644
--- a/test/CodeGenCUDA/filter-decl.cu
+++ b/test/CodeGenCUDA/filter-decl.cu
@@ -1,7 +1,7 @@
-// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s
-// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s
+// RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s
+// RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 // CHECK-HOST-NOT: constantdata = global
 // CHECK-DEVICE: constantdata = global
diff --git a/test/CodeGenCUDA/kernel-call.cu b/test/CodeGenCUDA/kernel-call.cu
index f134624eec19..9b849db908f8 100644
--- a/test/CodeGenCUDA/kernel-call.cu
+++ b/test/CodeGenCUDA/kernel-call.cu
@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 __global__ void g1(int x) {}
 
diff --git a/test/CodeGenCUDA/launch-bounds.cu b/test/CodeGenCUDA/launch-bounds.cu
new file mode 100644
index 000000000000..ed4c2bfc8870
--- /dev/null
+++ b/test/CodeGenCUDA/launch-bounds.cu
@@ -0,0 +1,30 @@
+// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+#define MAX_THREADS_PER_BLOCK 256
+#define MIN_BLOCKS_PER_MP     2
+
+// Test both max threads per block and Min cta per sm.
+extern "C" {
+__global__ void
+__launch_bounds__( MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP )
+Kernel1()
+{
+}
+}
+
+// CHECK: !{{[0-9]+}} = metadata !{void ()* @Kernel1, metadata !"maxntidx", i32 256}
+// CHECK: !{{[0-9]+}} = metadata !{void ()* @Kernel1, metadata !"minctasm", i32 2}
+
+// Test only max threads per block. Min cta per sm defaults to 0, and
+// CodeGen doesn't output a zero value for minctasm.
+extern "C" {
+__global__ void
+__launch_bounds__( MAX_THREADS_PER_BLOCK )
+Kernel2()
+{
+}
+}
+
+// CHECK: !{{[0-9]+}} = metadata !{void ()* @Kernel2, metadata !"maxntidx", i32 256}
diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu
index 211692fcc7c0..11b92b587606 100644
--- a/test/CodeGenCUDA/ptx-kernels.cu
+++ b/test/CodeGenCUDA/ptx-kernels.cu
@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 // CHECK-LABEL: define void @device_function
 extern "C"