| // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \ |
| // RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s |
| |
| // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ |
| // RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST %s |
| |
| // Negative tests. |
| |
| // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \ |
| // RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s |
| |
| #include "Inputs/cuda.h" |
| |
| template <class T> |
| class A { |
| static int h_member; |
| __device__ static int d_member; |
| __constant__ static int c_member; |
| __managed__ static int m_member; |
| const static int const_member = 0; |
| }; |
| |
| template <class T> |
| int A<T>::h_member; |
| |
| template <class T> |
| __device__ int A<T>::d_member; |
| |
| template <class T> |
| __constant__ int A<T>::c_member; |
| |
| template <class T> |
| __managed__ int A<T>::m_member; |
| |
| template <class T> |
| const int A<T>::const_member; |
| |
| template class A<int>; |
| |
| //DEV-DAG: @_ZN1AIiE8d_memberE = internal addrspace(1) global i32 0, comdat, align 4 |
| //DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) constant i32 0, comdat, align 4 |
| //DEV-DAG: @_ZN1AIiE8m_memberE = internal addrspace(1) externally_initialized global ptr addrspace(1) null |
| //DEV-DAG: @_ZN1AIiE12const_memberE = internal addrspace(4) constant i32 0, comdat, align 4 |
| //DEV-NEG-NOT: @_ZN1AIiE8h_memberE |
| |
| //HOST-DAG: @_ZN1AIiE8h_memberE = weak_odr global i32 0, comdat, align 4 |
| //HOST-DAG: @_ZN1AIiE8d_memberE = internal global i32 undef, comdat, align 4 |
| //HOST-DAG: @_ZN1AIiE8c_memberE = internal global i32 undef, comdat, align 4 |
| //HOST-DAG: @_ZN1AIiE8m_memberE = internal externally_initialized global ptr null |
| //HOST-DAG: @_ZN1AIiE12const_memberE = weak_odr constant i32 0, comdat, align 4 |
