[OPENMP]Fix PR41826: symbols visibility in device code.

[ABataev]

Summary

Currently, we ignore all locality attributes/info when building for
the device and thus all symblos are externally visible and can be
preemted at the runtime. It may lead to incorrect results. We need to
follow the same logic, compiler uses for static/pie builds. But in some
cases changing of dso locality may lead to problems with codegen, so
instead mark external symbols as hidden instead in the device code.

[jdoerfert]

LGTM.

[JonChesterfield]

@ABataev @jhuber6 and I would like to change this to annotate variables with whatever visibility the user asked for.

Can you recall if the concern about premption here was for the deviceRTL, in which case it can be comprehensively dealt with by internalizing the deviceRTL symbols, or for user defined things in which case we want whatever the linker flag -Bsymbolic maps onto as an IR construct. (<- said construct is dso_local)

[ABataev]

In D70549#3256063, @JonChesterfield wrote:

@ABataev @jhuber6 and I would like to change this to annotate variables with whatever visibility the user asked for.

Can you recall if the concern about premption here was for the deviceRTL, in which case it can be comprehensively dealt with by internalizing the deviceRTL symbols, or for user defined things in which case we want whatever the linker flag -Bsymbolic maps onto as an IR construct.

Hi, check the original PR, it has all required details. Maybe add a check that if the device is same as the host, then add hidden visibility (to avoid problems with the standard libs etc.), otherwise keep it as is?

clang/lib/AST/Decl.cpp

// Everything not covered here has no linkage.
     return LinkageInfo::none();
   }
 
   // If we ended up with non-externally-visible linkage, visibility should
   // always be default.
   if (!isExternallyVisible(LV.getLinkage()))
     return LinkageInfo(LV.getLinkage(), DefaultVisibility, false);
 
+  // Mark the symbols as hidden when compiling for the device.
+  if (Context.getLangOpts().OpenMP && Context.getLangOpts().OpenMPIsDevice)
+    LV.mergeVisibility(HiddenVisibility, /*newExplicit=*/false);
+
   return LV;
 }
 
 LinkageInfo
 LinkageComputer::getLVForClassMember(const NamedDecl *D,
                                      LVComputationKind computation,
                                      bool IgnoreVarTypeLinkage) {
   // Only certain class members have linkage.  Note that fields don't

clang/lib/CodeGen/CodeGenModule.cpp

static bool shouldAssumeDSOLocal(const CodeGenModule &CGM,
   // Only handle COFF and ELF for now.
   if (!TT.isOSBinFormatELF())
     return false;
 
   // 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 && !LOpts.OpenMPIsDevice)
+  if (RM != llvm::Reloc::Static && !LOpts.PIE)
     return false;
 
   // A definition cannot be preempted from an executable.
   if (!GV->isDeclarationForLinker())
     return true;
 
   // Most PIC code sequences that assume that a symbol is local cannot produce a
   // 0 if it turns out the symbol is undefined. While this is ABI and relocation

clang/test/OpenMP/declare_target_codegen.cpp

 
 // SIMD-ONLY-NOT: {{__kmpc|__tgt}}
 // KMPC-ONLY-NOT: __tgt
 
 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
 // CHECK-DAG: Bake
 // CHECK-NOT: @{{hhh|ggg|fff|eee}} =
 // CHECK-DAG: @aaa = external global i32,
-// CHECK-DAG: @bbb ={{ dso_local | }}global i32 0,
+// CHECK-DAG: @bbb ={{ hidden | }}global i32 0,
 // CHECK-DAG: weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @bbb to i8*),
 // CHECK-DAG: @ccc = external global i32,
-// CHECK-DAG: @ddd ={{ dso_local | }}global i32 0,
+// CHECK-DAG: @ddd ={{ hidden | }}global i32 0,
 // CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global i32* null
 // CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global i32* null
 // CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global i32* null
 // CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global i32* null
 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
-// CHECK-DAG: @b ={{ dso_local | }}global i32 15,
-// CHECK-DAG: @d ={{ dso_local | }}global i32 0,
+// CHECK-DAG: @b ={{ hidden | }}global i32 15,
+// CHECK-DAG: @d ={{ hidden | }}global i32 0,
 // CHECK-DAG: @c = external global i32,
-// CHECK-DAG: @globals ={{ dso_local | }}global %struct.S zeroinitializer,
+// CHECK-DAG: @globals ={{ hidden | }}global %struct.S zeroinitializer,
 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
-// CHECK-DAG: @out_decl_target ={{ dso_local | }}global i32 0,
+// CHECK-DAG: @out_decl_target ={{ hidden | }}global i32 0,
 // CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+84]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+85]]_ctor to i8*)],
 // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
 
 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}})
 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}})
 // CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+78]]_ctor()
 

