This revision was automatically updated to reflect the committed changes.
Closed by commit rG37561ba89b7d: -fno-semantic-interposition: Don't set
dso_local on GlobalVariable (authored by MaskRay).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D102583/new/
https://reviews.llvm.org/D102583
Files:
clang/lib/CodeGen/CodeGenModule.cpp
clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c
clang/test/CodeGen/attr-weakref2.c
clang/test/CodeGen/semantic-interposition.c
clang/test/CodeGenCUDA/device-stub.cu
clang/test/CodeGenCUDA/device-var-linkage.cu
clang/test/CodeGenCUDA/managed-var.cu
clang/test/CodeGenCUDA/static-device-var-rdc.cu
clang/test/OpenMP/declare_target_only_one_side_compilation.cpp
Index: clang/test/OpenMP/declare_target_only_one_side_compilation.cpp
===================================================================
--- clang/test/OpenMP/declare_target_only_one_side_compilation.cpp
+++ 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
Index: clang/test/CodeGenCUDA/static-device-var-rdc.cu
===================================================================
--- clang/test/CodeGenCUDA/static-device-var-rdc.cu
+++ 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;
Index: clang/test/CodeGenCUDA/managed-var.cu
===================================================================
--- clang/test/CodeGenCUDA/managed-var.cu
+++ 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;
Index: clang/test/CodeGenCUDA/device-var-linkage.cu
===================================================================
--- clang/test/CodeGenCUDA/device-var-linkage.cu
+++ 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;
Index: clang/test/CodeGenCUDA/device-stub.cu
===================================================================
--- clang/test/CodeGenCUDA/device-stub.cu
+++ 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;
Index: clang/test/CodeGen/semantic-interposition.c
===================================================================
--- clang/test/CodeGen/semantic-interposition.c
+++ 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()
Index: clang/test/CodeGen/attr-weakref2.c
===================================================================
--- clang/test/CodeGen/attr-weakref2.c
+++ 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) {
Index: clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c
===================================================================
--- clang/test/CodeGen/SystemZ/builtins-systemz-zvector-constrained.c
+++ 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
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ 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<llvm::Function>(GV) && GV->canBenefitFromLocalAlias()))
return false;
return !(CGM.getLangOpts().SemanticInterposition ||
CGM.getLangOpts().HalfNoSemanticInterposition);
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
