[PATCH] D98193: [CUDA][HIP] Allow non-ODR use of host var in device
Yaxun Liu via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Tue Apr 6 12:04:26 PDT 2021
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
More information about the cfe-commits
mailing list