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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
