[llvm] [InferAddressSpaces] Support address space inference from load values (PR #171019)

via llvm-commits llvm-commits at lists.llvm.org
Sun Dec 7 00:47:48 PST 2025


github-actions[bot] wrote:

<!--PREMERGE ADVISOR COMMENT: Linux-->
# :penguin: Linux x64 Test Results

* 166850 tests passed
* 2913 tests skipped
* 2 tests failed

## Failed Tests
(click on a test name to see its output)

### Clang
<details>
<summary>Clang.CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu</summary>

```
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 5
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/clang -cc1 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/build/lib/clang/22/include -nostdsysteminc -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu -o - | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --check-prefixes=CHECK /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/clang -cc1 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/build/lib/clang/22/include -nostdsysteminc -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu -o -
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --check-prefixes=CHECK /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
# note: command had no output on stdout or stderr
# RUN: at line 6
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/clang -cc1 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/build/lib/clang/22/include -nostdsysteminc -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu -o - | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --check-prefixes=CHECK-SPIRV /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/clang -cc1 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/build/lib/clang/22/include -nostdsysteminc -triple spirv64-amd-amdhsa -fcuda-is-device -emit-llvm -x hip /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu -o -
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --check-prefixes=CHECK-SPIRV /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
# note: command had no output on stdout or stderr
# RUN: at line 7
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/clang -cc1 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/build/lib/clang/22/include -nostdsysteminc -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu -disable-O0-optnone -o - | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -S -O2 | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu --check-prefixes=OPT
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/clang -cc1 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/build/lib/clang/22/include -nostdsysteminc -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu -disable-O0-optnone -o -
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -S -O2
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu --check-prefixes=OPT
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu:438:14: error: OPT-NEXT: expected string not found in input
# | // OPT-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
# |              ^
# | <stdin>:67:52: note: scanning from here
# |  %0 = load ptr, ptr addrspace(1) %s.coerce, align 8
# |                                                    ^
# | <stdin>:67:52: note: with "TMP0" equal to "%0"
# |  %0 = load ptr, ptr addrspace(1) %s.coerce, align 8
# |                                                    ^
# | <stdin>:74:2: note: possible intended match here
# |  %4 = load float, ptr %3, align 4
# |  ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |             .
# |             .
# |             .
# |            62: } 
# |            63:  
# |            64: ; Function Attrs: mustprogress nofree noinline norecurse nosync nounwind willreturn memory(readwrite, inaccessiblemem: none, target_mem0: none, target_mem1: none) 
# |            65: define dso_local amdgpu_kernel void @_Z7kernel5P1S(ptr addrspace(1) noundef readonly captures(none) %s.coerce) local_unnamed_addr #2 { 
# |            66: entry: 
# |            67:  %0 = load ptr, ptr addrspace(1) %s.coerce, align 8 
# | next:438'0                                                        X error: no match found
# | next:438'1                                                          with "TMP0" equal to "%0"
# |            68:  %1 = addrspacecast ptr %0 to ptr addrspace(1) 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            69:  %2 = load i32, ptr addrspace(1) %1, align 4 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            70:  %inc = add nsw i32 %2, 1 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            71:  store i32 %inc, ptr addrspace(1) %1, align 4 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            72:  %y = getelementptr inbounds nuw i8, ptr addrspace(1) %s.coerce, i64 8 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            73:  %3 = load ptr, ptr addrspace(1) %y, align 8 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            74:  %4 = load float, ptr %3, align 4 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | next:438'2      ?                                 possible intended match
# |            75:  %add = fadd contract float %4, 1.000000e+00 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            76:  store float %add, ptr %3, align 4 
# | next:438'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            77:  ret void 
# | next:438'0     ~~~~~~~~~~
# |            78: } 
# | next:438'0     ~~
# |            79:  
# | next:438'0     ~
# |             .
# |             .
# |             .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

```
</details>

### LLVM
<details>
<summary>LLVM.CodeGen/NVPTX/lower-byval-args.ll</summary>

```
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# note: command had no output on stdout or stderr
# RUN: at line 3
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# note: command had no output on stdout or stderr
# RUN: at line 4
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -passes=nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -passes=nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_60
# note: command had no output on stdout or stderr
# RUN: at line 5
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-lower-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-lower-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,LOWER-ARGS,SM_70
# note: command had no output on stdout or stderr
# RUN: at line 6
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-copy-byval-args -S | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,COPY
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/opt -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-copy-byval-args -S
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=COMMON,COPY
# note: command had no output on stdout or stderr
# RUN: at line 7
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll -mcpu=sm_60 -mattr=ptx77 | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=PTX,PTX_60
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mcpu=sm_60 -mattr=ptx77
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll --check-prefixes=PTX,PTX_60
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll:265:13: error: PTX-NEXT: expected string not found in input
# | ; PTX-NEXT: cvta.local.u64 %SP, %SPL;
# |             ^
# | <stdin>:153:31: note: scanning from here
# |  mov.b64 %SPL, __local_depot5;
# |                               ^
# | <stdin>:162:2: note: possible intended match here
# |  cvta.local.u64 %rd5, %rd4;
# |  ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |             .
# |             .
# |             .
# |           148:  .reg .b64 %SPL; 
# |           149:  .reg .b32 %r<3>; 
# |           150:  .reg .b64 %rd<6>; 
# |           151:  
# |           152: // %bb.0: // %entry 
# |           153:  mov.b64 %SPL, __local_depot5; 
# | next:265'0                                   X error: no match found
# |           154:  ld.param.b64 %rd1, [escape_ptr_gep_store_param_0]; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           155:  cvta.to.global.u64 %rd2, %rd1; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           156:  add.u64 %rd3, %SPL, 0; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~
# |           157:  ld.param.b32 %r1, [escape_ptr_gep_store_param_1+4]; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           158:  st.local.b32 [%rd3+4], %r1; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           159:  ld.param.b32 %r2, [escape_ptr_gep_store_param_1]; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           160:  st.local.b32 [%rd3], %r2; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           161:  add.s64 %rd4, %rd3, 4; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~
# |           162:  cvta.local.u64 %rd5, %rd4; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | next:265'1      ?                           possible intended match
# |           163:  st.global.b64 [%rd2], %rd5; 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           164:  ret; 
# | next:265'0     ~~~~~~
# |           165:  // -- End function 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~
# |           166: } 
# | next:265'0     ~~
# |           167:  // .globl escape_ptrtoint // -- Begin function escape_ptrtoint 
# | next:265'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |             .
# |             .
# |             .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

```
</details>

If these failures are unrelated to your changes (for example tests are broken or flaky at HEAD), please open an issue at https://github.com/llvm/llvm-project/issues and add the `infrastructure` label.

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


More information about the llvm-commits mailing list