[NVPTX] NVPTX byval lowering copies aggregate arguments to stack even for read-only field access

I am looking at NVPTX codegen for a case where a struct is passed by value:

```c
struct MyStruct { int *X; int *Y; int XX[20]; };
__device__ __noinline__ void call(int x, int y, int *q, MyStruct z) {
q[0] = z.XX[x] + z.X[x] + z.Y[y];
}

__global__ void kernel(int x, int y, int *q, MyStruct z) {
call(x, y, q, z);
}
```
In the NVPTX instruction selection lowering stage, during LowerCall, the byval argument is copied into stack memory before its fields are accessed.

From the SASS, it looks like the incoming object is first spilled into local stack space and then later reloaded for field access, rather than accessing the byval object directly from its incoming location.

In this case, the object is only read from, and its address is never taken. So I am wondering:

Is there already an optimization that should avoid materializing this byval argument on the stack?
Is the stack copy required by NVPTX lowering for correctness, or is it just a conservative implementation detail?
If it is not required for correctness, would avoiding this stack materialization be a reasonable backend improvement?
The relevant SASS appears to show a stack frame being used for the byval argument, with stores to local memory followed by loads from that same stack region before the final global stores.

If helpful, I can also provide:

the LLVM IR
the PTX output
the exact compiler flags / LLVM revision used
a small reproducer