| Age | Commit message (Collapse) | Author | Lines |
|---|
|
|
|
rustc now generates the coverage map and can support (limited)
coverage report generation, at the function level.
Example:
$ BUILD=$HOME/rust/build/x86_64-unknown-linux-gnu
$ $BUILD/stage1/bin/rustc -Zinstrument-coverage \
$HOME/rust/src/test/run-make-fulldeps/instrument-coverage/main.rs
$ LLVM_PROFILE_FILE="main.profraw" ./main
called
$ $BUILD/llvm/bin/llvm-profdata merge -sparse main.profraw -o main.profdata
$ $BUILD/llvm/bin/llvm-cov show --instr-profile=main.profdata main
1| 1|pub fn will_be_called() {
2| 1| println!("called");
3| 1|}
4| |
5| 0|pub fn will_not_be_called() {
6| 0| println!("should not have been called");
7| 0|}
8| |
9| 1|fn main() {
10| 1| let less = 1;
11| 1| let more = 100;
12| 1|
13| 1| if less < more {
14| 1| will_be_called();
15| 1| } else {
16| 1| will_not_be_called();
17| 1| }
18| 1|}
|
|
|
|
This commit adds a new lint - `improper_ctypes_definitions` - which
functions identically to `improper_ctypes`, but on `extern "C" fn`
definitions (as opposed to `improper_ctypes`'s `extern "C" {}`
declarations).
Signed-off-by: David Wood <david@davidtw.co>
|
|
|
|
|
|
rustllvm relies on the `LLVMRustStringWriteImpl` symbol existing, but
this symbol was previously defined in a *downstream* crate
(rustc_codegen_llvm, which depends on rustc_llvm.
While this somehow worked under the old 'separate bootstrap step for
codegen' scheme, it meant that rustc_llvm could not actually be built by
itself, since it relied linking to the downstream rustc_codegen_llvm
crate.
Now that librustc_codegen_llvm is just a normal crate, we actually try
to build a standalone rustc_llvm when we run tests. This commit moves
`LLVMRustStringWriteImpl` into rustc_llvm (technically the rustllvm
directory, which has its contents built by rustc_llvm). This ensures
that we can build each crate in the graph by itself, without requiring
that any downstream crates be linked in as well.
|
|
rustbuild
Remove some random unnecessary lint `allow`s
|
|
|
|
Now that procedural macros no longer link transitively to libsyntax,
this shouldn't be needed any more! This commit is an experiment in
removing all dynamic libraries from rustc except for librustc_driver
itself. Let's see how far we can get with that!
|
|
|
|
|
|
|
|
|
|
|
|
|
|
Co-authored-by: nikomatsakis
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
Note this isn't useful, yet. More changes will be necessary to be able to
actually codegen for this machine. As such, it is not enabled by default.
This patch is on its own for the benefit of the reviewers.
|
|
|
|
This permits easier iteration without having to worry about warnings
being denied.
Fixes #49517
|
|
|
|
It looks like LLVM also removed it in llvm-mirror/llvm@f45adc29d in favor of the
name "GNU64". This was added in the thought that we'd need such a variant when
adding mips64 support but we ended up not needing it! For now let's just
removing the various support on the Rust side of things.
|
|
Refs #46437 as it also removes LLVMRustWriteDebugLocToString()
|
|
|
|
|
|
|
|
|
|
Fixes #41701.
|
|
|
|
This commit bumps the master branch's version to 1.21.0 and also updates the
bootstrap compiler from the freshly minted beta release.
|
|
|
|
Enable wasm LLVM backend
Enables compilation to WebAssembly with the LLVM backend using the target triple "wasm32-unknown-unknown". This is the beginning of my work on #38804.
**edit:** The new new target is now wasm32-experimental-emscripten instead of wasm32-unknown-unknown.
|
|
|
|
The new target is wasm32-experimental-emscripten. Adds a new
configuration option to opt in to building experimental LLVM backends
such as the WebAssembly backend. The target name was chosen to be
similar to the existing wasm32-unknown-emscripten target so that the
build and tests would work with minimal other code changes. When/if the
new target replaces the old target, simply renaming it should just work.
|
|
These are now no longer necessary with `-Z force-unstable-if-unmarked`
|
|
This requires an updated LLVM with D31999 and D32000 to build libcore.
A basic hello world builds and runs successfully on the hexagon simulator.
|
|
|
|
These are all now no longer needed that we've only got rustbuild in tree.
|
|
Remove not(stage0) from deny(warnings)
Historically this was done to accommodate bugs in lints, but there hasn't been a
bug in a lint since this feature was added which the warnings affected. Let's
completely purge warnings from all our stages by denying warnings in all stages.
This will also assist in tracking down `stage0` code to be removed whenever
we're updating the bootstrap compiler.
|
|
As per the LLVM style guide, use CamelCase for all locals and classes,
and camelCase for all non-FFI functions.
Also, make names of variables of commonly used types more consistent.
Fixes #38688.
|
|
Historically this was done to accommodate bugs in lints, but there hasn't been a
bug in a lint since this feature was added which the warnings affected. Let's
completely purge warnings from all our stages by denying warnings in all stages.
This will also assist in tracking down `stage0` code to be removed whenever
we're updating the bootstrap compiler.
|
|
PTX support, take 2
- You can generate PTX using `--emit=asm` and the right (custom) target. Which
then you can run on a NVIDIA GPU.
- You can compile `core` to PTX. [Xargo] also works and it can compile some
other crates like `collections` (but I doubt all of those make sense on a GPU)
[Xargo]: https://github.com/japaric/xargo
- You can create "global" functions, which can be "called" by the host, using
the `"ptx-kernel"` ABI, e.g. `extern "ptx-kernel" fn kernel() { .. }`. Every
other function is a "device" function and can only be called by the GPU.
- Intrinsics like `__syncthreads()` and `blockIdx.x` are available as
`"platform-intrinsics"`. These intrinsics are *not* in the `core` crate but
any Rust user can create "bindings" to them using an `extern
"platform-intrinsics"` block. See example at the end.
- Trying to emit PTX with `-g` (debuginfo); you get an LLVM error. But I don't
think PTX can contain debuginfo anyway so `-g` should be ignored and a warning
should be printed ("`-g` doesn't work with this target" or something).
- "Single source" support. You *can't* write a single source file that contains
both host and device code. I think that should be possible to implement that
outside the compiler using compiler plugins / build scripts.
- The equivalent to CUDA `__shared__` which it's used to declare memory that's
shared between the threads of the same block. This could be implemented using
attributes: `#[shared] static mut SCRATCH_MEMORY: [f32; 64]` but hasn't been
implemented yet.
- Built-in targets. This PR doesn't add targets to the compiler just yet but one
can create custom targets to be able to emit PTX code (see the example at the
end). The idea is to have people experiment with this feature before
committing to it (built-in targets are "insta-stable")
- All functions must be "inlined". IOW, the `.rlib` must always contain the LLVM
bitcode of all the functions of the crate it was produced from. Otherwise, you
end with "undefined references" in the final PTX code but you won't get *any*
linker error because no linker is involved. IOW, you'll hit a runtime error
when loading the PTX into the GPU. The workaround is to use `#[inline]` on
non-generic functions and to never use `#[inline(never)]` but this may not
always be possible because e.g. you could be relying on third party code.
- Should `--emit=asm` generate a `.ptx` file instead of a `.s` file?
TL;DR Use Xargo to turn a crate into a PTX module (a `.s` file). Then pass that
PTX module, as a string, to the GPU and run it.
The full code is in [this repository]. This section gives an overview of how to
run Rust code on a NVIDIA GPU.
[this repository]: https://github.com/japaric/cuda
- Create a custom target. Here's the 64-bit NVPTX target (NOTE: the comments
are not valid because this is supposed to be a JSON file; remove them before
you use this file):
``` js
// nvptx64-nvidia-cuda.json
{
"arch": "nvptx64", // matches LLVM
"cpu": "sm_20", // "oldest" compute capability supported by LLVM
"data-layout": "e-i64:64-v16:16-v32:32-n16:32:64",
"llvm-target": "nvptx64-nvidia-cuda",
"max-atomic-width": 0, // LLVM errors with any other value :-(
"os": "cuda", // matches LLVM
"panic-strategy": "abort",
"target-endian": "little",
"target-pointer-width": "64",
"target-vendor": "nvidia", // matches LLVM -- not required
}
```
(There's a 32-bit target specification in the linked repository)
- Write a kernel
``` rust
extern "platform-intrinsic" {
fn nvptx_block_dim_x() -> i32;
fn nvptx_block_idx_x() -> i32;
fn nvptx_thread_idx_x() -> i32;
}
/// Copies an array of `n` floating point numbers from `src` to `dst`
pub unsafe extern "ptx-kernel" fn memcpy(dst: *mut f32,
src: *const f32,
n: usize) {
let i = (nvptx_block_dim_x() as isize)
.wrapping_mul(nvptx_block_idx_x() as isize)
.wrapping_add(nvptx_thread_idx_x() as isize);
if (i as usize) < n {
*dst.offset(i) = *src.offset(i);
}
}
```
- Emit PTX code
```
$ xargo rustc --target nvptx64-nvidia-cuda --release -- --emit=asm
Compiling core v0.0.0 (file://..)
(..)
Compiling nvptx-builtins v0.1.0 (https://github.com/japaric/nvptx-builtins)
Compiling kernel v0.1.0
$ cat target/nvptx64-nvidia-cuda/release/deps/kernel-*.s
//
// Generated by LLVM NVPTX Back-End
//
.version 3.2
.target sm_20
.address_size 64
// .globl memcpy
.visible .entry memcpy(
.param .u64 memcpy_param_0,
.param .u64 memcpy_param_1,
.param .u64 memcpy_param_2
)
{
.reg .pred %p<2>;
.reg .s32 %r<5>;
.reg .s64 %rd<12>;
ld.param.u64 %rd7, [memcpy_param_2];
mov.u32 %r1, %ntid.x;
mov.u32 %r2, %ctaid.x;
mul.wide.s32 %rd8, %r2, %r1;
mov.u32 %r3, %tid.x;
cvt.s64.s32 %rd9, %r3;
add.s64 %rd10, %rd9, %rd8;
setp.ge.u64 %p1, %rd10, %rd7;
@%p1 bra LBB0_2;
ld.param.u64 %rd3, [memcpy_param_0];
ld.param.u64 %rd4, [memcpy_param_1];
cvta.to.global.u64 %rd5, %rd4;
cvta.to.global.u64 %rd6, %rd3;
shl.b64 %rd11, %rd10, 2;
add.s64 %rd1, %rd6, %rd11;
add.s64 %rd2, %rd5, %rd11;
ld.global.u32 %r4, [%rd2];
st.global.u32 [%rd1], %r4;
LBB0_2:
ret;
}
```
- Run it on the GPU
``` rust
// `kernel.ptx` is the `*.s` file we got in the previous step
const KERNEL: &'static str = include_str!("kernel.ptx");
driver::initialize()?;
let device = Device(0)?;
let ctx = device.create_context()?;
let module = ctx.load_module(KERNEL)?;
let kernel = module.function("memcpy")?;
let h_a: Vec<f32> = /* create some random data */;
let h_b = vec![0.; N];
let d_a = driver::allocate(bytes)?;
let d_b = driver::allocate(bytes)?;
// Copy from host to GPU
driver::copy(h_a, d_a)?;
// Run `memcpy` on the GPU
kernel.launch(d_b, d_a, N)?;
// Copy from GPU to host
driver::copy(d_b, h_b)?;
// Verify
assert_eq!(h_a, h_b);
// `d_a`, `d_b`, `h_a`, `h_b` are dropped/freed here
```
---
cc @alexcrichton @brson @rkruppe
> What has changed since #34195?
- `core` now can be compiled into PTX. Which makes it very easy to turn `no_std`
crates into "kernels" with the help of Xargo.
- There's now a way, the `"ptx-kernel"` ABI, to generate "global" functions. The
old PR required a manual step (it was hack) to "convert" "device" functions
into "global" functions. (Only "global" functions can be launched by the host)
- Everything is unstable. There are not "insta stable" built-in targets this
time (\*). The users have to use a custom target to experiment with this
feature. Also, PTX instrinsics, like `__syncthreads` and `blockIdx.x`, are now
implemented as `"platform-intrinsics"` so they no longer live in the `core`
crate.
(\*) I'd actually like to have in-tree targets because that makes this target
more discoverable, removes the need to lug around .json files, etc.
However, bundling a target with the compiler immediately puts it in the path
towards stabilization. Which gives us just two cycles to find and fix any
problem with the target specification. Afterwards, it becomes hard to tweak
the specification because that could be a breaking change.
A possible solution could be "unstable built-in targets". Basically, to use an
unstable target, you'll have to also pass `-Z unstable-options` to the compiler.
And unstable targets, being unstable, wouldn't be available on stable.
> Why should this be merged?
- To let people experiment with the feature out of tree. Having easy access to
the feature (in every nightly) allows this. I also think that, as it is, it
should be possible to start prototyping type-safe single source support using
build scripts, macros and/or plugins.
- It's a straightforward implementation. No different that adding support for
any other architecture.
|
