ssquare08 added inline comments.
================
Comment at: clang/lib/CodeGen/TargetInfo.cpp:9439
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
+ // Make sure any variable with OpenMP declare target is visible to runtime
+ // except for those with hidden visibility
----------------
jhuber6 wrote:
> jhuber6 wrote:
> > ssquare08 wrote:
> > > jhuber6 wrote:
> > > > Just spitballing, is it possible to do this when we make the global
> > > > instead?
> > > This is something I was wondering as well. In
> > > CodeGenModule::GetOrCreateLLVMGlobal, when it creates a new global
> > > variable, it always uses the llvm::GlobalValue::ExternalLinkage. Seems
> > > like this changes somewhere later to internal for static globals. Do you
> > > know where that would be?
> > I'm not exactly sure, I remember deleting some code in D117806 that did
> > something like that, albeit incorrectly. But I'm not sure if you'd have the
> > necessary information to check whether or not there are updates attached to
> > it. We don't want to externalize things if we don't need to, otherwise we'd
> > get a lot of our device runtime variables with external visibility that now
> > can't be optimized out.
> Were you able to find a place for this when we generate the variable? You
> should be able to do something similar to the patch above if it's a declare
> target static to force it to have external visibility, but as mentioned
> before I would prefer we only do this if necessary which might take some
> extra analysis.
If you are asking about the GV, it is created in
'CodeGenModule::GetOrCreateLLVMGlobal' with external linkage always.
```
auto *GV = new llvm::GlobalVariable(
getModule(), Ty, false, llvm::GlobalValue::ExternalLinkage, nullptr,
MangledName, nullptr, llvm::GlobalVariable::NotThreadLocal,
getContext().getTargetAddressSpace(DAddrSpace));
```
The linkage, however, changes in 'CodeGenModule::EmitGlobalVarDefinition' based
on the information VarDecl
```
llvm::GlobalValue::LinkageTypes Linkage =
getLLVMLinkageVarDefinition(D, GV->isConstant());
```
Maybe you are suggesting changing the linkage information in 'VarDecl' itself?
================
Comment at: clang/test/OpenMP/declare_target_visibility_codegen.cpp:11-12
// HOST: @z = global i32 0
-// HOST-NOT: @.omp_offloading.entry.c
-// HOST-NOT: @.omp_offloading.entry.x
+// HOST: @.omp_offloading.entry.c__static__{{[0-9a-z]+_[0-9a-z]+}}
+// HOST: @.omp_offloading.entry.x__static__{{[0-9a-z]+_[0-9a-z]+}}
// HOST-NOT: @.omp_offloading.entry.y
----------------
jhuber6 wrote:
> ssquare08 wrote:
> > ssquare08 wrote:
> > > jhuber6 wrote:
> > > > If there are no updates between the host and device we can keep these
> > > > static without emitting an offloading entry.
> > > That 's a good point. I'll fix that.
> > I thought about this more and I think the behavior for these declare
> > target static globals should be the same as the other declare target.
> > Checking for update is not enough because users could also map these
> > variables. For update, it could be mapped with a pointer or the users could
> > pass address of these variables to an external function. Please let me know
> > what you think of these cases below:
> > ```
> > #pragma omp declare target
> > static int x[10];
> > #pragma omp end declare target
> >
> > //case 1
> > #pragma omp target update to(x)
> >
> > //case 2
> > int* y = &x[2];
> > #pragma omp target update to(y[0])
> >
> > //case 3
> > #pragma omp target map(always to:x)
> > {
> > x[0]= 111;
> > }
> >
> > //case 4
> > #pragma omp target
> > {
> > foo(&x[3]);
> > }
> > ```
> We should still be able to do this if there are either no updates at all in
> the module, or if the declare type is `nohost`. Doing anything more
> complicated would require some optimizations between the host and device we
> can't do yet. I'm making this point because making these statics external is
> a performance regression so we should only do it when needed. To that end we
> may even want a flag that entirely disables this feature.
I'll add a check to see if there are any updates in the module.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D129694/new/
https://reviews.llvm.org/D129694
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
