215 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \3// RUN: -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \4// RUN: -cuid=123 | FileCheck -check-prefix=DEV %s5// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \6// RUN: -std=c++17 -O3 -emit-llvm -o - -cuid=123 | FileCheck -check-prefix=HOST %s7 8// Negative tests.9 10// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \11// RUN: -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \12// RUN: | FileCheck -check-prefix=DEV-NEG %s13// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \14// RUN: -std=c++17 -O3 -emit-llvm -o - | FileCheck -check-prefix=HOST-NEG %s15 16#include "Inputs/cuda.h"17 18// DEV-DAG: @v119__device__ int v1;20 21// DEV-DAG: @v222__constant__ int v2;23 24// Check device variables used by neither host nor device functioins are not kept.25 26// DEV-NEG-NOT: @_ZL2v327static __device__ int v3;28 29// Check device variables used by host functions are kept.30 31// DEV-DAG: @u132__device__ int u1;33 34// DEV-DAG: @u235__constant__ int u2;36 37// Check host-used static device var is in llvm.compiler.used.38// DEV-DAG: @_ZL2u339static __device__ int u3;40 41// Check device-used static device var is emitted but is not in llvm.compiler.used.42// DEV-DAG: @_ZL2u443static __device__ int u4;44 45// Check device variables with used attribute are always kept.46// DEV-DAG: @u547__device__ __attribute__((used)) int u5;48 49// Test external device variable ODR-used by host code is not emitted or registered.50// DEV-NEG-NOT: @ext_var51extern __device__ int ext_var;52 53// DEV-DAG: @inline_var = linkonce_odr addrspace(1) externally_initialized global i32 054__device__ inline int inline_var;55 56template<typename T>57using func_t = T (*) (T, T);58 59template <typename T>60__device__ T add_func (T x, T y)61{62 return x + y;63}64 65// DEV-DAG: @_Z10p_add_funcIiE = linkonce_odr addrspace(1) externally_initialized global ptr @_Z8add_funcIiET_S0_S0_66template <typename T>67__device__ func_t<T> p_add_func = add_func<T>;68 69// Check non-constant constexpr variables ODR-used by host code only is not emitted.70// DEV-NEG-NOT: constexpr_var1a71// DEV-NEG-NOT: constexpr_var1b72constexpr int constexpr_var1a = 1;73inline constexpr int constexpr_var1b = 1;74 75// Check constant constexpr variables ODR-used by host code only.76// Device-side constexpr variables accessed by host code should be externalized and kept.77// DEV-DAG: @_ZL15constexpr_var2a = addrspace(4) externally_initialized constant i32 278// DEV-DAG: @constexpr_var2b = linkonce_odr addrspace(4) externally_initialized constant i32 279__constant__ constexpr int constexpr_var2a = 2;80inline __constant__ constexpr int constexpr_var2b = 2;81 82void use(func_t<int> p);83__host__ __device__ void use(const int *p);84 85// Check static device variable in host function.86// DEV-DAG: @_ZZ4fun1vE11static_var1 = addrspace(1) externally_initialized global i32 387void fun1() {88 static __device__ int static_var1 = 3;89 use(&u1);90 use(&u2);91 use(&u3);92 use(&ext_var);93 use(&inline_var);94 use(p_add_func<int>);95 use(&constexpr_var1a);96 use(&constexpr_var1b);97 use(&constexpr_var2a);98 use(&constexpr_var2b);99 use(&static_var1);100}101 102// Check static variable in host device function.103// DEV-DAG: @_ZZ4fun2vE11static_var2 = internal addrspace(1) global i32 4104// DEV-DAG: @_ZZ4fun2vE11static_var3 = addrspace(1) global i32 4105__host__ __device__ void fun2() {106 static int static_var2 = 4;107 static __device__ int static_var3 = 4;108 use(&static_var2);109 use(&static_var3);110}111 112__global__ void kern1(int **x) {113 *x = &u4;114 fun2();115}116 117// Check static variables of lambda functions.118 119// Lambda functions are implicit host device functions.120// Default static variables in lambda functions should be treated121// as host variables on host side, therefore should not be forced122// to be emitted on device.123 124// DEV-DAG: @_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var2 = addrspace(1) externally_initialized global i32 5125// DEV-NEG-NOT: @_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var1126namespace TestStaticVarInLambda {127class A {128public:129 A(char *);130};131void fun() {132 (void) [](char *c) {133 static A var1(c);134 static __device__ int var2 = 5;135 (void) var1;136 (void) var2;137 };138}139}140 141// Check implicit constant variable ODR-used by host code is not emitted.142 143// AST contains instantiation of al<ar>, which triggers AST instantiation144// of x::al<ar>::am, which triggers AST instatiation of x::ap<ar>,145// which triggers AST instantiation of aw<ar>::c, which has type146// ar. ar has base class x which has member ah. x::ah is initialized147// with function pointer pointing to ar:as, which returns an object148// of type ou. The constexpr aw<ar>::c is an implicit constant variable149// which is ODR-used by host function x::ap<ar>. An incorrect implementation150// will force aw<ar>::c to be emitted on device side, which will trigger151// emit of x::as and further more ctor of ou and variable o.152// The ODR-use of aw<ar>::c in x::ap<ar> should be treated as a host variable153// instead of device variable.154 155// DEV-NEG-NOT: _ZN16TestConstexprVar1oE156namespace TestConstexprVar {157char o;158class ou {159public:160 ou(char) { __builtin_strlen(&o); }161};162template < typename ao > struct aw { static constexpr ao c; };163class x {164protected:165 typedef ou (*y)(const x *);166 constexpr x(y ag) : ah(ag) {}167 template < bool * > struct ak;168 template < typename > struct al {169 static bool am;170 static ak< &am > an;171 };172 template < typename ao > static x ap() { (void)aw< ao >::c; return x(nullptr); }173 y ah;174};175template < typename ao > bool x::al< ao >::am(&ap< ao >);176class ar : x {177public:178 constexpr ar() : x(as) {}179 static ou as(const x *) { return 0; }180 al< ar > av;181};182}183 184// Check the exact list of variables to ensure @_ZL2u4 is not among them.185// DEV: @llvm.compiler.used = {{[^@]*}} @_Z10p_add_funcIiE186// DEV-SAME: {{^[^@]*}} @_ZL15constexpr_var2a187// DEV-SAME: {{^[^@]*}} @_ZL2u3188// DEV-SAME: {{^[^@]*}} @_ZZ4fun1vE11static_var1189// DEV-SAME: {{^[^@]*}} @_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var2190// DEV-SAME: {{^[^@]*}} @__hip_cuid_{{[0-9a-f]+}}191// DEV-SAME: {{^[^@]*}} @constexpr_var2b192// DEV-SAME: {{^[^@]*}} @inline_var193// DEV-SAME: {{^[^@]*}} @u1194// DEV-SAME: {{^[^@]*}} @u2195// DEV-SAME: {{^[^@]*}} @u5196// DEV-SAME: {{^[^@]*$}}197 198// HOST-DAG: hipRegisterVar{{.*}}@u1199// HOST-DAG: hipRegisterVar{{.*}}@u2200// HOST-DAG: hipRegisterVar{{.*}}@_ZL2u3201// HOST-DAG: hipRegisterVar{{.*}}@constexpr_var2b202// HOST-DAG: hipRegisterVar{{.*}}@u5203// HOST-DAG: hipRegisterVar{{.*}}@inline_var204// HOST-DAG: hipRegisterVar{{.*}}@_Z10p_add_funcIiE205// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZ4fun1vE11static_var1206// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZ4fun2vE11static_var2207// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZ4fun2vE11static_var3208// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var2209// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var1210// HOST-NEG-NOT: hipRegisterVar{{.*}}@ext_var211// HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZL2u4212// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1a213// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1b214// HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var2a215