[PATCH] D98193: [CUDA][HIP] Allow non-ODR use of host var in device

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added inline comments.


================
Comment at: clang/lib/Sema/SemaExpr.cpp:17050
+    auto Target = SemaRef.IdentifyCUDATarget(FD);
+    if (Var && Var->isFileVarDecl() && !Var->hasAttr<CUDADeviceAttr>() &&
+        !Var->hasAttr<CUDAConstantAttr>() && !Var->hasAttr<CUDASharedAttr>() &&
----------------
rsmith wrote:
> I suspect you want `hasGlobalStorage` rather than `isFileVarDecl` here (that is, preserve the condition from the deleted code), in order to disallow use of host-side local static variables from device-side functions:
> 
> ```
> __host__ void f() {
>   static int n;
>   struct X {
>     __device__ void g() { ++n; }
>   };
>   // ...
> }
> ```
For function scope static variable, if the smallest enclosing function is device or device host function, the static variable without `__device__` or `__constant__` attribute is allowed and the variable is emitted at device side (https://godbolt.org/z/PY5d3WGas). In that case, a device function is allowed to access that static variable even if it does not have `__device__` or `__constant__` attribute.

I will make changes to handle the function scope static variable.


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

https://reviews.llvm.org/D98193