| // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s |
| // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s |
| |
| #include "Inputs/cuda.h" |
| |
| // This has to be at the top of the file as that's where file-scope |
| // asm ends up. |
| // CHECK-HOST: module asm "file scope asm is host only" |
| // CHECK-DEVICE-NOT: module asm "file scope asm is host only" |
| __asm__("file scope asm is host only"); |
| |
| // CHECK-HOST: constantdata = internal global |
| // CHECK-DEVICE: constantdata = {{(dso_local )?}}externally_initialized global |
| __constant__ char constantdata[256]; |
| |
| // CHECK-HOST: devicedata = internal global |
| // CHECK-DEVICE: devicedata = {{(dso_local )?}}externally_initialized global |
| __device__ char devicedata[256]; |
| |
| // CHECK-HOST: shareddata = internal global |
| // CHECK-DEVICE: shareddata = {{(dso_local )?}}global |
| __shared__ char shareddata[256]; |
| |
| // CHECK-HOST: hostdata = {{(dso_local )?}}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) {} |