clang/test/OpenMP/nvptx_allocate_codegen.cpp

 extern const omp_allocator_handle_t omp_const_mem_alloc;
 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
 extern const omp_allocator_handle_t omp_pteam_mem_alloc;
 extern const omp_allocator_handle_t omp_thread_mem_alloc;
 
 // CHECK-DAG: @{{.+}}St1{{.+}}b{{.+}} = external global i32,
-// CHECK-DAG: @a ={{ dso_local | }}global i32 0,
-// CHECK-DAG: @b ={{ dso_local | }}addrspace(4) global i32 0,
-// CHECK-DAG: @c ={{ dso_local | }}global i32 0,
-// CHECK-DAG: @d ={{ dso_local | }}global %struct.St1 zeroinitializer,
-// CHECK-DAG: @{{.+}}ns{{.+}}a{{.+}} ={{ dso_local | }}addrspace(3) global i32 0,
+// CHECK-DAG: @a ={{ hidden | }}global i32 0,
+// CHECK-DAG: @b ={{ hidden | }}addrspace(4) global i32 0,
+// CHECK-DAG: @c ={{ hidden | }}global i32 0,
+// CHECK-DAG: @d ={{ hidden | }}global %struct.St1 zeroinitializer,
+// CHECK-DAG: @{{.+}}ns{{.+}}a{{.+}} ={{ hidden | }}addrspace(3) global i32 0,
 // CHECK-DAG: @{{.+}}main{{.+}}a{{.*}} = internal global i32 0,
 // CHECK-DAG: @{{.+}}ST{{.+}}m{{.+}} = external global i32,
 // CHECK-DAG: @bar_c = internal global i32 0,
 // CHECK-DAG: @bar_b = internal addrspace(3) global double 0.000000e+00,
 struct St{
  int a;
 };
 

 
 // CHECK: define {{.*}}i32 @{{.+}}foo{{.+}}()
 // CHECK-NOT: alloca i32,
 
 extern template int ST<int>::m;
 
 void baz(float &);
 
