diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1007,9 +1007,9 @@ // 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(GV) && GV->canBenefitFromLocalAlias())) return false; return !(CGM.getLangOpts().SemanticInterposition || CGM.getLangOpts().HalfNoSemanticInterposition); diff --git a/clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c b/clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c --- a/clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c +++ b/clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c @@ -66,28 +66,17 @@ // 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 diff --git a/clang/test/CodeGen/attr-weakref2.c b/clang/test/CodeGen/attr-weakref2.c --- a/clang/test/CodeGen/attr-weakref2.c +++ b/clang/test/CodeGen/attr-weakref2.c @@ -8,7 +8,7 @@ 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) { @@ -25,7 +25,7 @@ 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) { diff --git a/clang/test/CodeGen/semantic-interposition.c b/clang/test/CodeGen/semantic-interposition.c --- a/clang/test/CodeGen/semantic-interposition.c +++ b/clang/test/CodeGen/semantic-interposition.c @@ -8,7 +8,7 @@ /// 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() diff --git a/clang/test/CodeGenCUDA/device-stub.cu b/clang/test/CodeGenCUDA/device-stub.cu --- a/clang/test/CodeGenCUDA/device-stub.cu +++ b/clang/test/CodeGenCUDA/device-stub.cu @@ -61,17 +61,17 @@ #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; @@ -95,12 +95,12 @@ // 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; diff --git a/clang/test/CodeGenCUDA/device-var-linkage.cu b/clang/test/CodeGenCUDA/device-var-linkage.cu --- a/clang/test/CodeGenCUDA/device-var-linkage.cu +++ b/clang/test/CodeGenCUDA/device-var-linkage.cu @@ -13,17 +13,17 @@ #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 @@ -36,16 +36,16 @@ // 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; diff --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu --- a/clang/test/CodeGenCUDA/managed-var.cu +++ b/clang/test/CodeGenCUDA/managed-var.cu @@ -27,21 +27,21 @@ 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 @@ -50,16 +50,16 @@ // 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; diff --git a/clang/test/CodeGenCUDA/static-device-var-rdc.cu b/clang/test/CodeGenCUDA/static-device-var-rdc.cu --- a/clang/test/CodeGenCUDA/static-device-var-rdc.cu +++ b/clang/test/CodeGenCUDA/static-device-var-rdc.cu @@ -51,14 +51,14 @@ // 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; @@ -69,11 +69,11 @@ 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; diff --git a/clang/test/OpenMP/declare_target_only_one_side_compilation.cpp b/clang/test/OpenMP/declare_target_only_one_side_compilation.cpp --- a/clang/test/OpenMP/declare_target_only_one_side_compilation.cpp +++ b/clang/test/OpenMP/declare_target_only_one_side_compilation.cpp @@ -67,9 +67,9 @@ // 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