133 lines · plain
1// REQUIRES: x86-registered-target2// REQUIRES: amdgpu-registered-target3 4// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \5// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s6 7// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \8// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST %s9 10// Negative tests.11 12// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \13// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s14 15// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \16// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST-NEG %s17 18#include "Inputs/cuda.h"19 20// Test function scope static device variable, which should not be externalized.21// DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 122// DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 4223// DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 4324 25// Check a static device variable referenced by host function is externalized.26// DEV-DAG: @_ZL1x ={{.*}} addrspace(1) externally_initialized global i32 027// HOST-DAG: @_ZL1x = internal global i32 undef28// HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00"29 30static __device__ int x;31 32// Check a static device variables referenced only by device functions and kernels33// is not externalized.34// DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 035static __device__ int x2;36 37// Check a static device variable referenced by host device function is externalized.38// DEV-DAG: @_ZL2x3 ={{.*}} addrspace(1) externally_initialized global i32 039static __device__ int x3;40 41// Check a static device variable referenced in file scope is externalized.42// DEV-DAG: @_ZL2x4 ={{.*}} addrspace(1) externally_initialized global i32 043static __device__ int x4;44int& x4_ref = x4;45 46// Check a static device variable in anonymous namespace.47// DEV-DAG: @_ZN12_GLOBAL__N_12x5E ={{.*}} addrspace(1) externally_initialized global i32 048namespace {49static __device__ int x5;50}51 52// Check a static constant variable referenced by host is externalized.53// DEV-DAG: @_ZL1y ={{.*}} addrspace(4) externally_initialized constant i32 054// HOST-DAG: @_ZL1y = internal global i32 undef55// HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00"56 57static __constant__ int y;58 59// Test static host variable, which should not be externalized nor registered.60// HOST-DAG: @_ZL1z = internal global i32 061// DEV-NEG-NOT: @_ZL1z62static int z;63 64// Test implicit static constant variable, which should not be externalized.65// HOST-DAG: @_ZL2z2 = internal constant i32 45666// DEV-DAG: @_ZL2z2 = internal addrspace(4) constant i32 45667 68static constexpr int z2 = 456;69 70// Test static device variable in inline function, which should not be71// externalized nor registered.72// DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat73 74// Check a static device variable referenced by host function only is externalized.75// DEV-DAG: @_ZL1w ={{.*}} addrspace(1) externally_initialized global i32 076// HOST-DAG: @_ZL1w = internal global i32 undef77// HOST-DAG: @[[DEVNAMEW:[0-9]+]] = {{.*}}c"_ZL1w\00"78 79static __device__ int w;80 81// Test non-ODR-use of static device var should not be emitted or registered.82// DEV-NEG-NOT: @_ZL1u83// HOST-NEG-NOT: @_ZL1u84 85static __device__ int u;86 87inline __device__ void devfun(const int ** b) {88 const static int p = 2;89 b[0] = &p;90 b[1] = &x2;91}92 93__global__ void kernel(int *a, const int **b) {94 const static int w = 1;95 const static __constant__ int local_static_constant = 42;96 const static __device__ int local_static_device = 43;97 a[0] = x;98 a[1] = y;99 a[2] = x2;100 a[3] = x3;101 a[4] = x4;102 a[5] = x5;103 a[6] = sizeof(u);104 b[0] = &w;105 b[1] = &z2;106 b[2] = &local_static_constant;107 b[3] = &local_static_device;108 devfun(b);109}110 111__host__ __device__ void hdf(int *a) {112 a[0] = x3;113}114 115int* getDeviceSymbol(int *x);116 117void foo(const int **a) {118 getDeviceSymbol(&x);119 getDeviceSymbol(&x5);120 getDeviceSymbol(&y);121 getDeviceSymbol(&w);122 z = 123;123 a[0] = &z2;124 decltype(u) tmp;125}126 127// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x, {{.*}}@[[DEVNAMEX]]128// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y, {{.*}}@[[DEVNAMEY]]129// HOST-DAG: __hipRegisterVar({{.*}}@_ZL1w, {{.*}}@[[DEVNAMEW]]130// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL1u131// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w132// HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p133