-fno-semantic-interposition: Don't set dso_local on GlobalVariable

[MaskRay]

Summary

clang -fpic -fno-semantic-interposition may set dso_local on variables for -fpic.

GCC folks consider there are 'address interposition' and 'semantic interposition',
and 'disabling semantic interposition' can optimize function calls but
cannot change variable references to use local aliases
(https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100483).

This patch removes dso_local for variables in
clang -fpic -fno-semantic-interposition mode so that the built shared objects can
work with copy relocations.

Example:

// a.c
int var;
int *addr() { return var; }

// old: cannot be interposed
movslq  .Lvar$local(%rip), %rax
// new: can be interposed
movq    var@GOTPCREL(%rip), %rax
movslq  (%rax), %rax

The local alias lowering for GlobalVariables is kept in case there is a
future option allowing local aliases.

[rnk]

lgtm

[thakis]

Looks like this breaks check-clang:
http://45.33.8.238/linux/47049/step_7.txt
http://45.33.8.238/macm1/9850/step_7.txt

PTAL!

clang/lib/CodeGen/CodeGenModule.cpp

static bool shouldAssumeDSOLocal(const CodeGenModule &CGM,
   // If this is not an executable, don't assume anything is local.
   const auto &CGOpts = CGM.getCodeGenOpts();
   llvm::Reloc::Model RM = CGOpts.RelocationModel;
   const auto &LOpts = CGM.getLangOpts();
   if (RM != llvm::Reloc::Static && !LOpts.PIE) {
     // On ELF, if -fno-semantic-interposition is specified and the target
     // supports local aliases, there will be neither CC1
     // -fsemantic-interposition nor -fhalf-no-semantic-interposition. Set
-    // dso_local if using a local alias is preferable (can avoid GOT
-    // indirection).
-    if (!GV->canBenefitFromLocalAlias())
+    // dso_local on the function if using a local alias is preferable (can avoid
+    // PLT indirection).
+    if (!(isa<llvm::Function>(GV) && GV->canBenefitFromLocalAlias()))
       return false;
     return !(CGM.getLangOpts().SemanticInterposition ||
              CGM.getLangOpts().HalfNoSemanticInterposition);
   }
 
   // A definition cannot be preempted from an executable.
   if (!GV->isDeclarationForLinker())
     return true;

clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c

void test_core(void) {
   // CHECK-ASM: vgeg %{{.*}}, 0(%{{.*}},%{{.*}}), 1
 
   vec_scatter_element(vd, vul, ptrd, 0);
   // CHECK-ASM: vsceg %{{.*}}, 0(%{{.*}},%{{.*}}), 0
   vec_scatter_element(vd, vul, ptrd, 1);
   // CHECK-ASM: vsceg %{{.*}}, 0(%{{.*}},%{{.*}}), 1
 
   vd = vec_xl(idx, cptrd);
-  // CHECK-ASM-NEXT: lgfrl   %r3, idx
-  // CHECK-ASM-NEXT: lgrl    %r4, cptrd
-  // CHECK-ASM-NEXT: vl      %v0, 0(%r3,%r4){{$}}
+  // CHECK-ASM-NEXT: lgf     %r5, 0(%r3)
+  // CHECK-ASM-NEXT: lg      %r13, 0(%r4)
+  // CHECK-ASM-NEXT: vl      %v0, 0(%r5,%r13){{$}}
   // CHECK-ASM-NEXT: vst
 
   vd = vec_xld2(idx, cptrd);
-  // CHECK-ASM-NEXT: lgfrl   %r3, idx
-  // CHECK-ASM-NEXT: lgrl    %r4, cptrd
-  // CHECK-ASM-NEXT: vl      %v0, 0(%r3,%r4){{$}}
-  // CHECK-ASM-NEXT: vst
+  // CHECK-ASM:      vst
 
   vec_xst(vd, idx, ptrd);
-  // CHECK-ASM-NEXT: vl
-  // CHECK-ASM-NEXT: lgfrl   %r3, idx
-  // CHECK-ASM-NEXT: lgrl    %r4, ptrd
-  // CHECK-ASM-NEXT: vst     %v0, 0(%r3,%r4){{$}}
 
   vec_xstd2(vd, idx, ptrd);
-  // CHECK-ASM-NEXT: vl
-  // CHECK-ASM-NEXT: lgfrl   %r3, idx
-  // CHECK-ASM-NEXT: lgrl    %r4, ptrd
-  // CHECK-ASM-NEXT: vst     %v0, 0(%r3,%r4){{$}}
 
   vd = vec_splat(vd, 0);
   // CHECK: shufflevector <2 x double> %{{.*}}, <2 x double> poison, <2 x i32> zeroinitializer
   // CHECK-ASM: vrepg
   vd = vec_splat(vd, 1);
   // CHECK: shufflevector <2 x double> %{{.*}}, <2 x double> undef, <2 x i32> <i32 1, i32 1>
   // CHECK-ASM: vrepg
 

clang/test/CodeGen/attr-weakref2.c

 // RUN: %clang_cc1 -emit-llvm -triple i386-linux-gnu -o %t %s
 // RUN: FileCheck --input-file=%t %s
 
 // CHECK: @test1_f = extern_weak global i32
 extern int test1_f;
 static int test1_g __attribute__((weakref("test1_f")));
 int test1_h(void) {
   return test1_g;
 }
 
-// CHECK: @test2_f = dso_local global i32 0, align 4
+// CHECK: @test2_f = global i32 0, align 4
 int test2_f;
 static int test2_g __attribute__((weakref("test2_f")));
 int test2_h(void) {
   return test2_g;
 }
 
 // CHECK: @test3_f = external global i32
 extern int test3_f;
 static int test3_g __attribute__((weakref("test3_f")));
 int test3_foo(void) {
   return test3_f;
 }
 int test3_h(void) {
   return test3_g;
 }
 
-// CHECK: @test4_f = dso_local global i32 0, align 4
+// CHECK: @test4_f = global i32 0, align 4
 extern int test4_f;
 static int test4_g __attribute__((weakref("test4_f")));
 int test4_h(void) {
   return test4_g;
 }
 int test4_f;
 
 // CHECK: @test5_f = external global i32

clang/test/CodeGen/semantic-interposition.c

 /// -fno-semantic-interposition is the default and local aliases (via dso_local) are allowed.
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm -mrelocation-model pic -pic-level 1 %s -o - | FileCheck %s --check-prefixes=CHECK,NOMETADATA
 
 /// -fsemantic-interposition sets a module metadata.
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm -mrelocation-model pic -pic-level 1 -fsemantic-interposition %s -o - | FileCheck %s --check-prefixes=PREEMPT,METADATA
 
 /// Traditional half-baked behavior: interprocedural optimizations are allowed
 /// but local aliases are not used.
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm -mrelocation-model pic -pic-level 1 -fhalf-no-semantic-interposition %s -o - | FileCheck %s --check-prefixes=PREEMPT,NOMETADATA
 
-// CHECK: @var = dso_local global i32 0, align 4
+// CHECK: @var = global i32 0, align 4
 // CHECK: @ext_var = external global i32, align 4
 // CHECK: @ifunc = ifunc i32 (), bitcast (i8* ()* @ifunc_resolver to i32 ()*)
 // CHECK: define dso_local i32 @func()
 // CHECK: declare i32 @ext()
 
 // PREEMPT: @var = global i32 0, align 4
 // PREEMPT: @ext_var = external global i32, align 4
 // PREEMPT: @ifunc = ifunc i32 (), bitcast (i8* ()* @ifunc_resolver to i32 ()*)

clang/test/CodeGenCUDA/device-stub.cu

 // RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -aux-triple amdgcn -emit-llvm %s \
 // RUN:     -fcuda-include-gpubinary %t -o - -x hip\
 // RUN:   | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=ALL,WIN
 
 #include "Inputs/cuda.h"
 
 #ifndef NOGLOBALS
 // NORDC-DAG: @device_var = internal global i32
-// RDC-DAG: @device_var = dso_local global i32
+// RDC-DAG: @device_var = global i32
 // WIN-DAG: @"?device_var@@3HA" = internal global i32
 __device__ int device_var;
 
 // NORDC-DAG: @constant_var = internal global i32
-// RDC-DAG: @constant_var = dso_local global i32
+// RDC-DAG: @constant_var = global i32
 // WIN-DAG: @"?constant_var@@3HA" = internal global i32
 __constant__ int constant_var;
 
 // NORDC-DAG: @shared_var = internal global i32
-// RDC-DAG: @shared_var = dso_local global i32
+// RDC-DAG: @shared_var = global i32
 // WIN-DAG: @"?shared_var@@3HA" = internal global i32
 __shared__ int shared_var;
 
 // Make sure host globals don't get internalized...
 // LNX-DAG: @host_var ={{.*}} global i32
 // WIN-DAG: @"?host_var@@3HA" = dso_local global i32
 int host_var;
 // ... and that extern vars remain external.
 // LNX-DAG: @ext_host_var = external global i32
 // WIN-DAG: @"?ext_host_var@@3HA" = external dso_local global i32
 extern int ext_host_var;
 
 // external device-side variables -> extern references to their shadows.
 // LNX-DAG: @ext_device_var = external global i32
 // WIN-DAG: @"?ext_device_var@@3HA" = external dso_local global i32
 extern __device__ int ext_device_var;
 // LNX-DAG: @ext_device_var = external global i32
 // WIN-DAG: @"?ext_constant_var@@3HA" = external dso_local global i32
 extern __constant__ int ext_constant_var;
 
 // external device-side variables with definitions should generate
 // definitions for the shadows.
 // NORDC-DAG: @ext_device_var_def = internal global i32 undef,
-// RDC-DAG: @ext_device_var_def = dso_local global i32 undef,
+// RDC-DAG: @ext_device_var_def = global i32 undef,
 // WIN-DAG: @"?ext_device_var_def@@3HA" = internal global i32 undef
 extern __device__ int ext_device_var_def;
 __device__ int ext_device_var_def = 1;
 // NORDC-DAG: @ext_device_var_def = internal global i32 undef,
-// RDC-DAG: @ext_device_var_def = dso_local global i32 undef,
+// RDC-DAG: @ext_device_var_def = global i32 undef,
 // WIN-DAG: @"?ext_constant_var_def@@3HA" = internal global i32 undef
 __constant__ int ext_constant_var_def = 2;
 
 #if __cplusplus > 201402L
 // NORDC17: @inline_var = internal global i32 undef, comdat, align 4{{$}}
 // RDC17: @inline_var = linkonce_odr global i32 undef, comdat, align 4{{$}}
 // NORDC17-NOT: @inline_var2 =
 // RDC17-NOT: @inline_var2 =

clang/test/CodeGenCUDA/device-var-linkage.cu

 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
 // RUN:   -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=DEV,NORDC %s
 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device \
 // RUN:   -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=DEV,RDC %s
 // RUN: %clang_cc1 -triple nvptx \
 // RUN:   -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=HOST,NORDC-H %s
 // RUN: %clang_cc1 -triple nvptx \
 // RUN:   -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=HOST,RDC-H %s
 
 #include "Inputs/cuda.h"
 
-// DEV-DAG: @v1 = dso_local addrspace(1) externally_initialized global i32 0
+// DEV-DAG: @v1 = addrspace(1) externally_initialized global i32 0
 // NORDC-H-DAG: @v1 = internal global i32 undef
-// RDC-H-DAG: @v1 = dso_local global i32 undef
+// RDC-H-DAG: @v1 = global i32 undef
 __device__ int v1;
-// DEV-DAG: @v2 = dso_local addrspace(4) externally_initialized global i32 0
+// DEV-DAG: @v2 = addrspace(4) externally_initialized global i32 0
 // NORDC-H-DAG: @v2 = internal global i32 undef
-// RDC-H-DAG: @v2 = dso_local global i32 undef
+// RDC-H-DAG: @v2 = global i32 undef
 __constant__ int v2;
-// DEV-DAG: @v3 = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// DEV-DAG: @v3 = addrspace(1) externally_initialized global i32 addrspace(1)* null
 // NORDC-H-DAG: @v3 = internal externally_initialized global i32* null
-// RDC-H-DAG: @v3 = dso_local externally_initialized global i32* null
+// RDC-H-DAG: @v3 = externally_initialized global i32* null
 __managed__ int v3;
 
 // DEV-DAG: @ev1 = external addrspace(1) global i32
 // HOST-DAG: @ev1 = external global i32
 extern __device__ int ev1;
 // DEV-DAG: @ev2 = external addrspace(4) global i32
 // HOST-DAG: @ev2 = external global i32
 extern __constant__ int ev2;
 // DEV-DAG: @ev3 = external addrspace(1) externally_initialized global i32 addrspace(1)*
 // HOST-DAG: @ev3 = external externally_initialized global i32*
 extern __managed__ int ev3;
 
-// NORDC-DAG: @_ZL3sv1 = dso_local addrspace(1) externally_initialized global i32 0
-// RDC-DAG: @_ZL3sv1.static.[[HASH:.*]] = dso_local addrspace(1) externally_initialized global i32 0
+// NORDC-DAG: @_ZL3sv1 = addrspace(1) externally_initialized global i32 0
+// RDC-DAG: @_ZL3sv1.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
 // HOST-DAG: @_ZL3sv1 = internal global i32 undef
 static __device__ int sv1;
-// NORDC-DAG: @_ZL3sv2 = dso_local addrspace(4) externally_initialized global i32 0
-// RDC-DAG: @_ZL3sv2.static.[[HASH]] = dso_local addrspace(4) externally_initialized global i32 0
+// NORDC-DAG: @_ZL3sv2 = addrspace(4) externally_initialized global i32 0
+// RDC-DAG: @_ZL3sv2.static.[[HASH]] = addrspace(4) externally_initialized global i32 0
 // HOST-DAG: @_ZL3sv2 = internal global i32 undef
 static __constant__ int sv2;
-// NORDC-DAG: @_ZL3sv3 = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
-// RDC-DAG: @_ZL3sv3.static.[[HASH]] = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// NORDC-DAG: @_ZL3sv3 = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-DAG: @_ZL3sv3.static.[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
 // HOST-DAG: @_ZL3sv3 = internal externally_initialized global i32* null
 static __managed__ int sv3;
 
 __device__ __host__ int work(int *x);
 
 __device__ __host__ int fun1() {
   return work(&ev1) + work(&ev2) + work(&ev3) + work(&sv1) + work(&sv2) + work(&sv3);
 }

clang/test/CodeGenCUDA/managed-var.cu

 // RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s
 
 #include "Inputs/cuda.h"
 
 struct vec {
   float x,y,z;
 };
 
-// DEV-DAG: @x.managed = dso_local addrspace(1) externally_initialized global i32 1, align 4
-// DEV-DAG: @x = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// DEV-DAG: @x.managed = addrspace(1) externally_initialized global i32 1, align 4
+// DEV-DAG: @x = addrspace(1) externally_initialized global i32 addrspace(1)* null
 // NORDC-DAG: @x.managed = internal global i32 1
-// RDC-DAG: @x.managed = dso_local global i32 1
+// RDC-DAG: @x.managed = global i32 1
 // NORDC-DAG: @x = internal externally_initialized global i32* null
-// RDC-DAG: @x = dso_local externally_initialized global i32* null
+// RDC-DAG: @x = externally_initialized global i32* null
 // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00"
 __managed__ int x = 1;
 
-// DEV-DAG: @v.managed = dso_local addrspace(1) externally_initialized global [100 x %struct.vec] zeroinitializer, align 4
-// DEV-DAG: @v = dso_local addrspace(1) externally_initialized global [100 x %struct.vec] addrspace(1)* null
+// DEV-DAG: @v.managed = addrspace(1) externally_initialized global [100 x %struct.vec] zeroinitializer, align 4
+// DEV-DAG: @v = addrspace(1) externally_initialized global [100 x %struct.vec] addrspace(1)* null
 __managed__ vec v[100];
 
-// DEV-DAG: @v2.managed = dso_local addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> <{ %struct.vec { float 1.000000e+00, float 1.000000e+00, float 1.000000e+00 }, [99 x %struct.vec] zeroinitializer }>, align 4
-// DEV-DAG: @v2 = dso_local addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> addrspace(1)* null
+// DEV-DAG: @v2.managed = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> <{ %struct.vec { float 1.000000e+00, float 1.000000e+00, float 1.000000e+00 }, [99 x %struct.vec] zeroinitializer }>, align 4
+// DEV-DAG: @v2 = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> addrspace(1)* null
 __managed__ vec v2[100] = {{1, 1, 1}};
 
 // DEV-DAG: @ex.managed = external addrspace(1) global i32, align 4
 // DEV-DAG: @ex = external addrspace(1) externally_initialized global i32 addrspace(1)*
 // HOST-DAG: @ex.managed = external global i32
 // HOST-DAG: @ex = external externally_initialized global i32*
 extern __managed__ int ex;
 
-// NORDC-D-DAG: @_ZL2sx.managed = dso_local addrspace(1) externally_initialized global i32 1, align 4
-// NORDC-D-DAG: @_ZL2sx = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
-// RDC-D-DAG: @_ZL2sx.static.[[HASH:.*]].managed = dso_local addrspace(1) externally_initialized global i32 1, align 4
-// RDC-D-DAG: @_ZL2sx.static.[[HASH]] = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// NORDC-D-DAG: @_ZL2sx.managed = addrspace(1) externally_initialized global i32 1, align 4
+// NORDC-D-DAG: @_ZL2sx = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-D-DAG: @_ZL2sx.static.[[HASH:.*]].managed = addrspace(1) externally_initialized global i32 1, align 4
+// RDC-D-DAG: @_ZL2sx.static.[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
 // HOST-DAG: @_ZL2sx.managed = internal global i32 1
 // HOST-DAG: @_ZL2sx = internal externally_initialized global i32* null
 // NORDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx\00"
 // RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH:.*]]\00"
 
-// POSTFIX:  @_ZL2sx.static.[[HASH:.*]] = dso_local addrspace(1) externally_initialized global i32 addrspace(1)* null
+// POSTFIX:  @_ZL2sx.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
 // POSTFIX: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH]]\00"
 static __managed__ int sx = 1;
 
 // DEV-DAG: @llvm.compiler.used
 // DEV-SAME-DAG: @x.managed
 // DEV-SAME-DAG: @x
 // DEV-SAME-DAG: @v.managed
 // DEV-SAME-DAG: @v

clang/test/CodeGenCUDA/static-device-var-rdc.cu

 // Test function scope static device variable, which should not be externalized.
 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
 
 
 // HOST-DAG: @_ZL1x = internal global i32 undef
 // HOST-DAG: @_ZL1y = internal global i32 undef
 
 // Test normal static device variables
-// INT-DEV-DAG: @_ZL1x = dso_local addrspace(1) externally_initialized global i32 0
+// INT-DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0
 // INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00"
 
 // Test externalized static device variables
-// EXT-DEV-DAG: @_ZL1x.static.[[HASH:.*]] = dso_local addrspace(1) externally_initialized global i32 0
+// EXT-DEV-DAG: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
 // EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH:.*]]\00"
 
-// POSTFIX: @_ZL1x.static.[[HASH:.*]] = dso_local addrspace(1) externally_initialized global i32 0
+// POSTFIX: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
 // POSTFIX: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH]]\00"
 
 static __device__ int x;
 
 // Test static device variables not used by host code should not be externalized
 // DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 0
 
 static __device__ int x2;
 
 // Test normal static device variables
-// INT-DEV-DAG: @_ZL1y = dso_local addrspace(4) externally_initialized global i32 0
+// INT-DEV-DAG: @_ZL1y = addrspace(4) externally_initialized global i32 0
 // INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00"
 
 // Test externalized static device variables
-// EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = dso_local addrspace(4) externally_initialized global i32 0
+// EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = addrspace(4) externally_initialized global i32 0
 // EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y.static.[[HASH]]\00"
 
 static __constant__ int y;
 
 // Test static host variable, which should not be externalized nor registered.
 // HOST-DAG: @_ZL1z = internal global i32 0
 // DEV-NEG-NOT: @_ZL1z
 static int z;

clang/test/OpenMP/declare_target_only_one_side_compilation.cpp

 // DEVICE-DAG: @G3 = hidden {{.*}}global i32 0, align 4
 // DEVICE-DAG: @_ZL2G4 = internal {{.*}}global i32 0, align 4
 // DEVICE-DAG: @G5 = hidden {{.*}}global i32 0, align 4
 // DEVICE-DAG: @_ZL2G6 = internal {{.*}}global i32 0, align 4
 // DEVICE-NOT: ref
 // DEVICE-NOT: llvm.used
 // DEVICE-NOT: omp_offload
 
-// HOST-DAG: @G7 = dso_local global i32 0, align 4
+// HOST-DAG: @G7 = global i32 0, align 4
 // HOST-DAG: @_ZL2G8 = internal global i32 0, align 4
-// HOST-DAG: @G9 = dso_local global i32 0, align 4
+// HOST-DAG: @G9 = global i32 0, align 4
 // HOST-DAG: @_ZL3G10 = internal global i32 0, align 4
-// HOST-DAG: @G11 = dso_local global i32 0, align 4
+// HOST-DAG: @G11 = global i32 0, align 4
 // HOST-DAG: @_ZL3G12 = internal global i32 0, align 4