diff options
author | Zancanaro; Carlo <czan8762@plang3.cs.usyd.edu.au> | 2012-09-24 09:58:17 +1000 |
---|---|---|
committer | Zancanaro; Carlo <czan8762@plang3.cs.usyd.edu.au> | 2012-09-24 09:58:17 +1000 |
commit | 222e2a7620e6520ffaf4fc4e69d79c18da31542e (patch) | |
tree | 7bfbc05bfa3b41c8f9d2e56d53a0bc3e310df239 /clang/test/CodeGenCUDA | |
parent | 3d206f03985b50beacae843d880bccdc91a9f424 (diff) |
Add the clang library to the repo (with some of my changes, too).
Diffstat (limited to 'clang/test/CodeGenCUDA')
-rw-r--r-- | clang/test/CodeGenCUDA/device-stub.cu | 13 | ||||
-rw-r--r-- | clang/test/CodeGenCUDA/filter-decl.cu | 40 | ||||
-rw-r--r-- | clang/test/CodeGenCUDA/kernel-call.cu | 13 | ||||
-rw-r--r-- | clang/test/CodeGenCUDA/ptx-kernels.cu | 12 |
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(); +} |