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

[Michael Liao via Phabricator via llvm-commits]

hliao added a comment.

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

> In D89525#2336002 <https://reviews.llvm.org/D89525#2336002>, @yaxunl wrote:
>
>> 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.
>
> But not in OpenCL.

For OpenCL, since it won't allow generic pointer as arguments, there never be a case a generic pointer from OpenCL kernel function.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D89525