STR
$ cargo new --lib kernel && cd $_
$ edit src/lib.rs && cat $_
#![no_std]
extern {
fn bar();
}
fn foo() {
unsafe {
bar()
}
}$ cat nvptx64-nvidia-cuda.json
{
"arch": "nvptx64",
"cpu": "sm_20",
"data-layout": "e-i64:64-v16:16-v32:32-n16:32:64",
"llvm-target": "nvptx64-nvidia-cuda",
"max-atomic-width": 0,
"os": "cuda",
"panic-strategy": "abort",
"target-endian": "little",
"target-pointer-width": "64"
}$ edit Cargo.toml && cat $_
[profile.dev]
debug = false # cf. rust-lang/rust#38785
$ cargo install xargo --vers 0.3.0 || true
$ xargo rustc --target nvptx64-nvidia-cuda -- --emit=asm
$ cat $(find target/nvptx64-nvidia-cuda/debug -name '*.s')
.version 3.2
.target sm_20
.address_size 64
.extern .func bar
()
;
.func _ZN6kernel3foo17h9be095784726d4d0E()
{
bra.uni LBB0_1;
LBB0_1:
{ // callseq 0
.reg .b32 temp_param_reg;
call.uni
bar,
(
);
} // callseq 0
bra.uni LBB0_2;
LBB0_2:
ret;
}
Note that bar is called by foo but it's not defined in the PTX module. Trying to load this PTX module, with cuModuleLoadData, generates a runtime error: CUDA_ERROR_NO_BINARY_FOR_GPU.
Expected behavior
Other targets that produce an executable would fail to compile this with an "undefined reference" error which is actually raised by the linker. The NVPTX targets don't involve a linker but they shouldn't produce incomplete PTX modules like the one above; instead the should raise an error like the other targets do.
STR
Note that
baris called byfoobut it's not defined in the PTX module. Trying to load this PTX module, withcuModuleLoadData, generates a runtime error:CUDA_ERROR_NO_BINARY_FOR_GPU.Expected behavior
Other targets that produce an executable would fail to compile this with an "undefined reference" error which is actually raised by the linker. The NVPTX targets don't involve a linker but they shouldn't produce incomplete PTX modules like the one above; instead the should raise an error like the other targets do.