[clang] [NVPTX][clang] Ensure CLZ(0) is defined on NVPTX (PR #185630)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 10 05:32:37 PDT 2026
github-actions[bot] wrote:
<!--PREMERGE ADVISOR COMMENT: Linux-->
# :penguin: Linux x64 Test Results
* 86841 tests passed
* 1309 tests skipped
* 1 test failed
## Failed Tests
(click on a test name to see its output)
### Clang
<details>
<summary>Clang.Headers/gpuintrin.c</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 2
/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 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/Inputs/include -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/../../lib/Headers/ -triple amdgcn-amd-amdhsa -emit-llvm /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c -o - | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c --check-prefix=AMDGPU
# 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 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/Inputs/include -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/../../lib/Headers/ -triple amdgcn-amd-amdhsa -emit-llvm /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c -o -
# 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/Headers/gpuintrin.c --check-prefix=AMDGPU
# 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/23/include -nostdsysteminc -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/Inputs/include -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/../../lib/Headers/ -triple nvptx64-nvidia-cuda -target-feature +ptx63 -emit-llvm /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c -o - | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c --check-prefix=NVPTX
# 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 -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/Inputs/include -internal-isystem /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/../../lib/Headers/ -triple nvptx64-nvidia-cuda -target-feature +ptx63 -emit-llvm /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c -o -
# 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/Headers/gpuintrin.c --check-prefix=NVPTX
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c:1112:16: error: NVPTX-NEXT: expected string not found in input
# | // NVPTX-NEXT: [[TMP4:%.*]] = call i64 @llvm.ctlz.i64(i64 [[TMP3]], i1 true)
# | ^
# | <stdin>:496:38: note: scanning from here
# | %3 = load i64, ptr %__below, align 8
# | ^
# | <stdin>:496:38: note: with "TMP3" equal to "%3"
# | %3 = load i64, ptr %__below, align 8
# | ^
# | <stdin>:542:2: note: possible intended match here
# | %14 = call i64 @llvm.ctlz.i64(i64 %13, i1 true)
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c:1729:13: error: undefined variable: LOOP1
# | // NVPTX: [[LOOP1]] = distinct !{[[LOOP1]], [[META2:![0-9]+]]}
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c:1729:36: error: undefined variable: LOOP1
# | // NVPTX: [[LOOP1]] = distinct !{[[LOOP1]], [[META2:![0-9]+]]}
# | ^
# | <stdin>:632:34: note: possible intended match here
# | declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32) #3
# | ^
# |
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/clang/test/Headers/gpuintrin.c
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 491: %2 = load i64, ptr %__below, align 8
# | 492: %tobool = icmp ne i64 %2, 0
# | 493: br i1 %tobool, label %cond.true, label %cond.false
# | 494:
# | 495: cond.true: ; preds = %for.body
# | 496: %3 = load i64, ptr %__below, align 8
# | next:1112'0 X error: no match found
# | next:1112'1 with "TMP3" equal to "%3"
# | 497: %4 = call i64 @llvm.ctlz.i64(i64 %3, i1 false)
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 498: %cast = trunc i64 %4 to i32
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 499: %sub2 = sub nsw i32 63, %cast
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 500: br label %cond.end
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~
# | 501:
# | next:1112'0 ~
# | .
# | .
# | .
# | 537: %cmp12 = icmp ult i32 %11, %12
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 538: br i1 %cmp12, label %for.body13, label %for.end
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 539:
# | next:1112'0 ~
# | 540: for.body13: ; preds = %for.cond11
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 541: %13 = load i64, ptr %__below, align 8
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 542: %14 = call i64 @llvm.ctlz.i64(i64 %13, i1 true)
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | next:1112'2 ? possible intended match
# | 543: %cast14 = trunc i64 %14 to i32
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 544: %iszero = icmp eq i64 %13, 0
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 545: %clzg = select i1 %iszero, i32 0, i32 %cast14
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 546: %sub15 = sub nsw i32 63, %clzg
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 547: %sh_prom16 = zext i32 %sub15 to i64
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | 575:
# | next:1112'0 ~
# | 576: ; Function Attrs: convergent noinline noreturn nounwind optnone
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 577: define internal void @__gpu_exit() #1 {
# | next:1112'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 578: entry:
# | 579: call void @llvm.nvvm.exit()
# | 580: ret void
# | check:1729'0 X error: match failed for invalid pattern
# | check:1729'1 undefined variable: LOOP1
# | check:1729'2 undefined variable: LOOP1
# | 581: }
# | check:1729'0 ~~
# | 582:
# | check:1729'0 ~
# | 583: ; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none)
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 584: declare noundef range(i32 1, -2147483648) i32 @llvm.nvvm.read.ptx.sreg.nctaid.x() #2
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 585:
# | check:1729'0 ~
# | .
# | .
# | .
# | 627:
# | check:1729'0 ~
# | 628: ; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none)
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 629: declare i32 @llvm.cttz.i32(i32, i1 immarg) #2
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 630:
# | check:1729'0 ~
# | 631: ; Function Attrs: convergent nocallback nounwind memory(inaccessiblemem: readwrite)
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 632: declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32) #3
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:1729'3 ? possible intended match
# | 633:
# | check:1729'0 ~
# | 634: ; Function Attrs: convergent nocallback nounwind memory(inaccessiblemem: readwrite)
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 635: declare i32 @llvm.nvvm.vote.ballot.sync(i32, i1) #3
# | check:1729'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 636:
# | check:1729'0 ~
# | 637: ; Function Attrs: convergent nocallback nounwind
# | check:1729'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/185630
More information about the cfe-commits
mailing list