5 files changed, 100 insertions, 0 deletions
diff --git a/clang/test/SemaCUDA/config-type.cu b/clang/test/SemaCUDA/config-type.cu
new file mode 100644
index 0000000..a469d38
--- /dev/null
+++ b/clang/test/SemaCUDA/config-type.cu
@@ -0,0 +1,3 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+void cudaConfigureCall(unsigned gridSize, unsigned blockSize); // expected-error {{must have scalar return type}}
diff --git a/clang/test/SemaCUDA/cuda.h b/clang/test/SemaCUDA/cuda.h
new file mode 100644
index 0000000..26a8df0
--- /dev/null
+++ b/clang/test/SemaCUDA/cuda.h
@@ -0,0 +1,19 @@
+/* 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))
+
+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/clang/test/SemaCUDA/function-target.cu b/clang/test/SemaCUDA/function-target.cu
new file mode 100644
index 0000000..c7a55e2
--- /dev/null
+++ b/clang/test/SemaCUDA/function-target.cu
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "cuda.h"
+
+__host__ void h1h(void);
+__device__ void h1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ function}}
+__host__ __device__ void h1hd(void);
+__global__ void h1g(void);
+
+struct h1ds { // expected-note {{requires 1 argument}}
+  __device__ h1ds(); // expected-note {{candidate constructor not viable: call to __device__ function from __host__ function}}
+};
+
+__host__ void h1(void) {
+  h1h();
+  h1d(); // expected-error {{no matching function}}
+  h1hd();
+  h1g<<<1, 1>>>();
+  h1ds x; // expected-error {{no matching constructor}}
+}
+
+__host__ void d1h(void); // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
+__device__ void d1d(void);
+__host__ __device__ void d1hd(void);
+__global__ void d1g(void); // expected-note {{'d1g' declared here}}
+
+__device__ void d1(void) {
+  d1h(); // expected-error {{no matching function}}
+  d1d();
+  d1hd();
+  d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}}
+}
+
+__host__ void hd1h(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+__device__ void hd1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+__host__ __device__ void hd1hd(void);
+__global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
+
+__host__ __device__ void hd1(void) {
+  hd1h(); // expected-error {{no matching function}}
+  hd1d(); // expected-error {{no matching function}}
+  hd1hd();
+  hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
+}
diff --git a/clang/test/SemaCUDA/kernel-call.cu b/clang/test/SemaCUDA/kernel-call.cu
new file mode 100644
index 0000000..91b1d49
--- /dev/null
+++ b/clang/test/SemaCUDA/kernel-call.cu
@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "cuda.h"
+
+__global__ void g1(int x) {}
+
+template <typename T> void t1(T arg) {
+  g1<<<arg, arg>>>(1);
+}
+
+void h1(int x) {}
+int h2(int x) { return 1; }
+
+int main(void) {
+  g1<<<1, 1>>>(42);
+  g1(42); // expected-error {{call to global function g1 not configured}}
+  g1<<<1>>>(42); // expected-error {{too few execution configuration arguments to kernel function call}}
+  g1<<<1, 1, 0, 0, 0>>>(42); // expected-error {{too many execution configuration arguments to kernel function call}}
+
+  t1(1);
+
+  h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function h1}}
+
+  int (*fp)(int) = h2;
+  fp<<<1, 1>>>(42); // expected-error {{must have void return type}}
+}
diff --git a/clang/test/SemaCUDA/qualifiers.cu b/clang/test/SemaCUDA/qualifiers.cu
new file mode 100644
index 0000000..1346d65
--- /dev/null
+++ b/clang/test/SemaCUDA/qualifiers.cu
@@ -0,0 +1,8 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "cuda.h"
+
+__global__ void g1(int x) {}
+__global__ int g2(int x) { // expected-error {{must have void return type}}
+  return 1;
+}