Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

NVPTX: No "undefined reference" error is raised when it should be #38786

Closed
japaric opened this issue Jan 2, 2017 · 0 comments
Closed

NVPTX: No "undefined reference" error is raised when it should be #38786

japaric opened this issue Jan 2, 2017 · 0 comments
Labels
C-bug Category: This is a bug. O-NVPTX Target: the NVPTX LLVM backend for running rust on GPUs, https://llvm.org/docs/NVPTXUsage.html

Comments

@japaric
Copy link
Member

japaric commented Jan 2, 2017

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.

@japaric japaric added the O-NVPTX Target: the NVPTX LLVM backend for running rust on GPUs, https://llvm.org/docs/NVPTXUsage.html label Jan 2, 2017
@Mark-Simulacrum Mark-Simulacrum added the C-bug Category: This is a bug. label Jul 26, 2017
bors added a commit that referenced this issue Feb 1, 2019
NVPTX target specification

This change adds a built-in `nvptx64-nvidia-cuda` GPGPU no-std target specification and a basic PTX assembly smoke tests.

The approach is taken here and the target spec is based on `ptx-linker`, a project started about 1.5 years ago. Key feature: bitcode object files being linked with LTO into the final module on the linker's side.

Prior to this change, the linker used a `ld` linker-flavor, but I think, having the special CLI convention is a more reliable way.

Questions about further progress on reliable CUDA workflow with Rust:
1. Is it possible to create a test suite `codegen-asm` to verify end-to-end integration with LLVM backend?
1. How would it be better to organise no-std `compile-fail` tests: add `#![no_std]` where possible and mark others as `ignore-nvptx` directive, or alternatively, introduce `compile-fail-no-std` test suite?
1. Can we have the `ptx-linker` eventually be integrated as `rls` or `clippy`? Hopefully, this should allow to statically link against LLVM used in Rust and get rid of the [current hacky solution](https://github.com/denzp/rustc-llvm-proxy).
1. Am I missing some methods from `rustc_codegen_ssa::back::linker::Linker` that can be useful for bitcode-only linking?

Currently, there are no major public CUDA projects written in Rust I'm aware of, but I'm expecting to have a built-in target will create a solid foundation for further experiments and awesome crates.

Related to #38789
Fixes #38787
Fixes #38786
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
C-bug Category: This is a bug. O-NVPTX Target: the NVPTX LLVM backend for running rust on GPUs, https://llvm.org/docs/NVPTXUsage.html
Projects
None yet
Development

No branches or pull requests

2 participants