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 6 7 // Negative tests. 8 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 11 12 #include "Inputs/cuda.h" 13 14 // Test const var initialized with address of a const var. 15 // Both are promoted to device side. 16 17 // DEV-DAG: @_ZN5Test1L1aE = internal addrspace(4) constant i32 1 18 // DEV-DAG: @_ZN5Test11B2p1E = addrspace(4) externally_initialized constant ptr addrspacecast (ptr addrspace(4) @_ZN5Test1L1aE to ptr) 19 // DEV-DAG: @_ZN5Test11B2p2E = addrspace(4) externally_initialized constant ptr addrspacecast (ptr addrspace(4) @_ZN5Test1L1aE to ptr) 20 // DEV-DAG: @_ZN5Test12b2E = addrspace(1) externally_initialized global i32 1 21 // HOST-DAG: @_ZN5Test1L1aE = internal constant i32 1 22 // HOST-DAG: @_ZN5Test11B2p1E = constant ptr @_ZN5Test1L1aE 23 // HOST-DAG: @_ZN5Test11B2p2E = internal constant ptr undef 24 // HOST-DAG: @_ZN5Test12b1E = global i32 1 25 // HOST-DAG: @_ZN5Test12b2E = internal global i32 undef 26 namespace Test1 { 27 const int a = 1; 28 29 struct B { 30 static const int *const p1; 31 static __device__ const int *const p2; 32 }; 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; 37 } 38 39 // Test const var initialized with address of a non-cost var. 40 // Neither is promoted to device side. 41 42 // DEV-NEG-NOT: @_ZN5Test2L1aE 43 // DEV-NEG-NOT: @_ZN5Test21B1pE 44 // HOST-DAG: @_ZN5Test21aE = global i32 1 45 // HOST-DAG: @_ZN5Test21B1pE = constant ptr @_ZN5Test21aE 46 47 namespace Test2 { 48 int a = 1; 49 50 struct B { 51 static int *const p; 52 }; 53 int *const B::p = &a; 54 } 55