[PATCH] D129694: [OPENMP] Make declare target static global externally visible

Sunil Shrestha via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Tue Jul 26 18:21:35 PDT 2022


ssquare08 added inline comments.


================
Comment at: clang/test/OpenMP/declare_target_only_one_side_compilation.cpp:61
 // DEVICE-DAG: @G1 = {{.*}}global i32 0, align 4
-// DEVICE-DAG: @_ZL2G2 = internal {{.*}}global i32 0, align 4
+// DEVICE-DAG: @_ZL2G2 = {{.*}}global i32 0, align 4
 // DEVICE-DAG: @G3 = {{.*}}global i32 0, align 4
----------------
I wasn't expecting this to change. For some reason G2 gets the `OMPDeclareTargetDeclAttr::DT_Any` attribute instead of `OMPDeclareTargetDeclAttr::DT_NoHost` and because of that the visibility changes.  @jdoerfert, is `OMPDeclareTargetDeclAttr::DT_Any` attribute expected here? 


================
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:
> 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]);
}
```


================
Comment at: clang/test/OpenMP/target_update_messages.cpp:17
 
-static int y;
-#pragma omp declare target(y)
-
-void yyy() {
-#pragma omp target update to(y) // expected-error {{the host cannot update a declare target variable that is not externally visible.}}
-}
-
 int __attribute__((visibility("hidden"))) z;
 #pragma omp declare target(z)
----------------
jhuber6 wrote:
> jdoerfert wrote:
> > There is no test to show you can actually write the update now, is there?
> We should probably take the deleted code above and put it in an OpenMP runtime test to make sure it actually works now.
I have now added a test as suggested.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D129694/new/

https://reviews.llvm.org/D129694



More information about the cfe-commits mailing list