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