[PATCH] D89525: [amdgpu] Enhance AMDGPU AA.

[Yaxun Liu via Phabricator via llvm-commits]

yaxunl added a comment.

In D89525#2335991 <https://reviews.llvm.org/D89525#2335991>, @rampitec wrote:

> In D89525#2335989 <https://reviews.llvm.org/D89525#2335989>, @yaxunl wrote:
>
>> In D89525#2335972 <https://reviews.llvm.org/D89525#2335972>, @rampitec wrote:
>>
>>> In D89525#2335949 <https://reviews.llvm.org/D89525#2335949>, @yaxunl wrote:
>>>
>>>> I think they are correct for OpenCL, since in OpenCL shared var can only be declared in kernel function or passed by kernel arg.
>>>>
>>>> However I am not sure whether a constant pointer can pointer to shared memory, i.e, whether the address of a shared variable is compile time constant, or whether the following is valid code:
>>>>
>>>>   __shared__ int a;
>>>>   
>>>>   __constant__ int *b = &a;
>>>>
>>>> Currently clang allows it but nvcc does not https://godbolt.org/z/9W8vee
>>>>
>>>> I tends to agree with nvcc's treatment since this allows more flexible way of implementing shared variable supports in backend. @tra for advice
>>>
>>> But you are not checking for a constant pointer here!
>>
>> In HIP `__constant__` is a variable attribute, not the address space of the pointee. `__constant__ int *` means a pointer itself in constant address space and pointing to generic/flat address space.
>
> Where do you check for this specifically in this block:
>
>   } else if (const Argument *Arg = dyn_cast<Argument>(ObjA)) {
>      const Function *F = Arg->getParent();
>      switch (F->getCallingConv()) {
>      case CallingConv::AMDGPU_KERNEL:
>        // In the kernel function, kernel arguments won't alias to (local)
>        // variables in shared or private address space.
>        return NoAlias;

I was talking about semantic check in language. Here is the IR. In IR a kernel arg can pointing to constant or global addr due to promotion. Originally all kernel arg of HIP points to generic addr space only.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D89525