[NVPTX] Failure in PTX when generating code for `i128` regresses `libc` tests

[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.

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
}

$ clang --target=nvptx64-nvidia-cuda -march=sm_89 bug.bc

This will generate an error and the following PTX.

//
// 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.

[llvmbot]

@llvm/issue-subscribers-libc

Author: Joseph Huber (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.

//
// 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&lt;4&gt;;

// %bb.0:                               // %entry
	bra.uni 	$L__BB0_1;
$L__BB0_1:                              // %while.cond
                                        // =&gt;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.

[Artem-B]

Ouch. How did we end up with u4?
https://godbolt.org/z/baMhM7W6E

@AlexMaclean ^^^ did we miss a typo somewhere?

[Artem-B]

Looks like it may be due to a somewhat broken input which we didn't expect.
In case of the reproducer, we end up trying to lower LDV_i64_v2, but the TotalSize is just 8, not 128.
Looks like we incorrectly convert loading of i8 into an extending load. Or assuming too much about the underlying value we're loading.

[jhuber6]

Looks like it may be due to a somewhat broken input which we didn't expect. In case of the reproducer, we end up trying to lower LDV_i64_v2, but the TotalSize is just 8, not 128. Looks like we incorrectly convert loading of i8 into an extending load. Or assuming too much about the underlying value we're loading.

Could be a quirk of the reduced output, but it shows up in a lot of the libc tests.

[AlexMaclean]

I'm looking at this now as well. We are using the resulting value type instead of the memory VT when lowering load/stores to our custom vector nodes.