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

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

rampitec added a comment.

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;


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D89525