summaryrefslogtreecommitdiff
path: root/clang/test/SemaCUDA
diff options
context:
space:
mode:
authorCarlo Zancanaro <carlo@pc-4w14-0.cs.usyd.edu.au>2012-10-15 17:10:06 +1100
committerCarlo Zancanaro <carlo@pc-4w14-0.cs.usyd.edu.au>2012-10-15 17:10:06 +1100
commitbe1de4be954c80875ad4108e0a33e8e131b2f2c0 (patch)
tree1fbbecf276bf7c7bdcbb4dd446099d6d90eaa516 /clang/test/SemaCUDA
parentc4626a62754862d20b41e8a46a3574264ea80e6d (diff)
parentf1bd2e48c5324d3f7cda4090c87f8a5b6f463ce2 (diff)
Merge branch 'master' of ssh://bitbucket.org/czan/honours
Diffstat (limited to 'clang/test/SemaCUDA')
-rw-r--r--clang/test/SemaCUDA/config-type.cu3
-rw-r--r--clang/test/SemaCUDA/cuda.h19
-rw-r--r--clang/test/SemaCUDA/function-target.cu44
-rw-r--r--clang/test/SemaCUDA/kernel-call.cu26
-rw-r--r--clang/test/SemaCUDA/qualifiers.cu8
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;
+}