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