55 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \3// RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV %s4// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \5// RUN: -emit-llvm -o - | FileCheck -check-prefix=HOST %s6 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 %s11 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 118// 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 121// HOST-DAG: @_ZN5Test1L1aE = internal constant i32 122// HOST-DAG: @_ZN5Test11B2p1E = constant ptr @_ZN5Test1L1aE23// HOST-DAG: @_ZN5Test11B2p2E = internal constant ptr undef24// HOST-DAG: @_ZN5Test12b1E = global i32 125// HOST-DAG: @_ZN5Test12b2E = internal global i32 undef26namespace Test1 {27const int a = 1;28 29struct B {30 static const int *const p1;31 static __device__ const int *const p2;32};33const int *const B::p1 = &a;34__device__ const int *const B::p2 = &a;35int 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: @_ZN5Test2L1aE43// DEV-NEG-NOT: @_ZN5Test21B1pE44// HOST-DAG: @_ZN5Test21aE = global i32 145// HOST-DAG: @_ZN5Test21B1pE = constant ptr @_ZN5Test21aE46 47namespace Test2 {48int a = 1;49 50struct B {51 static int *const p;52};53int *const B::p = &a;54}55