[PATCH] D94814: [HIP] Support `__managed__` attribute

[Yaxun Liu via Phabricator via llvm-commits]

yaxunl marked 9 inline comments as done.
yaxunl added inline comments.


================
Comment at: clang/lib/Sema/SemaDeclAttr.cpp:7728
+  case ParsedAttr::AT_HIPManaged:
+    handleSimpleAttributeWithExclusions<HIPManagedAttr, CUDAGlobalAttr>(S, D,
+                                                                        AL);
----------------
aaron.ballman wrote:
> aaron.ballman wrote:
> > tra wrote:
> > > The code changes in the patch appear to treat `__managed__` variable as a superset of a `__device__` var. If that's indeed the case, adding an implicit `__device__` attribute here would help to simplify the code. This way the existing code can handle generic  `__device__` var functionality without additional changes, and would use `__managed__` checks for the cases specific for managed vars only.
> > > 
> > I think you're missing changes to the CUDA global attr to check for mutual exclusions with `__managed__` as well. Also, I think this won't do the right thing for redeclarations, like:
> > ```
> > __device__ extern int i;
> > __managed__ extern int i;
> > ```
> > This way the existing code can handle generic __device__ var functionality without additional changes, and would use __managed__ checks for the cases specific for managed vars only.
> 
> Another alternative to consider is to not create a new semantic attribute named `HIPManagedAttr` but to instead add a new spelling to `CUDADevice` in `Attr.td`, giving the class an `Accessor` to distinguish which spelling the user wrote in code, and use that accessor for the specific cases for managed vars.
I will add an implicit device attr for managed attr. I think this will have a more readable AST.


================
Comment at: clang/lib/Sema/SemaDeclAttr.cpp:7728
+  case ParsedAttr::AT_HIPManaged:
+    handleSimpleAttributeWithExclusions<HIPManagedAttr, CUDAGlobalAttr>(S, D,
+                                                                        AL);
----------------
yaxunl wrote:
> aaron.ballman wrote:
> > aaron.ballman wrote:
> > > tra wrote:
> > > > The code changes in the patch appear to treat `__managed__` variable as a superset of a `__device__` var. If that's indeed the case, adding an implicit `__device__` attribute here would help to simplify the code. This way the existing code can handle generic  `__device__` var functionality without additional changes, and would use `__managed__` checks for the cases specific for managed vars only.
> > > > 
> > > I think you're missing changes to the CUDA global attr to check for mutual exclusions with `__managed__` as well. Also, I think this won't do the right thing for redeclarations, like:
> > > ```
> > > __device__ extern int i;
> > > __managed__ extern int i;
> > > ```
> > > This way the existing code can handle generic __device__ var functionality without additional changes, and would use __managed__ checks for the cases specific for managed vars only.
> > 
> > Another alternative to consider is to not create a new semantic attribute named `HIPManagedAttr` but to instead add a new spelling to `CUDADevice` in `Attr.td`, giving the class an `Accessor` to distinguish which spelling the user wrote in code, and use that accessor for the specific cases for managed vars.
> I will add an implicit device attr for managed attr. I think this will have a more readable AST.
Actually `__managed__` does not need exclude with `__global__` since `__global__` cannot be used with variable. However, `__managed__` need to exclude with `__constant__` and `__shared__`. I will add handling for that.

`__managed__` should behave the same as `__device__` regarding compatibility with `extern`. It is allowed to be with `extern`. Currently for `extern __device__` var, on device side, it is treated as a normal device variable with definition; on host side, the shadow variable is emitted with internal linkage. This is fine for the default -fno-gpu-rdc mode. However, this does not work for -fgpu-rdc mode. For -fgpu-rdc mode, both device variable and shadow variable should be declarations. Since this issue is orthogonal to `__managed__`, I will create a separate patch to fix it.


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

https://reviews.llvm.org/D94814