[PATCH] D91928: [nvptx] Skip alloca for read-only byval arguments.

[Michael Liao via Phabricator via llvm-commits]

hliao added a comment.

In D91928#2410164 <https://reviews.llvm.org/D91928#2410164>, @jlebar wrote:

>> I don't believe there's any exception to prove deduction [of the readonly attribute] wrong.
>
> Understood.
>
>> The address space inference here only refers to the one in the backend directly after this argument lowering gpass.
>
> Also understood.
>
> This isn't speaking to my concern, though.
>
> Suppose we have
>
>   __global__ void foo(int x, const int* y, int* out, bool flag) {
>     int* ptr = flag ? &x : y;
>     *out = *ptr;
>   }
>
> In this case we can say with confidence that x is readonly.
>
> But address space inference cannot infer the address space of ptr (how could it?).  Therefore we will do a generic load, which is wrong.

I see your point. PTX doesn't state the generic addressing could be performed on that parameter space. But, that case could be excluded with the extra check on how that parameter space pointer is used. In case it's not used in PHI or SELECT and cannot ensure the result is also a pointer to the parameter space, we could skip alloca insertion.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D91928