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

[Michael Liao via Phabricator via llvm-commits]

hliao added a comment.

In D89525#2335949 <https://reviews.llvm.org/D89525#2335949>, @yaxunl wrote:

> In D89525#2333924 <https://reviews.llvm.org/D89525#2333924>, @hliao wrote:
>
>> @yaxunl could you double-check that OpenCL also follows that rule.
>> @nhaehnle could you check whether that potentially breaks graphics.
>
> 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

I don't think that should be valid regard to the lifetime definition of shared variables, which have a maximal lifetime of the kernel execution. Instead, global CONSTANT and GLOBAL variables have longer lifetime than kernel execution. It's not valid to assign a pointer of LOCAL variables to a CONSTANT or GLOBAL in the static initializer (when no kernel execution is started). Within kernel/function body, CONSTANT cannot be assigned anymore. That's not a valid case.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D89525