4 files changed, 61 insertions, 0 deletions
diff --git a/test/PCH/cuda-kernel-call.cu b/test/PCH/cuda-kernel-call.cu
new file mode 100644
index 0000000000..ef12c59207
--- /dev/null
+++ b/test/PCH/cuda-kernel-call.cu
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -emit-pch -o %t %s
+// RUN: %clang_cc1 -include-pch %t -fsyntax-only %s 
+
+#ifndef HEADER
+#define HEADER
+// Header.
+
+#include "../SemaCUDA/cuda.h"
+
+void kcall(void (*kp)()) {
+  kp<<<1, 1>>>();
+}
+
+__global__ void kern() {
+}
+
+#else
+// Using the header.
+
+void test() {
+  kcall(kern);
+  kern<<<1, 1>>>();
+}
+
+#endif
diff --git a/test/Parser/cuda-kernel-call.cu b/test/Parser/cuda-kernel-call.cu
new file mode 100644
index 0000000000..f95ae9e619
--- /dev/null
+++ b/test/Parser/cuda-kernel-call.cu
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+void foo(void) {
+  foo<<<1;      // expected-error {{expected '>>>'}} expected-note {{to match this '<<<'}}
+
+  foo<<<1,1>>>; // expected-error {{expected '('}}
+
+  foo<<<>>>();  // expected-error {{expected expression}}
+}
diff --git a/test/SemaCUDA/cuda.h b/test/SemaCUDA/cuda.h
index c503747820..e3aeb99ed2 100644
--- a/test/SemaCUDA/cuda.h
+++ b/test/SemaCUDA/cuda.h
@@ -1,7 +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;
+  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/SemaCUDA/kernel-call.cu b/test/SemaCUDA/kernel-call.cu
new file mode 100644
index 0000000000..6d51695522
--- /dev/null
+++ b/test/SemaCUDA/kernel-call.cu
@@ -0,0 +1,15 @@
+// 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);
+}
+
+int main(void) {
+  g1<<<1, 1>>>(42);
+
+  t1(1);
+}