-// CHECK: define{{ dso_local | }}void @{{.+}}bar{{.+}}()
+// CHECK: define{{ hidden | }}void @{{.+}}bar{{.+}}()
 void bar() {
   // CHECK: alloca float,
   float bar_a;
   // CHECK: alloca double,
   double bar_b;
   int bar_c;
 #pragma omp allocate(bar_c) allocator(omp_cgroup_mem_alloc)
   // CHECK: call void [[OUTLINED:@.+]](i32* %{{.+}}, i32* %{{.+}})

clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp

 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix SIMD-ONLY
 
 #ifndef HEADER
 #define HEADER
 
 // SIMD-ONLY-NOT: {{__kmpc|__tgt}}
 
 // DEVICE-DAG: [[C_ADDR:.+]] = internal global i32 0,
-// DEVICE-DAG: [[CD_ADDR:@.+]] ={{ dso_local | }}global %struct.S zeroinitializer,
+// DEVICE-DAG: [[CD_ADDR:@.+]] ={{ hidden | }}global %struct.S zeroinitializer,
 // HOST-DAG: @[[C_ADDR:.+]] = internal global i32 0,
-// HOST-DAG: @[[CD_ADDR:.+]] ={{ dso_local | }}global %struct.S zeroinitializer,
+// HOST-DAG: @[[CD_ADDR:.+]] ={{ hidden | }}global %struct.S zeroinitializer,
 
 #pragma omp declare target
 int foo() { return 0; }
 #pragma omp end declare target
 int bar() { return 0; }
 #pragma omp declare target (bar)
 int baz() { return 0; }
 
 #pragma omp declare target
 int doo() { return 0; }
 #pragma omp end declare target
 int car() { return 0; }
 #pragma omp declare target (bar)
 int caz() { return 0; }
 
-// DEVICE-DAG: define{{ dso_local | }}i32 [[FOO:@.*foo.*]]()
-// DEVICE-DAG: define{{ dso_local | }}i32 [[BAR:@.*bar.*]]()
-// DEVICE-DAG: define{{ dso_local | }}i32 [[BAZ:@.*baz.*]]()
-// DEVICE-DAG: define{{ dso_local | }}i32 [[DOO:@.*doo.*]]()
-// DEVICE-DAG: define{{ dso_local | }}i32 [[CAR:@.*car.*]]()
-// DEVICE-DAG: define{{ dso_local | }}i32 [[CAZ:@.*caz.*]]()
+// DEVICE-DAG: define{{ hidden | }}i32 [[FOO:@.*foo.*]]()
+// DEVICE-DAG: define{{ hidden | }}i32 [[BAR:@.*bar.*]]()
+// DEVICE-DAG: define{{ hidden | }}i32 [[BAZ:@.*baz.*]]()
+// DEVICE-DAG: define{{ hidden | }}i32 [[DOO:@.*doo.*]]()
+// DEVICE-DAG: define{{ hidden | }}i32 [[CAR:@.*car.*]]()
+// DEVICE-DAG: define{{ hidden | }}i32 [[CAZ:@.*caz.*]]()
 
 static int c = foo() + bar() + baz();
 #pragma omp declare target (c)
 // HOST-DAG: @[[C_CTOR:__omp_offloading__.+_c_l44_ctor]] = private constant i8 0
 // DEVICE-DAG: define internal void [[C_CTOR:@__omp_offloading__.+_c_l44_ctor]]()
 // DEVICE-DAG: call i32 [[FOO]]()
 // DEVICE-DAG: call i32 [[BAR]]()
 // DEVICE-DAG: call i32 [[BAZ]]()

clang/test/OpenMP/nvptx_target_codegen.cpp

#pragma omp parallel
   // CHECK: [[TERMINATE]]
   // CHECK: call void @__kmpc_kernel_deinit(
   // CHECK: call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
   // CHECK: br label {{%?}}[[EXIT]]
   //
   // CHECK: [[EXIT]]
   // CHECK: ret void
 
-  // CHECK: define{{ dso_local | }}i32 [[BAZ]](i32 [[F:%.*]], double* dereferenceable{{.*}})
+  // CHECK: define{{ hidden | }}i32 [[BAZ]](i32 [[F:%.*]], double* dereferenceable{{.*}})
   // CHECK: alloca i32,
   // CHECK: [[LOCAL_F_PTR:%.+]] = alloca i32,
   // CHECK: [[ZERO_ADDR:%.+]] = alloca i32,
   // CHECK: [[BND_ZERO_ADDR:%.+]] = alloca i32,
   // CHECK: store i32 0, i32* [[BND_ZERO_ADDR]]
   // CHECK: store i32 0, i32* [[ZERO_ADDR]]
   // CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* [[UNKNOWN]]
   // CHECK: [[PAR_LEVEL:%.+]] = call i16 @__kmpc_parallel_level(%struct.ident_t* [[UNKNOWN]], i32 [[GTID]])

clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp

struct T1 {
   char c;
   T1() : a(12), f(15) {}
   T1 &operator+(T1 &b) { f += b.a; return *this;}
 };
 
 #pragma omp declare target
 T a = T();
 T f = a;
-// CHECK: define{{ dso_local | }}void @{{.+}}foo{{.+}}([[T]]* byval([[T]]) align {{.+}})
+// CHECK: define{{ hidden | }}void @{{.+}}foo{{.+}}([[T]]* byval([[T]]) align {{.+}})
 void foo(T a = T()) {
   return;
 }
-// CHECK: define{{ dso_local | }}[6 x i64] @{{.+}}bar{{.+}}()
+// CHECK: define{{ hidden | }}[6 x i64] @{{.+}}bar{{.+}}()
 T bar() {
 // CHECK:      bitcast [[T]]* %{{.+}} to [6 x i64]*
 // CHECK-NEXT: load [6 x i64], [6 x i64]* %{{.+}},
 // CHECK-NEXT: ret [6 x i64]
   return T();
 }
-// CHECK: define{{ dso_local | }}void @{{.+}}baz{{.+}}()
+// CHECK: define{{ hidden | }}void @{{.+}}baz{{.+}}()
 void baz() {
 // CHECK:      call [6 x i64] @{{.+}}bar{{.+}}()
 // CHECK-NEXT: bitcast [[T]]* %{{.+}} to [6 x i64]*
 // CHECK-NEXT: store [6 x i64] %{{.+}}, [6 x i64]* %{{.+}},
   T t = bar();
 }
 T1 a1 = T1();
 T1 f1 = a1;
-// CHECK: define{{ dso_local | }}void @{{.+}}foo1{{.+}}([[T1]]* byval([[T1]]) align {{.+}})
+// CHECK: define{{ hidden | }}void @{{.+}}foo1{{.+}}([[T1]]* byval([[T1]]) align {{.+}})
 void foo1(T1 a = T1()) {
   return;
 }
-// CHECK: define{{ dso_local | }}[[T1]] @{{.+}}bar1{{.+}}()
+// CHECK: define{{ hidden | }}[[T1]] @{{.+}}bar1{{.+}}()
 T1 bar1() {
 // CHECK:      load [[T1]], [[T1]]*
 // CHECK-NEXT: ret [[T1]]
   return T1();
 }
-// CHECK: define{{ dso_local | }}void @{{.+}}baz1{{.+}}()
+// CHECK: define{{ hidden | }}void @{{.+}}baz1{{.+}}()
 void baz1() {
 // CHECK: call [[T1]] @{{.+}}bar1{{.+}}()
   T1 t = bar1();
 }
 #pragma omp end declare target
 
 BIGTYPE foo(BIGTYPE f) {
 #pragma omp target map(f)
   f = 1;
   return f;
 }
 
 // CHECK: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l75([[BIGTYPE:.+]]*
 // CHECK: store [[BIGTYPE]] {{0xL00000000000000003FFF000000000000|0xM3FF00000000000000000000000000000}}, [[BIGTYPE]]* %