1 | // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s |
2 | // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s |
3 | |
4 | #include "Inputs/cuda.h" |
5 | |
6 | // This has to be at the top of the file as that's where file-scope |
7 | // asm ends up. |
8 | // CHECK-HOST: module asm "file scope asm is host only" |
9 | // CHECK-DEVICE-NOT: module asm "file scope asm is host only" |
10 | __asm__("file scope asm is host only"); |
11 | |
12 | // CHECK-HOST: constantdata = internal global |
13 | // CHECK-DEVICE: constantdata = {{(dso_local )?}}externally_initialized global |
14 | __constant__ char constantdata[256]; |
15 | |
16 | // CHECK-HOST: devicedata = internal global |
17 | // CHECK-DEVICE: devicedata = {{(dso_local )?}}externally_initialized global |
18 | __device__ char devicedata[256]; |
19 | |
20 | // CHECK-HOST: shareddata = internal global |
21 | // CHECK-DEVICE: shareddata = {{(dso_local )?}}global |
22 | __shared__ char shareddata[256]; |
23 | |
24 | // CHECK-HOST: hostdata = {{(dso_local )?}}global |
25 | // CHECK-DEVICE-NOT: hostdata = global |
26 | char hostdata[256]; |
27 | |
28 | // CHECK-HOST: define{{.*}}implicithostonlyfunc |
29 | // CHECK-DEVICE-NOT: define{{.*}}implicithostonlyfunc |
30 | void implicithostonlyfunc(void) {} |
31 | |
32 | // CHECK-HOST: define{{.*}}explicithostonlyfunc |
33 | // CHECK-DEVICE-NOT: define{{.*}}explicithostonlyfunc |
34 | __host__ void explicithostonlyfunc(void) {} |
35 | |
36 | // CHECK-HOST-NOT: define{{.*}}deviceonlyfunc |
37 | // CHECK-DEVICE: define{{.*}}deviceonlyfunc |
38 | __device__ void deviceonlyfunc(void) {} |
39 | |
40 | // CHECK-HOST: define{{.*}}hostdevicefunc |
41 | // CHECK-DEVICE: define{{.*}}hostdevicefunc |
42 | __host__ __device__ void hostdevicefunc(void) {} |
43 | |
44 | // CHECK-HOST: define{{.*}}globalfunc |
45 | // CHECK-DEVICE: define{{.*}}globalfunc |
46 | __global__ void globalfunc(void) {} |
47 | |