================ @@ -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 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
