NVPTX support for new asm!

[westernmagic]

This PR implements the new asm! syntax for the nvptx64-nvidia-cuda target.

r? @Amanieu

[rust-highfive]

Thanks for the pull request, and welcome! The Rust team is excited to review your changes, and you should hear from @davidtwco (or someone else) soon.

If any changes to this PR are deemed necessary, please add them as extra commits. This ensures that the reviewer can see what has changed since they last reviewed the code. Due to the way GitHub handles out-of-date commits, this should also make it reasonably obvious what issues have or haven't been addressed. Large or tricky changes may require several passes of review and changes.

Please see the contribution instructions for more information.

[Amanieu]

You are missing an nvptx-modifiers.rs test which tests that register names are rendered properly in the output asm.

[westernmagic]

You are missing an nvptx-modifiers.rs test which tests that register names are rendered properly in the output asm.

As discussed on zulip, NVPTX does not support any modifiers, therefore no tests are needed.

[Amanieu]

The code looks good!

Can you update the specification in src/doc/unstable-book/src/library-features/asm.md to include the new register classes?

[westernmagic]

This is false, I checked on NVPTX and values are not zero-extended. The upper bits are undefined. https://rust.godbolt.org/z/9yU64A

They are, implicitly:

A destination register wider than the specified type may be used. The value loaded is sign-extended to the destination register width for signed integers, and is zero-extended to the destination register width for unsigned and bit-size types. See Table 25 for a description of these relaxed type-checking rules.

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#operand-size-exceeding-instruction-type-size

Since clang gives the registers bit sizes, they are all zero extended in this implementation.

Here's a small clang experiment as proof:

#include <cstdint>
#include <cstdio>

__global__ void foo(uint32_t x, uint64_t * y) {
	asm volatile (
		"mov.u64 %0, %1;"
		: "=l"(*y) : "l"(x)
	);
}

__host__ int main() {
	uint64_t y = UINT64_MAX;
	uint64_t * yd;
	cudaMalloc(&yd, sizeof(uint64_t));
	cudaMemcpy(yd, &y, sizeof(uint64_t), cudaMemcpyHostToDevice);
	foo<<<1, 1>>>(0xBEEF, yd);
	cudaMemcpy(&y, yd, sizeof(uint64_t), cudaMemcpyDeviceToHost);
	cudaFree(yd);
	printf("%lx\n", y);
}

If you wish, I can add the links I referenced to the documentation

[Amanieu]

Again, look at the assembly output from godbolt: https://rust.godbolt.org/z/-8yrgo

.visible .func  (.param .b64 func_retval0) foo(
        .param .b64 foo_param_0
)
{

        ld.param.u64    %rd2, [foo_param_0];
        mov.u64 %rd1, %rd2
        st.param.b64    [func_retval0+0], %rd1;
        ret;

}

The input value is not zero-extended anywhere in this code.

[westernmagic]

Ah, I see now. Looking at this though, the problem seems to be more complicated though: https://rust.godbolt.org/z/ojcbJV

[Amanieu]

This is the same as every other architecture: when you put a value into a register that is smaller than the register size, the upper bits are UNDEFINED. Since they are undefined, their exact value depends on what the optimizer decides to do.

[Amanieu]

@bors r+

[bors]

📌 Commit 8706f76020e84863e8afc4e25dc014decc0f5a2f has been approved by Amanieu

[bors]

☔ The latest upstream changes (presumably #72516) made this pull request unmergeable. Please resolve the merge conflicts.

[Amanieu]

@bors r+

[bors]

📌 Commit e18054d has been approved by Amanieu