[llvm-branch-commits] [llvm] [AMDGPU] Remove the pass `AMDGPUPromoteKernelArguments` (PR #137655)

[Shilei Tian via llvm-branch-commits]

================
@@ -11,11 +10,9 @@ define amdgpu_kernel void @ptr_nest_3(ptr addrspace(1) nocapture readonly %Arg)
 ; CHECK-NEXT:  entry:
 ; CHECK-NEXT:    [[I:%.*]] = tail call i32 @llvm.amdgcn.workitem.id.x()
 ; CHECK-NEXT:    [[P1:%.*]] = getelementptr inbounds ptr, ptr addrspace(1) [[ARG:%.*]], i32 [[I]]
-; CHECK-NEXT:    [[P2:%.*]] = load ptr, ptr addrspace(1) [[P1]], align 8, !amdgpu.noclobber [[META0:![0-9]+]]
-; CHECK-NEXT:    [[P2_GLOBAL:%.*]] = addrspacecast ptr [[P2]] to ptr addrspace(1)
-; CHECK-NEXT:    [[P3:%.*]] = load ptr, ptr addrspace(1) [[P2_GLOBAL]], align 8, !amdgpu.noclobber [[META0]]
-; CHECK-NEXT:    [[P3_GLOBAL:%.*]] = addrspacecast ptr [[P3]] to ptr addrspace(1)
-; CHECK-NEXT:    store float 0.000000e+00, ptr addrspace(1) [[P3_GLOBAL]], align 4
+; CHECK-NEXT:    [[P2:%.*]] = load ptr, ptr addrspace(1) [[P1]], align 8
+; CHECK-NEXT:    [[P3:%.*]] = load ptr, ptr [[P2]], align 8
----------------
shiltian wrote:

The parent PR in the stack (i.e. https://github.com/llvm/llvm-project/pull/137488) can cover the case of a direct use of pointer kernel argument, as shown in the `P2` check line, but it can't do a good job for indirection.

The case under discussion here is something like:

```
struct S {
  int a;
  void *b;
};
```

The direct load/store of `S *` would be casted to the correct AS1, which is fine. However, the load/store of `S::b` is not, as shown in the `P3` of the check line.

> A pointer passed from host cannot be anything but global and be valid. So, this is a surprising change.

That's why I said it is not 100% right but probably 100% practically meaningful.

For example, I can do something like:

```

struct S {
  int a;
  void *b;
};

__global__ void kernel1(S *p) {
  __shared__ ss[2];
  p->b = (void *)ss;
}

__global__ void kernel2(S *p) {
  int *p = (int *)p->b;
  *p = 1;
}

int foo() {
  S *p;
  hipMalloc(&p, ...);
  kernel1<<<...>>>(p);
  kernel2<<<...>>>(p);
}
```

Is this practically correct? No. Is this legal code? I think it is. It will just causes a runtime crash.

https://github.com/llvm/llvm-project/pull/137655