[clang] [CIR][HIP] Add Stub body emission test coverage and Fix kernelHandle storage (PR #179823)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Feb 5 22:02:40 PST 2026
github-actions[bot] wrote:
<!--PREMERGE ADVISOR COMMENT: Linux-->
# :penguin: Linux x64 Test Results
* 87023 tests passed
* 758 tests skipped
* 2 tests failed
## Failed Tests
(click on a test name to see its output)
### Clang
<details>
<summary>Clang.CIR/CodeGenCUDA/kernel-call.cu</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 4
/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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -target-sdk-version=9.2 -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -x cuda -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir
# 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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -target-sdk-version=9.2 -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -x cuda -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir
# note: command had no output on stdout or stderr
# RUN: at line 6
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=CUDA-NEW
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=CUDA-NEW
# note: command had no output on stdout or stderr
# RUN: at line 8
/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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -fhip-new-launch-api -x hip -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir
# 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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -fhip-new-launch-api -x hip -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir
# note: command had no output on stdout or stderr
# RUN: at line 10
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=HIP-NEW
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=HIP-NEW
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu:54:13: error: HIP-NEW: expected string not found in input
# | // HIP-NEW: cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !void
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:1:1: note: scanning from here
# | !rec_hipStream = !cir.record<struct "hipStream" incomplete>
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:14:2: note: possible intended match here
# | cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !cir.func<(!s32i, !cir.float)> {alignment = 8 : i64} loc(#loc9)
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu:58:13: error: HIP-NEW: expected string not found in input
# | // HIP-NEW: cir.get_global @_Z6kernelif : !cir.ptr<!void>
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:36:170: note: scanning from here
# | %14 = cir.call @__hipPopCallConfiguration(%10, %11, %12, %13) : (!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_hipStream>>) -> !s32i loc(#loc9)
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:37:8: note: possible intended match here
# | %15 = cir.get_global @_Z6kernelif : !cir.ptr<!cir.func<(!s32i, !cir.float)>> loc(#loc9)
# | ^
# |
# | Input file: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | 1: !rec_hipStream = !cir.record<struct "hipStream" incomplete>
# | check:54'0 X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
# | 2: !s32i = !cir.int<s, 32>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~
# | 3: !u32i = !cir.int<u, 32>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~
# | 4: !u64i = !cir.int<u, 64>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~
# | 5: !void = !cir.void
# | check:54'0 ~~~~~~~~~~~~~~~~~~
# | 6: #loc3 = loc("/home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu":60:24)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | 9: #loc6 = loc("/home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu":60:37)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 10: !rec_dim3 = !cir.record<struct "dim3" {!u32i, !u32i, !u32i}>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 11: #loc10 = loc(fused[#loc3, #loc4])
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 12: #loc11 = loc(fused[#loc5, #loc6])
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 13: module @"/home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu" attributes {cir.lang = #cir.lang<cxx>, cir.module_asm = [], cir.triple = "x86_64-unknown-linux-gnu", dlti.dl_spec = #dlti.dl_spec<!llvm.ptr<270> = dense<32> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr = dense<64> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.mangling_mode" = "e", "dlti.legal_int_widths" = array<i32: 8, 16, 32, 64>, "dlti.stack_alignment" = 128 : i64>} {
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 14: cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !cir.func<(!s32i, !cir.float)> {alignment = 8 : i64} loc(#loc9)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:54'1 ? possible intended match
# | 15: cir.func private dso_local @__hipPopCallConfiguration(!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_hipStream>>) -> !s32i loc(#loc)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 16: cir.func private dso_local @hipLaunchKernel(!cir.ptr<!void>, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>>, !u64i, !cir.ptr<!rec_hipStream>) -> !u32i loc(#loc)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 17: cir.func no_inline dso_local @_Z21__device_stub__kernelif(%arg0: !s32i loc(fused[#loc3, #loc4]), %arg1: !cir.float loc(fused[#loc5, #loc6])) {
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 18: %0 = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] {alignment = 4 : i64} loc(#loc10)
# | 19: %1 = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["y", init] {alignment = 4 : i64} loc(#loc11)
# | .
# | .
# | .
# | 31: cir.store %9, %8 : !cir.ptr<!void>, !cir.ptr<!cir.ptr<!void>> loc(#loc9)
# | 32: %10 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["grid_dim"] {alignment = 8 : i64} loc(#loc9)
# | 33: %11 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["block_dim"] {alignment = 8 : i64} loc(#loc9)
# | 34: %12 = cir.alloca !u64i, !cir.ptr<!u64i>, ["shared_mem"] {alignment = 8 : i64} loc(#loc9)
# | 35: %13 = cir.alloca !cir.ptr<!rec_hipStream>, !cir.ptr<!cir.ptr<!rec_hipStream>>, ["stream"] {alignment = 8 : i64} loc(#loc9)
# | 36: %14 = cir.call @__hipPopCallConfiguration(%10, %11, %12, %13) : (!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_hipStream>>) -> !s32i loc(#loc9)
# | check:58'0 X~~~~~~~~~~~ error: no match found
# | 37: %15 = cir.get_global @_Z6kernelif : !cir.ptr<!cir.func<(!s32i, !cir.float)>> loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:58'1 ? possible intended match
# | 38: %16 = cir.cast bitcast %15 : !cir.ptr<!cir.func<(!s32i, !cir.float)>> -> !cir.ptr<!void> loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 39: %17 = cir.load %12 : !cir.ptr<!u64i>, !u64i loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 40: %18 = cir.load %13 : !cir.ptr<!cir.ptr<!rec_hipStream>>, !cir.ptr<!rec_hipStream> loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 41: %19 = cir.load align(8) %10 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc8)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 42: %20 = cir.load align(8) %11 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc8)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1
--
```
</details>
<details>
<summary>Clang.CIR/CodeGenCUDA/kernel-call.cu</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 4
/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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -target-sdk-version=9.2 -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -x cuda -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir
# 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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -target-sdk-version=9.2 -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -x cuda -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir
# note: command had no output on stdout or stderr
# RUN: at line 6
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=CUDA-NEW
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=CUDA-NEW
# note: command had no output on stdout or stderr
# RUN: at line 8
/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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -fhip-new-launch-api -x hip -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir
# 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/23/include -nostdsysteminc -triple x86_64-unknown-linux-gnu -fhip-new-launch-api -x hip -emit-cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu -o /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir
# note: command had no output on stdout or stderr
# RUN: at line 10
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=HIP-NEW
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck --input-file=/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu --check-prefix=HIP-NEW
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu:54:13: error: HIP-NEW: expected string not found in input
# | // HIP-NEW: cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !void
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:1:1: note: scanning from here
# | !rec_hipStream = !cir.record<struct "hipStream" incomplete>
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:14:2: note: possible intended match here
# | cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !cir.func<(!s32i, !cir.float)> {alignment = 8 : i64} loc(#loc9)
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu:58:13: error: HIP-NEW: expected string not found in input
# | // HIP-NEW: cir.get_global @_Z6kernelif : !cir.ptr<!void>
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:36:170: note: scanning from here
# | %14 = cir.call @__hipPopCallConfiguration(%10, %11, %12, %13) : (!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_hipStream>>) -> !s32i loc(#loc9)
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir:37:8: note: possible intended match here
# | %15 = cir.get_global @_Z6kernelif : !cir.ptr<!cir.func<(!s32i, !cir.float)>> loc(#loc9)
# | ^
# |
# | Input file: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/clang/test/CIR/CodeGenCUDA/Output/kernel-call.cu.tmp.hip.cir
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | 1: !rec_hipStream = !cir.record<struct "hipStream" incomplete>
# | check:54'0 X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
# | 2: !s32i = !cir.int<s, 32>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~
# | 3: !u32i = !cir.int<u, 32>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~
# | 4: !u64i = !cir.int<u, 64>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~
# | 5: !void = !cir.void
# | check:54'0 ~~~~~~~~~~~~~~~~~~
# | 6: #loc3 = loc("/home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu":60:24)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | 9: #loc6 = loc("/home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu":60:37)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 10: !rec_dim3 = !cir.record<struct "dim3" {!u32i, !u32i, !u32i}>
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 11: #loc10 = loc(fused[#loc3, #loc4])
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 12: #loc11 = loc(fused[#loc5, #loc6])
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 13: module @"/home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/CIR/CodeGenCUDA/kernel-call.cu" attributes {cir.lang = #cir.lang<cxx>, cir.module_asm = [], cir.triple = "x86_64-unknown-linux-gnu", dlti.dl_spec = #dlti.dl_spec<!llvm.ptr<270> = dense<32> : vector<4xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, i64 = dense<64> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, f80 = dense<128> : vector<2xi64>, !llvm.ptr = dense<64> : vector<4xi64>, i1 = dense<8> : vector<2xi64>, i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, f128 = dense<128> : vector<2xi64>, "dlti.endianness" = "little", "dlti.mangling_mode" = "e", "dlti.legal_int_widths" = array<i32: 8, 16, 32, 64>, "dlti.stack_alignment" = 128 : i64>} {
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 14: cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !cir.func<(!s32i, !cir.float)> {alignment = 8 : i64} loc(#loc9)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:54'1 ? possible intended match
# | 15: cir.func private dso_local @__hipPopCallConfiguration(!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_hipStream>>) -> !s32i loc(#loc)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 16: cir.func private dso_local @hipLaunchKernel(!cir.ptr<!void>, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>>, !u64i, !cir.ptr<!rec_hipStream>) -> !u32i loc(#loc)
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 17: cir.func no_inline dso_local @_Z21__device_stub__kernelif(%arg0: !s32i loc(fused[#loc3, #loc4]), %arg1: !cir.float loc(fused[#loc5, #loc6])) {
# | check:54'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 18: %0 = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] {alignment = 4 : i64} loc(#loc10)
# | 19: %1 = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["y", init] {alignment = 4 : i64} loc(#loc11)
# | .
# | .
# | .
# | 31: cir.store %9, %8 : !cir.ptr<!void>, !cir.ptr<!cir.ptr<!void>> loc(#loc9)
# | 32: %10 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["grid_dim"] {alignment = 8 : i64} loc(#loc9)
# | 33: %11 = cir.alloca !rec_dim3, !cir.ptr<!rec_dim3>, ["block_dim"] {alignment = 8 : i64} loc(#loc9)
# | 34: %12 = cir.alloca !u64i, !cir.ptr<!u64i>, ["shared_mem"] {alignment = 8 : i64} loc(#loc9)
# | 35: %13 = cir.alloca !cir.ptr<!rec_hipStream>, !cir.ptr<!cir.ptr<!rec_hipStream>>, ["stream"] {alignment = 8 : i64} loc(#loc9)
# | 36: %14 = cir.call @__hipPopCallConfiguration(%10, %11, %12, %13) : (!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_hipStream>>) -> !s32i loc(#loc9)
# | check:58'0 X~~~~~~~~~~~ error: no match found
# | 37: %15 = cir.get_global @_Z6kernelif : !cir.ptr<!cir.func<(!s32i, !cir.float)>> loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:58'1 ? possible intended match
# | 38: %16 = cir.cast bitcast %15 : !cir.ptr<!cir.func<(!s32i, !cir.float)>> -> !cir.ptr<!void> loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 39: %17 = cir.load %12 : !cir.ptr<!u64i>, !u64i loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 40: %18 = cir.load %13 : !cir.ptr<!cir.ptr<!rec_hipStream>>, !cir.ptr<!rec_hipStream> loc(#loc9)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 41: %19 = cir.load align(8) %10 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc8)
# | check:58'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 42: %20 = cir.load align(8) %11 : !cir.ptr<!rec_dim3>, !rec_dim3 loc(#loc8)
# | check:58'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/179823
More information about the cfe-commits
mailing list