ssahasra wrote: > > > I'm a little confused why want to do adjustments like this in CodeGen; > > > shouldn't Sema know the address-space of variables? > > > > I think it needs to be tracked in sema because of `constexpr` (which can > > happen at compile time or runtime) and AST matchers, etc. > > To be clear, this patch just aims to adjust the emitted AS, not the > user-facing AS (which remains `__shared__` for the time being).
That claim is not entirely true. The AMDGPU target _specifies_ that the address space of the builtin type `__amdgpu_named_workgroup_barrier_t` is the new BARRIER address space (13). The declaration may look like it is being declared in the `__shared__` address space, which frankly, I find difficult to think through in itself: ``` __shared__ __amdgpu_named_workgroup_barrier_t bar; ``` > Does Sema need to be aware of the AS used by the IR at any point ? If so I'll > of course look into fixing it, I just didn't think it was the case when I > made this. I followed the breadcrumbs of how we do it for other similar cases > and built from there Sema does need to check the `address_space` attribute. So maybe it should be aware that this type has a specific address space which cannot be changed? I am not really familiar about the division of concerns to be used here, but from first principles, it feels like the address space of a barrier on AMDGPU is known as soon as its type is known, which means it should be known to Sema. https://github.com/llvm/llvm-project/pull/195612 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
