llvm::LLVMConstIntGetZExtValue doesn't accept values with more than 64 bits.
This fixes an LLVM assertion error when compiling libcore with stage1:
src/llvm/include/llvm/ADT/APInt.h:1336:
uint64_t llvm::APInt::getZExtValue() const:
Assertion `getActiveBits() <= 64 && "Too many bits for uint64_t"' failed.
Fixes rebase fallout, makes code correct in presence of 128-bit constants.
This commit includes manual merge conflict resolution changes from a rebase by @est31.
This commit introduces 128-bit integers. Stage 2 builds and produces a working compiler which
understands and supports 128-bit integers throughout.
The general strategy used is to have rustc_i128 module which provides aliases for iu128, equal to
iu64 in stage9 and iu128 later. Since nowhere in rustc we rely on large numbers being supported,
this strategy is good enough to get past the first bootstrap stages to end up with a fully working
128-bit capable compiler.
In order for this strategy to work, number of locations had to be changed to use associated
max_value/min_value instead of MAX/MIN constants as well as the min_value (or was it max_value?)
had to be changed to use xor instead of shift so both 64-bit and 128-bit based consteval works
(former not necessarily producing the right results in stage1).
This commit includes manual merge conflict resolution changes from a rebase by @est31.
appveyor: Attempt to debug flaky test runs
This commit is an attempt to debug #38620 since we're unable to reproduce it
locally. It follows the [advice] of those with AppVeyor to use the `handle.exe`
tool to try to debug what processes have a handle to the file open.
This won't be guaranteed to actually help us, but hopefully it'll diagnose
something at some point?
[advice]: http://help.appveyor.com/discussions/questions/2898
Check *all* errors in LLVMRustArchiveIterator* API
Incrementing the `Archive::child_iterator` fetches and validates the next child.
This can trigger an error, which we previously checked on the *next* call to `LLVMRustArchiveIteratorNext()`.
This means we ignore the last error if we stop iterating halfway through.
This is harmless (we don't access the child, after all) but LLVM 4.0 calls `abort()` if *any* error goes unchecked, even a success value.
This means that basically any rustc invocation that opens an archive and searches through it would die.
The solution implemented here is to change the order of operations, such that
advancing the iterator and fetching the newly-validated iterator happens in the same `Next()` call.
This keeps the error handling behavior as before but ensures all `Error`s get checked.
rustbuild: Move pretty test suites to host-only
In an ongoing effort to optimize the runtime of the Android cross builder this
commit updates the pretty test suites to run only for host platforms, not for
target platforms as well. This means we'll still keep running all the suites but
we'll only run them for configured hosts, not for configured targets. This
notably means that we won't be running these suites on Android or musl targets,
for example.
Use "an" before "i32"
(Minor typo fix.)
Since the word `i32` starts with a vowel, the indefinite article should use "an", not "a" \[[1](http://www.dictionary.com/browse/an)\]. (Previously there was one instance of "an i32" and two instances of "a i32", so at least something is wrong!) Since I believe that "an" is the correct form, I aligned everything with that.
Fix default terminfo code for reset (sg0 -> sgr0)
Resetting the terminal should first try `sgr0` (as per the comment), not
`sg0` which I believe to be a typo.
This will at least fix rustc output in Emacs terminals (e.g., ansi-term)
with `TERM=eterm-color` which does not provide the next fallback `sgr`. In
such a terminal, the final fallback `op` (`\e[39;49`) is used which
resets only colors, not all attributes. This causes all text to be
printed in bold from the first string printed in bold by rustc onwards,
including the terminal prompt and the output from all following commands.
The typo seems to have been introduced by #29999
Add missing apostrophe.
(Minor typo fix.)
The "support" in this case is possessed by the "programmer", and that ownership should be indicated by an apostrophe.
compiletest: Don't limit all suites on Android
On Android we only have one test thread for supposed problems with concurrency
and the remote debugger. Not all of our suites require one concurrency, however,
and suites like compile-fail or pretty can be much faster if they're
parallelized on Travis.
This commit only sets the test threads to one on Android for suites which
actually run code, and other suites aren't tampered with.
rustbuild: Compile rustc twice, not thrice
This commit switches the rustbuild build system to compiling the
compiler twice for a normal bootstrap rather than the historical three
times.
Rust is a bootstrapped language which means that a previous version of
the compiler is used to build the next version of the compiler. Over
time, however, we change many parts of compiler artifacts such as the
metadata format, symbol names, etc. These changes make artifacts from
one compiler incompatible from another compiler. Consequently if a
compiler wants to be able to use some artifacts then it itself must have
compiled the artifacts.
Historically the rustc build system has achieved this by compiling the
compiler three times:
* An older compiler (stage0) is downloaded to kick off the chain.
* This compiler now compiles a new compiler (stage1)
* The stage1 compiler then compiles another compiler (stage2)
* Finally, the stage2 compiler needs libraries to link against, so it
compiles all the libraries again.
This entire process amounts in compiling the compiler three times.
Additionally, this process always guarantees that the Rust source tree
can compile itself because the stage2 compiler (created by a freshly
created compiler) would successfully compile itself again. This
property, ensuring Rust can compile itself, is quite important!
In general, though, this third compilation is not required for general
purpose development on the compiler. The third compiler (stage2) can
reuse the libraries that were created during the second compile. In
other words, the second compilation can produce both a compiler and the
libraries that compiler will use. These artifacts *must* be compatible
due to the way plugins work today anyway, and they were created by the
same source code so they *should* be compatible as well.
So given all that, this commit switches the default build process to
only compile the compiler two times, avoiding this third compilation
by copying artifacts from the previous one. Along the way a new entry in
the Travis matrix was also added to ensure that our full bootstrap can
succeed. This entry does not run tests, though, as it should not be
necessary.
To restore the old behavior of a full bootstrap (three compiles) you can
either pass:
./configure --enable-full-bootstrap
or if you're using config.toml:
[build]
full-bootstrap = true
Overall this will hopefully be an easy 33% win in build times of the
compiler. If we do 33% less work we should be 33% faster! This in turn
should affect cycle times and such on Travis and AppVeyor positively as
well as making it easier to work on the compiler itself.
std: Clamp max read/write sizes on Unix
Turns out that even though all these functions take a `size_t` they don't
actually work that well with anything larger than the maximum value of
`ssize_t`, the return value. Furthermore it looks like OSX rejects any
read/write requests larger than `INT_MAX - 1`. Handle all these cases by just
clamping the maximum size of a read/write on Unix to a platform-specific value.
Closes#38590
travis: Don't use -9 on gzip
I timed this locally and plain old `gzip` took 2m06s while `gzip -9` took a
whopping 6m23s to save a mere 4MB out of 1.2GB. Let's shave a few minutes off
the Android builder by turning down the compression level.
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.
Add some more repos to cargotest
From suggestions at https://users.rust-lang.org/t/what-stable-rust-applications-do-you-use-frequently/7618
This adds some applications which use stable Rust and come with their own lockfiles in their respective trees.
ripgrep, xsv, and bins have 33 unique dependencies between them.
I alphabetized the list by project name because that seems tidier.
r? @brson
appveyor: Attempt to debug flaky test runs
This commit is an attempt to debug #38620 since we're unable to reproduce it
locally. It follows the [advice] of those with AppVeyor to use the `handle.exe`
tool to try to debug what processes have a handle to the file open.
This won't be guaranteed to actually help us, but hopefully it'll diagnose
something at some point?
[advice]: http://help.appveyor.com/discussions/questions/2898
This commit is an attempt to debug #38620 since we're unable to reproduce it
locally. It follows the [advice] of those with AppVeyor to use the `handle.exe`
tool to try to debug what processes have a handle to the file open.
This won't be guaranteed to actually help us, but hopefully it'll diagnose
something at some point?
[advice]: http://help.appveyor.com/discussions/questions/2898
Incrementing the `Archive::child_iterator` fetches and validates the next child.
This can trigger an error, which we previously checked on the *next* call to `LLVMRustArchiveIteratorNext()`.
This means we ignore the last error if we stop iterating halfway through.
This is harmless (we don't access the child, after all) but LLVM 4.0 calls `abort()` if *any* error goes unchecked, even a success value.
This means that basically any rustc invocation that opens an archive and searches through it would die.
The solution implemented here is to change the order of operations, such that
advancing the iterator and fetching the newly-validated iterator happens in the same `Next()` call.
This keeps the error handling behavior as before but ensures all `Error`s get checked.
