| // REQUIRES: amdgpu-registered-target |
| // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \ |
| // RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV %s |
| // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-unknown-linux-gnu -x hip %s \ |
| // RUN: -emit-llvm -o - | FileCheck -check-prefix=HOST %s |
| |
| // Negative tests. |
| |
| // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \ |
| // RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV-NEG %s |
| |
| #include "Inputs/cuda.h" |
| |
| // Test const var initialized with address of a const var. |
| // Both are promoted to device side. |
| |
| // DEV-DAG: @_ZN5Test1L1aE = internal addrspace(4) constant i32 1 |
| // DEV-DAG: @_ZN5Test11B2p1E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*) |
| // DEV-DAG: @_ZN5Test11B2p2E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*) |
| // DEV-DAG: @_ZN5Test12b2E = addrspace(1) externally_initialized global i32 1 |
| // HOST-DAG: @_ZN5Test1L1aE = internal constant i32 1 |
| // HOST-DAG: @_ZN5Test11B2p1E = constant i32* @_ZN5Test1L1aE |
| // HOST-DAG: @_ZN5Test11B2p2E = internal constant i32* undef |
| // HOST-DAG: @_ZN5Test12b1E = global i32 1 |
| // HOST-DAG: @_ZN5Test12b2E = internal global i32 undef |
| namespace Test1 { |
| const int a = 1; |
| |
| struct B { |
| static const int *const p1; |
| static __device__ const int *const p2; |
| }; |
| const int *const B::p1 = &a; |
| __device__ const int *const B::p2 = &a; |
| int b1 = B::p1 == B::p2; |
| __device__ int b2 = B::p1 == B::p2; |
| } |
| |
| // Test const var initialized with address of a non-cost var. |
| // Neither is promoted to device side. |
| |
| // DEV-NEG-NOT: @_ZN5Test2L1aE |
| // DEV-NEG-NOT: @_ZN5Test21B1pE |
| // HOST-DAG: @_ZN5Test21aE = global i32 1 |
| // HOST-DAG: @_ZN5Test21B1pE = constant i32* @_ZN5Test21aE |
| |
| namespace Test2 { |
| int a = 1; |
| |
| struct B { |
| static int *const p; |
| }; |
| int *const B::p = &a; |
| } |