jhuber6 added a comment.
I still think we shouldn't bother making all the noise containing the original
name. Just mangle it and treat it like every other declare target variable
without introducing any extra complexity. These symbols never should've been
emitted in the first place so I'm not concerned if someone cracks open a binary
and sees some ugly names. CUDA and HIP just mangle the declaration directly as
far as I'm aware.
================
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:
> 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.
================
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
----------------
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.
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
