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

[Aaron Ballman via Phabricator via llvm-commits]

aaron.ballman added inline comments.


================
Comment at: clang/include/clang/Basic/AttrDocs.td:5426
+The ``__managed__`` attribute can be applied to a global variable declaration in HIP.
+A managed variable is emitted as a undefined global symbol in device binary and
+registered by ``__hipRegisterManagedVariable`` in init functions. HIP runtime allocates
----------------



================
Comment at: clang/include/clang/Basic/AttrDocs.td:5427-5429
+registered by ``__hipRegisterManagedVariable`` in init functions. HIP runtime allocates
+managed memory and use it to define the symbol when loading the device binary.
+A managed variable can be accessed in both device and host code.
----------------



================
Comment at: clang/lib/Sema/SemaDeclAttr.cpp:7728
+  case ParsedAttr::AT_HIPManaged:
+    handleSimpleAttributeWithExclusions<HIPManagedAttr, CUDAGlobalAttr>(S, D,
+                                                                        AL);
----------------
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;
```


================
Comment at: clang/lib/Sema/SemaDeclAttr.cpp:7728
+  case ParsedAttr::AT_HIPManaged:
+    handleSimpleAttributeWithExclusions<HIPManagedAttr, CUDAGlobalAttr>(S, D,
+                                                                        AL);
----------------
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.


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

https://reviews.llvm.org/D94814