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

[Michael Liao via Phabricator via llvm-commits]

hliao added a comment.

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

> In D89525#2336479 <https://reviews.llvm.org/D89525#2336479>, @hliao wrote:
>
>> 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 kernel function arguments, there never be such a case, a generic pointer argument.
>
> OK, that makes sense. Now we only need to make sure we never compile anything but HIP and OpenCL <= 2.0. For example that we will not support fortran maybe? And never extend any of the languages to allow different address space arguments.
>
> My point here is that it is language specific, but there is nothing language specific on the AA implementation. A possible solution is to define some attributes or metadata saying that a certain situation may not ever happen and then check it in AA. That property has to ve set by the language though as AA may not know the source semantics.

Unless we change how LDS is used in the future hardware, I don't see any change of the lifetime of LDS variables. Similar to PRIVATE variables, they are only valid once the kernel is executed. Any reference to them outside the kernel execution doesn't make any sense. All languages and models so far have the similar usage of LDS. Any reference to LDS variables outside kernel execution needs persistent LDS support. That should be a significant change that requires more component support, especially the hardware.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D89525