1 // REQUIRES: amdgpu-registered-target
2 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
3 // RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV %s
4 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
5 // RUN: -emit-llvm -o - | FileCheck -check-prefix=HOST %s
9 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
10 // RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV-NEG %s
12 #include "Inputs/cuda.h"
14 // Test const var initialized with address of a const var.
15 // Both are promoted to device side.
17 // DEV-DAG: @_ZN5Test1L1aE = internal addrspace(4) constant i32 1
18 // DEV-DAG: @_ZN5Test11B2p1E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*)
19 // DEV-DAG: @_ZN5Test11B2p2E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*)
20 // DEV-DAG: @_ZN5Test12b2E = addrspace(1) externally_initialized global i32 1
21 // HOST-DAG: @_ZN5Test1L1aE = internal constant i32 1
22 // HOST-DAG: @_ZN5Test11B2p1E = constant i32* @_ZN5Test1L1aE
23 // HOST-DAG: @_ZN5Test11B2p2E = internal constant i32* undef
24 // HOST-DAG: @_ZN5Test12b1E = global i32 1
25 // HOST-DAG: @_ZN5Test12b2E = internal global i32 undef
30 static const int *const p1;
31 static __device__ const int *const p2;
33 const int *const B::p1 = &a;
34 __device__ const int *const B::p2 = &a;
35 int b1 = B::p1 == B::p2;
36 __device__ int b2 = B::p1 == B::p2;
39 // Test const var initialized with address of a non-cost var.
40 // Neither is promoted to device side.
42 // DEV-NEG-NOT: @_ZN5Test2L1aE
43 // DEV-NEG-NOT: @_ZN5Test21B1pE
44 // HOST-DAG: @_ZN5Test21aE = global i32 1
45 // HOST-DAG: @_ZN5Test21B1pE = constant i32* @_ZN5Test21aE