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

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

rampitec added a comment.

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!



================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUAliasAnalysis.cpp:114
+      case CallingConv::AMDGPU_KERNEL:
+        // In the kernel function, kernel arguments won't alias to (local)
+        // variables in shared or private address space.
----------------
hliao wrote:
> rampitec wrote:
> > hliao wrote:
> > > rampitec wrote:
> > > > What is we pass a pointer to LDS allocated on the host side?
> > > No LDS pointer is passed through constant variables in user code. So far, LDS is allocated and assigned during codegen. Late, as LDS could be assigned per kernel and from 0, we could choose clone sub-function (called by more than one kernel functions) so that assignment could be done within linker, or add GOT like relocation if sub-function is not cloned. The late needs to prepare GOT entries on the host side but that load from GOT could be easily distinguished in the backend as it needs to use PseuodSource in MMO.
> > But that's not constant, it is just Argument?
> In HIP, LDS pointer cannot be passed as an argument to a kernel.
In OpenCL it is legal.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D89525