Is this a codegen bug in ptx kernels?

kjetilkjeka · February 25, 2022, 1:14pm

I think I found a bug in the code generated for nvptx kernels this week. It is related to passing structs to kernels as a function argument.

I already posted it in the tracking issue, but I believe it's quite invisible as the tracking issue is so old.

Can anyone help me confirm that this is in fact a bug, and that I am not misunderstanding something.

github.com/rust-lang/rust

Tracking issue for the "ptx-kernel" ABI

opened 11:49PM - 02 Jan 17 UTC

japaric

T-compiler B-unstable O-NVPTX C-tracking-issue

Introduced in #38559. Feature gate: `abi_ptx` This ABI is used to generate… ["global"] functions. All the other functions that *don't* use the "ptx-kernel" ABI are ["device"] functions. ["global"]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#global ["device"]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#device-function-qualifier The host can only make the device, the GPU, execute global functions, not device functions. IOW, a PTX module that only contains device functions is basically unusable. --- Note that this ABI *can't* be emulated using naked functions. ``` rust #![feature(abi_ptx)] #![feature(naked_functions)] #![no_std] #[naked] #[no_mangle] fn foo() {} #[no_mangle] fn bar() {} #[no_mangle] extern "ptx-kernel" fn baz() {} ``` Produces: ``` .version 3.2 .target sm_20 .address_size 64 .func foo() { .local .align 8 .b8 __local_depot0[8]; .reg .b64 %SP; .reg .b64 %SPL; mov.u64 %SPL, __local_depot0; bra.uni LBB0_1; LBB0_1: ret; } .func bar() { .local .align 8 .b8 __local_depot1[8]; .reg .b64 %SP; .reg .b64 %SPL; mov.u64 %SPL, __local_depot1; bra.uni LBB1_1; LBB1_1: ret; } .entry baz() { .local .align 8 .b8 __local_depot2[8]; .reg .b64 %SP; .reg .b64 %SPL; mov.u64 %SPL, __local_depot2; bra.uni LBB2_1; LBB2_1: ret; } ``` Note that only `baz` uses the `.entry` directive.

bjorn3 · February 25, 2022, 4:14pm

cc @RDambrosio016

system · May 26, 2022, 4:15pm

This topic was automatically closed 90 days after the last reply. New replies are no longer allowed.

Topic		Replies	Views
[Solved] Passing cuda-arch when compiling for nvptx64-nvidia-cuda compiler	5	942	May 14, 2022
Emitting ptrauth intrinsics in rustc_codegen_llvm for function calls compiler	2	529	June 10, 2021
New target support compiler	9	1036	March 25, 2019
AVX512 intrinsics / stdsimd	1	1020	August 27, 2020
Global symbols from statically linked system libraries libs	4	486	February 29, 2024

Is this a codegen bug in ptx kernels?

Related Topics