| Issue |
138034
|
| Summary |
[NVPTX] Failure in PTX when generating code for `i128` regresses `libc` tests
|
| Labels |
libc
|
| Assignees |
|
| Reporter |
jhuber6
|
A recent patch seems to have changed the behavior of `i128` on NVPTX targets. This is causing failures as shown in https://godbolt.org/z/cPq69rs1T. This is the reproducer.
```llvm
target triple = "nvptx64-nvidia-cuda"
define i128 @foo() {
entry:
br label %while.cond
while.cond: ; preds = %while.cond, %entry
%0 = load i8, ptr null, align 1
%conv = zext i8 %0 to i128
store i128 %conv, ptr null, align 16
br label %while.cond
}
```
```console
$ clang --target=nvptx64-nvidia-cuda -march=sm_89 bug.bc
```
This will generate an error and the following PTX.
```asm
//
// Generated by LLVM NVPTX Back-End
//
.version 8.7
.target sm_89
.address_size 64
// .globl foo // -- Begin function foo
// @foo
.visible .func (.param .align 16 .b8 func_retval0[16]) foo()
{
.reg .b64 %rd<4>;
// %bb.0: // %entry
bra.uni $L__BB0_1;
$L__BB0_1: // %while.cond
// =>This Inner Loop Header: Depth=1
mov.b64 %rd1, 0;
ld.v2.u4 {%rd2, %rd3}, [%rd1];
st.v2.u64 [%rd1], {%rd2, %rd3};
bra.uni $L__BB0_1;
// -- End function
}
```
```
ptxas reduced.s, line 20; error : Unexpected instruction types specified for 'ld'
ptxas fatal : Ptx assembly aborted due to errors
clang: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 21.0.0git
Target: nvptx64-nvidia-cuda
Thread model: posix
InstalledDir: /home/jhuber/Documents/llvm/clang/bin
Build config: +assertions
clang: note: diagnostic msg: Error generating preprocessed source(s) - no preprocessable inputs.
```
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
