4 files changed, 78 insertions, 0 deletions
diff --git a/clang/test/CodeGenCUDA/device-stub.cu b/clang/test/CodeGenCUDA/device-stub.cu
new file mode 100644
index 0000000..af73ea9
--- /dev/null
+++ b/clang/test/CodeGenCUDA/device-stub.cu
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
+
+#include "../SemaCUDA/cuda.h"
+
+// Test that we build the correct number of calls to cudaSetupArgument followed
+// by a call to cudaLaunch.
+
+// CHECK: define{{.*}}kernelfunc
+// CHECK: call{{.*}}cudaSetupArgument
+// CHECK: call{{.*}}cudaSetupArgument
+// CHECK: call{{.*}}cudaSetupArgument
+// CHECK: call{{.*}}cudaLaunch
+__global__ void kernelfunc(int i, int j, int k) {}
diff --git a/clang/test/CodeGenCUDA/filter-decl.cu b/clang/test/CodeGenCUDA/filter-decl.cu
new file mode 100644
index 0000000..b758632
--- /dev/null
+++ b/clang/test/CodeGenCUDA/filter-decl.cu
@@ -0,0 +1,40 @@
+// 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
+
+#include "../SemaCUDA/cuda.h"
+
+// CHECK-HOST-NOT: constantdata = global
+// CHECK-DEVICE: constantdata = global
+__constant__ char constantdata[256];
+
+// CHECK-HOST-NOT: devicedata = global
+// CHECK-DEVICE: devicedata = global
+__device__ char devicedata[256];
+
+// CHECK-HOST-NOT: shareddata = global
+// CHECK-DEVICE: shareddata = global
+__shared__ char shareddata[256];
+
+// CHECK-HOST: hostdata = global
+// CHECK-DEVICE-NOT: hostdata = global
+char hostdata[256];
+
+// CHECK-HOST: define{{.*}}implicithostonlyfunc
+// CHECK-DEVICE-NOT: define{{.*}}implicithostonlyfunc
+void implicithostonlyfunc(void) {}
+
+// CHECK-HOST: define{{.*}}explicithostonlyfunc
+// CHECK-DEVICE-NOT: define{{.*}}explicithostonlyfunc
+__host__ void explicithostonlyfunc(void) {}
+
+// CHECK-HOST-NOT: define{{.*}}deviceonlyfunc
+// CHECK-DEVICE: define{{.*}}deviceonlyfunc
+__device__ void deviceonlyfunc(void) {}
+
+// CHECK-HOST: define{{.*}}hostdevicefunc
+// CHECK-DEVICE: define{{.*}}hostdevicefunc
+__host__  __device__ void hostdevicefunc(void) {}
+
+// CHECK-HOST: define{{.*}}globalfunc
+// CHECK-DEVICE: define{{.*}}globalfunc
+__global__ void globalfunc(void) {}
diff --git a/clang/test/CodeGenCUDA/kernel-call.cu b/clang/test/CodeGenCUDA/kernel-call.cu
new file mode 100644
index 0000000..f134624
--- /dev/null
+++ b/clang/test/CodeGenCUDA/kernel-call.cu
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
+
+#include "../SemaCUDA/cuda.h"
+
+__global__ void g1(int x) {}
+
+int main(void) {
+  // CHECK: call{{.*}}cudaConfigureCall
+  // CHECK: icmp
+  // CHECK: br
+  // CHECK: call{{.*}}g1
+  g1<<<1, 1>>>(42);
+}
diff --git a/clang/test/CodeGenCUDA/ptx-kernels.cu b/clang/test/CodeGenCUDA/ptx-kernels.cu
new file mode 100644
index 0000000..ecca851
--- /dev/null
+++ b/clang/test/CodeGenCUDA/ptx-kernels.cu
@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 %s -triple ptx32-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
+
+#include "../SemaCUDA/cuda.h"
+
+// CHECK: define ptx_device{{.*}}device_function
+__device__ void device_function() {}
+
+// CHECK: define ptx_kernel{{.*}}global_function
+__global__ void global_function() {
+  // CHECK: call ptx_device{{.*}}device_function
+  device_function();
+}