From d3903d5f9c4e58cc3fa256ec5be52717b84e6308 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Sat, 19 Jan 2019 21:59:34 +0100 Subject: [PATCH 1/8] Create `nvptx64-nvidia-cuda` target specification --- src/bootstrap/lib.rs | 1 + src/bootstrap/sanity.rs | 4 +- src/librustc/ty/context.rs | 6 + src/librustc_codegen_ssa/back/link.rs | 1 + src/librustc_codegen_ssa/back/linker.rs | 125 ++++++++++++++++++ src/librustc_codegen_utils/lib.rs | 1 + src/librustc_codegen_utils/symbol_names.rs | 57 +++++--- src/librustc_target/spec/mod.rs | 4 + .../spec/nvptx64_nvidia_cuda.rs | 73 ++++++++++ src/test/run-make/nvptx-binary-crate/Makefile | 9 ++ src/test/run-make/nvptx-binary-crate/main.rs | 28 ++++ src/test/run-make/nvptx-dylib-crate/Makefile | 10 ++ src/test/run-make/nvptx-dylib-crate/dep.rs | 14 ++ src/test/run-make/nvptx-dylib-crate/kernel.rs | 67 ++++++++++ src/test/run-make/nvptx-emit-asm/Makefile | 9 ++ src/test/run-make/nvptx-emit-asm/kernel.rs | 41 ++++++ 16 files changed, 433 insertions(+), 17 deletions(-) create mode 100644 src/librustc_target/spec/nvptx64_nvidia_cuda.rs create mode 100644 src/test/run-make/nvptx-binary-crate/Makefile create mode 100644 src/test/run-make/nvptx-binary-crate/main.rs create mode 100644 src/test/run-make/nvptx-dylib-crate/Makefile create mode 100644 src/test/run-make/nvptx-dylib-crate/dep.rs create mode 100644 src/test/run-make/nvptx-dylib-crate/kernel.rs create mode 100644 src/test/run-make/nvptx-emit-asm/Makefile create mode 100644 src/test/run-make/nvptx-emit-asm/kernel.rs diff --git a/src/bootstrap/lib.rs b/src/bootstrap/lib.rs index 37451a74dfad6..b0bbf2395e2fc 100644 --- a/src/bootstrap/lib.rs +++ b/src/bootstrap/lib.rs @@ -831,6 +831,7 @@ impl Build { !target.contains("msvc") && !target.contains("emscripten") && !target.contains("wasm32") && + !target.contains("nvptx") && !target.contains("fuchsia") { Some(self.cc(target)) } else { diff --git a/src/bootstrap/sanity.rs b/src/bootstrap/sanity.rs index fe547a6b151c2..ff4fb85bbfad3 100644 --- a/src/bootstrap/sanity.rs +++ b/src/bootstrap/sanity.rs @@ -156,7 +156,7 @@ pub fn check(build: &mut Build) { panic!("the iOS target is only supported on macOS"); } - if target.contains("-none-") { + if target.contains("-none-") || target.contains("nvptx") { if build.no_std(*target).is_none() { let target = build.config.target_config.entry(target.clone()) .or_default(); @@ -165,7 +165,7 @@ pub fn check(build: &mut Build) { } if build.no_std(*target) == Some(false) { - panic!("All the *-none-* targets are no-std targets") + panic!("All the *-none-* and nvptx* targets are no-std targets") } } diff --git a/src/librustc/ty/context.rs b/src/librustc/ty/context.rs index 4c8f81411163c..c126bff42582a 100644 --- a/src/librustc/ty/context.rs +++ b/src/librustc/ty/context.rs @@ -1675,6 +1675,12 @@ impl<'a, 'gcx, 'tcx> TyCtxt<'a, 'gcx, 'tcx> { } false } + + /// Determine whether identifiers in the assembly have strict naming rules. + /// Currently, only NVPTX* targets need it. + pub fn has_strict_asm_symbol_naming(&self) -> bool { + self.gcx.sess.target.target.arch.contains("nvptx") + } } impl<'a, 'tcx> TyCtxt<'a, 'tcx, 'tcx> { diff --git a/src/librustc_codegen_ssa/back/link.rs b/src/librustc_codegen_ssa/back/link.rs index d03bb0a3d73a4..2a5ecf9a0593f 100644 --- a/src/librustc_codegen_ssa/back/link.rs +++ b/src/librustc_codegen_ssa/back/link.rs @@ -149,6 +149,7 @@ pub fn linker_and_flavor(sess: &Session) -> (PathBuf, LinkerFlavor) { LinkerFlavor::Ld => "ld", LinkerFlavor::Msvc => "link.exe", LinkerFlavor::Lld(_) => "lld", + LinkerFlavor::PtxLinker => "rust-ptx-linker", }), flavor)), (Some(linker), None) => { let stem = if linker.extension().and_then(|ext| ext.to_str()) == Some("exe") { diff --git a/src/librustc_codegen_ssa/back/linker.rs b/src/librustc_codegen_ssa/back/linker.rs index ad61f8f01d8d4..5e9aeed7107ac 100644 --- a/src/librustc_codegen_ssa/back/linker.rs +++ b/src/librustc_codegen_ssa/back/linker.rs @@ -83,6 +83,10 @@ impl LinkerInfo { LinkerFlavor::Lld(LldFlavor::Wasm) => { Box::new(WasmLd::new(cmd, sess, self)) as Box } + + LinkerFlavor::PtxLinker => { + Box::new(PtxLinker { cmd, sess }) as Box + } } } } @@ -1080,3 +1084,124 @@ fn exported_symbols(tcx: TyCtxt, crate_type: CrateType) -> Vec { symbols } + +/// Much simplified and explicit CLI for the NVPTX linker. The linker operates +/// with bitcode and uses LLVM backend to generate a PTX assembly. +pub struct PtxLinker<'a> { + cmd: Command, + sess: &'a Session, +} + +impl<'a> Linker for PtxLinker<'a> { + fn link_rlib(&mut self, path: &Path) { + self.cmd.arg("--rlib").arg(path); + } + + fn link_whole_rlib(&mut self, path: &Path) { + self.cmd.arg("--rlib").arg(path); + } + + fn include_path(&mut self, path: &Path) { + self.cmd.arg("-L").arg(path); + } + + fn debuginfo(&mut self) { + self.cmd.arg("--debug"); + } + + fn add_object(&mut self, path: &Path) { + self.cmd.arg("--bitcode").arg(path); + } + + fn args(&mut self, args: &[String]) { + self.cmd.args(args); + } + + fn optimize(&mut self) { + self.cmd.arg(match self.sess.opts.optimize { + OptLevel::No => "-O0", + OptLevel::Less => "-O1", + OptLevel::Default => "-O2", + OptLevel::Aggressive => "-O3", + OptLevel::Size => "-Os", + OptLevel::SizeMin => "-Os" + }); + } + + fn output_filename(&mut self, path: &Path) { + self.cmd.arg("-o").arg(path); + } + + fn finalize(&mut self) -> Command { + ::std::mem::replace(&mut self.cmd, Command::new("")) + } + + fn link_dylib(&mut self, _lib: &str) { + panic!("external dylibs not supported") + } + + fn link_rust_dylib(&mut self, _lib: &str, _path: &Path) { + panic!("external dylibs not supported") + } + + fn link_staticlib(&mut self, _lib: &str) { + panic!("staticlibs not supported") + } + + fn link_whole_staticlib(&mut self, _lib: &str, _search_path: &[PathBuf]) { + panic!("staticlibs not supported") + } + + fn framework_path(&mut self, _path: &Path) { + panic!("frameworks not supported") + } + + fn link_framework(&mut self, _framework: &str) { + panic!("frameworks not supported") + } + + fn position_independent_executable(&mut self) { + } + + fn full_relro(&mut self) { + } + + fn partial_relro(&mut self) { + } + + fn no_relro(&mut self) { + } + + fn build_static_executable(&mut self) { + } + + fn gc_sections(&mut self, _keep_metadata: bool) { + } + + fn pgo_gen(&mut self) { + } + + fn no_default_libraries(&mut self) { + } + + fn build_dylib(&mut self, _out_filename: &Path) { + } + + fn export_symbols(&mut self, _tmpdir: &Path, _crate_type: CrateType) { + } + + fn subsystem(&mut self, _subsystem: &str) { + } + + fn no_position_independent_executable(&mut self) { + } + + fn group_start(&mut self) { + } + + fn group_end(&mut self) { + } + + fn cross_lang_lto(&mut self) { + } +} diff --git a/src/librustc_codegen_utils/lib.rs b/src/librustc_codegen_utils/lib.rs index 1f590d46ed8c6..8e96f98540117 100644 --- a/src/librustc_codegen_utils/lib.rs +++ b/src/librustc_codegen_utils/lib.rs @@ -12,6 +12,7 @@ #![feature(nll)] #![allow(unused_attributes)] #![feature(rustc_diagnostic_macros)] +#![feature(in_band_lifetimes)] #![recursion_limit="256"] diff --git a/src/librustc_codegen_utils/symbol_names.rs b/src/librustc_codegen_utils/symbol_names.rs index 9267f14f24234..f2014f7421289 100644 --- a/src/librustc_codegen_utils/symbol_names.rs +++ b/src/librustc_codegen_utils/symbol_names.rs @@ -103,7 +103,7 @@ use rustc_mir::monomorphize::Instance; use syntax_pos::symbol::Symbol; -use std::fmt::Write; +use std::fmt::{self, Write}; use std::mem::discriminant; pub fn provide(providers: &mut Providers) { @@ -221,7 +221,7 @@ fn get_symbol_hash<'a, 'tcx>( } fn def_symbol_name<'a, 'tcx>(tcx: TyCtxt<'a, 'tcx, 'tcx>, def_id: DefId) -> ty::SymbolName { - let mut buffer = SymbolPathBuffer::new(); + let mut buffer = SymbolPathBuffer::new(tcx); item_path::with_forced_absolute_paths(|| { tcx.push_item_path(&mut buffer, def_id, false); }); @@ -317,7 +317,7 @@ fn compute_symbol_name<'a, 'tcx>(tcx: TyCtxt<'a, 'tcx, 'tcx>, instance: Instance let hash = get_symbol_hash(tcx, def_id, instance, instance_ty, substs); - let mut buf = SymbolPathBuffer::from_interned(tcx.def_symbol_name(def_id)); + let mut buf = SymbolPathBuffer::from_interned(tcx.def_symbol_name(def_id), tcx); if instance.is_vtable_shim() { buf.push("{{vtable-shim}}"); @@ -339,26 +339,28 @@ fn compute_symbol_name<'a, 'tcx>(tcx: TyCtxt<'a, 'tcx, 'tcx>, instance: Instance // // To be able to work on all platforms and get *some* reasonable output, we // use C++ name-mangling. -#[derive(Debug)] -struct SymbolPathBuffer { +struct SymbolPathBuffer<'a, 'tcx> { + tcx: TyCtxt<'a, 'tcx, 'tcx>, result: String, temp_buf: String, } -impl SymbolPathBuffer { - fn new() -> Self { +impl SymbolPathBuffer<'a, 'tcx> { + fn new(tcx: TyCtxt<'a, 'tcx, 'tcx>) -> Self { let mut result = SymbolPathBuffer { result: String::with_capacity(64), temp_buf: String::with_capacity(16), + tcx, }; result.result.push_str("_ZN"); // _Z == Begin name-sequence, N == nested result } - fn from_interned(symbol: ty::SymbolName) -> Self { + fn from_interned(symbol: ty::SymbolName, tcx: TyCtxt<'a, 'tcx, 'tcx>) -> Self { let mut result = SymbolPathBuffer { result: String::with_capacity(64), temp_buf: String::with_capacity(16), + tcx, }; result.result.push_str(&symbol.as_str()); result @@ -377,7 +379,7 @@ impl SymbolPathBuffer { } } -impl ItemPathBuffer for SymbolPathBuffer { +impl ItemPathBuffer for SymbolPathBuffer<'a, 'tcx> { fn root_mode(&self) -> &RootMode { const ABSOLUTE: &RootMode = &RootMode::Absolute; ABSOLUTE @@ -385,7 +387,7 @@ impl ItemPathBuffer for SymbolPathBuffer { fn push(&mut self, text: &str) { self.temp_buf.clear(); - let need_underscore = sanitize(&mut self.temp_buf, text); + let need_underscore = sanitize(&mut self.temp_buf, text, self.tcx); let _ = write!( self.result, "{}", @@ -398,12 +400,24 @@ impl ItemPathBuffer for SymbolPathBuffer { } } +// Manual Debug implementation to omit non-Debug `tcx` field. +impl fmt::Debug for SymbolPathBuffer<'_, '_> { + fn fmt(&self, fmt: &mut fmt::Formatter) -> fmt::Result { + fmt.debug_struct("SymbolPathBuffer") + .field("result", &self.result) + .field("temp_buf", &self.temp_buf) + .finish() + } +} + // Name sanitation. LLVM will happily accept identifiers with weird names, but // gas doesn't! // gas accepts the following characters in symbols: a-z, A-Z, 0-9, ., _, $ +// NVPTX assembly has more strict naming rules than gas, so additionally, dots +// are replaced with '$' there. // // returns true if an underscore must be added at the start -pub fn sanitize(result: &mut String, s: &str) -> bool { +pub fn sanitize(result: &mut String, s: &str, tcx: TyCtxt<'a, 'tcx, 'tcx>) -> bool { for c in s.chars() { match c { // Escape these with $ sequences @@ -416,12 +430,25 @@ pub fn sanitize(result: &mut String, s: &str) -> bool { ')' => result.push_str("$RP$"), ',' => result.push_str("$C$"), - // '.' doesn't occur in types and functions, so reuse it - // for ':' and '-' - '-' | ':' => result.push('.'), + '-' | ':' => if tcx.has_strict_asm_symbol_naming() { + // NVPTX doesn't support these characters in symbol names. + result.push('$') + } + else { + // '.' doesn't occur in types and functions, so reuse it + // for ':' and '-' + result.push('.') + }, + + '.' => if tcx.has_strict_asm_symbol_naming() { + result.push('$') + } + else { + result.push('.') + }, // These are legal symbols - 'a'..='z' | 'A'..='Z' | '0'..='9' | '_' | '.' | '$' => result.push(c), + 'a'..='z' | 'A'..='Z' | '0'..='9' | '_' | '$' => result.push(c), _ => { result.push('$'); diff --git a/src/librustc_target/spec/mod.rs b/src/librustc_target/spec/mod.rs index e47da3cff95b6..aeecce49b0c67 100644 --- a/src/librustc_target/spec/mod.rs +++ b/src/librustc_target/spec/mod.rs @@ -75,6 +75,7 @@ pub enum LinkerFlavor { Ld, Msvc, Lld(LldFlavor), + PtxLinker, } #[derive(Clone, Copy, Debug, Eq, Ord, PartialEq, PartialOrd, Hash, @@ -143,6 +144,7 @@ flavor_mappings! { ((LinkerFlavor::Gcc), "gcc"), ((LinkerFlavor::Ld), "ld"), ((LinkerFlavor::Msvc), "msvc"), + ((LinkerFlavor::PtxLinker), "ptx-linker"), ((LinkerFlavor::Lld(LldFlavor::Wasm)), "wasm-ld"), ((LinkerFlavor::Lld(LldFlavor::Ld64)), "ld64.lld"), ((LinkerFlavor::Lld(LldFlavor::Ld)), "ld.lld"), @@ -455,6 +457,8 @@ supported_targets! { ("x86_64-fortanix-unknown-sgx", x86_64_fortanix_unknown_sgx), ("x86_64-unknown-uefi", x86_64_unknown_uefi), + + ("nvptx64-nvidia-cuda", nvptx64_nvidia_cuda), } /// Everything `rustc` knows about how to compile for a specific target. diff --git a/src/librustc_target/spec/nvptx64_nvidia_cuda.rs b/src/librustc_target/spec/nvptx64_nvidia_cuda.rs new file mode 100644 index 0000000000000..ed5d0f2450623 --- /dev/null +++ b/src/librustc_target/spec/nvptx64_nvidia_cuda.rs @@ -0,0 +1,73 @@ +use spec::{LinkerFlavor, Target, TargetOptions, TargetResult, PanicStrategy, MergeFunctions}; +use spec::abi::Abi; + +pub fn target() -> TargetResult { + Ok(Target { + arch: "nvptx64".to_string(), + data_layout: "e-i64:64-i128:128-v16:16-v32:32-n16:32:64".to_string(), + llvm_target: "nvptx64-nvidia-cuda".to_string(), + + target_os: "cuda".to_string(), + target_vendor: "nvidia".to_string(), + target_env: String::new(), + + linker_flavor: LinkerFlavor::PtxLinker, + + target_endian: "little".to_string(), + target_pointer_width: "64".to_string(), + target_c_int_width: "32".to_string(), + + options: TargetOptions { + // The linker can be installed from `crates.io`. + linker: Some("rust-ptx-linker".to_string()), + + // With `ptx-linker` approach, it can be later overriden via link flags. + cpu: "sm_20".to_string(), + + // TODO(denzp): create tests for the atomics. + max_atomic_width: Some(64), + + // Unwinding on CUDA is neither feasible nor useful. + panic_strategy: PanicStrategy::Abort, + + // Needed to use `dylib` and `bin` crate types and the linker. + dynamic_linking: true, + executables: true, + + // Avoid using dylib because it contain metadata not supported + // by LLVM NVPTX backend. + only_cdylib: true, + + // Let the `ptx-linker` to handle LLVM lowering into MC / assembly. + obj_is_bitcode: true, + + // Convinient and predicable naming scheme. + dll_prefix: "".to_string(), + dll_suffix: ".ptx".to_string(), + exe_suffix: ".ptx".to_string(), + + // Disable MergeFunctions LLVM optimisation pass because it can + // produce kernel functions that call other kernel functions. + // This behavior is not supported by PTX ISA. + merge_functions: MergeFunctions::Disabled, + + // TODO(denzp): enable compilation tests for the target and + // create the tests for this. + abi_blacklist: vec![ + Abi::Cdecl, + Abi::Stdcall, + Abi::Fastcall, + Abi::Vectorcall, + Abi::Thiscall, + Abi::Aapcs, + Abi::Win64, + Abi::SysV64, + Abi::Msp430Interrupt, + Abi::X86Interrupt, + Abi::AmdGpuKernel, + ], + + .. Default::default() + }, + }) +} diff --git a/src/test/run-make/nvptx-binary-crate/Makefile b/src/test/run-make/nvptx-binary-crate/Makefile new file mode 100644 index 0000000000000..4c22dae265c14 --- /dev/null +++ b/src/test/run-make/nvptx-binary-crate/Makefile @@ -0,0 +1,9 @@ +-include ../../run-make-fulldeps/tools.mk + +ifeq ($(TARGET),nvptx64-nvidia-cuda) +all: + $(RUSTC) main.rs -Clink-arg=--arch=sm_60 --crate-type="bin" -O --target $(TARGET) + FileCheck main.rs --input-file $(TMPDIR)/main.ptx +else +all: +endif diff --git a/src/test/run-make/nvptx-binary-crate/main.rs b/src/test/run-make/nvptx-binary-crate/main.rs new file mode 100644 index 0000000000000..826bc3a47bbd6 --- /dev/null +++ b/src/test/run-make/nvptx-binary-crate/main.rs @@ -0,0 +1,28 @@ +#![no_std] +#![no_main] +#![deny(warnings)] +#![feature(abi_ptx, core_intrinsics)] + +// Check the overriden CUDA arch. +// CHECK: .target sm_60 +// CHECK: .address_size 64 + +// Verify that no extra function declarations are present. +// CHECK-NOT: .func + +// CHECK-LABEL: .visible .entry top_kernel( +#[no_mangle] +pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) { + // CHECK: add.s32 %{{r[0-9]+}}, %{{r[0-9]+}}, 5; + *b = *a + 5; +} + +// Verify that no extra function definitions are there. +// CHECK-NOT: .func +// CHECK-NOT: .entry + +#[panic_handler] +unsafe fn breakpoint_panic_handler(_: &::core::panic::PanicInfo) -> ! { + core::intrinsics::breakpoint(); + core::hint::unreachable_unchecked(); +} diff --git a/src/test/run-make/nvptx-dylib-crate/Makefile b/src/test/run-make/nvptx-dylib-crate/Makefile new file mode 100644 index 0000000000000..7284e9d1a7c99 --- /dev/null +++ b/src/test/run-make/nvptx-dylib-crate/Makefile @@ -0,0 +1,10 @@ +-include ../../run-make-fulldeps/tools.mk + +ifeq ($(TARGET),nvptx64-nvidia-cuda) +all: + $(RUSTC) dep.rs --crate-type="rlib" --target $(TARGET) + $(RUSTC) kernel.rs --crate-type="cdylib" -O --target $(TARGET) + FileCheck kernel.rs --input-file $(TMPDIR)/kernel.ptx +else +all: +endif diff --git a/src/test/run-make/nvptx-dylib-crate/dep.rs b/src/test/run-make/nvptx-dylib-crate/dep.rs new file mode 100644 index 0000000000000..57f3ee87cdb9d --- /dev/null +++ b/src/test/run-make/nvptx-dylib-crate/dep.rs @@ -0,0 +1,14 @@ +#![no_std] +#![deny(warnings)] + +#[inline(never)] +#[no_mangle] +pub fn wrapping_external_fn(a: u32) -> u32 { + a.wrapping_mul(a) +} + +#[inline(never)] +#[no_mangle] +pub fn panicking_external_fn(a: u32) -> u32 { + a * a +} diff --git a/src/test/run-make/nvptx-dylib-crate/kernel.rs b/src/test/run-make/nvptx-dylib-crate/kernel.rs new file mode 100644 index 0000000000000..a889e23018d81 --- /dev/null +++ b/src/test/run-make/nvptx-dylib-crate/kernel.rs @@ -0,0 +1,67 @@ +#![no_std] +#![deny(warnings)] +#![feature(abi_ptx, core_intrinsics)] + +extern crate dep; + +// Verify the default CUDA arch. +// CHECK: .target sm_20 +// CHECK: .address_size 64 + +// Make sure declarations are there. +// CHECK: .func (.param .b32 func_retval0) wrapping_external_fn +// CHECK: .func (.param .b32 func_retval0) panicking_external_fn +// CHECK: .func [[PANIC_HANDLER:_ZN4core9panicking5panic[a-zA-Z0-9]+]] +// CHECK: .func [[PANIC_FMT:_ZN4core9panicking9panic_fmt[a-zA-Z0-9]+]] + +// CHECK-LABEL: .visible .entry top_kernel( +#[no_mangle] +pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) { + // CHECK: call.uni (retval0), + // CHECK-NEXT: wrapping_external_fn + // CHECK: ld.param.b32 %[[LHS:r[0-9]+]], [retval0+0]; + let lhs = dep::wrapping_external_fn(*a); + + // CHECK: call.uni (retval0), + // CHECK-NEXT: panicking_external_fn + // CHECK: ld.param.b32 %[[RHS:r[0-9]+]], [retval0+0]; + let rhs = dep::panicking_external_fn(*a); + + // CHECK: add.s32 %[[RES:r[0-9]+]], %[[RHS]], %[[LHS]]; + // CHECK: st.global.u32 [%{{rd[0-9]+}}], %[[RES]]; + *b = lhs + rhs; +} + +// Verify that external function bodies are available. +// CHECK-LABEL: .func (.param .b32 func_retval0) wrapping_external_fn +// CHECK: { +// CHECK: st.param.b32 [func_retval0+0], %{{r[0-9]+}}; +// CHECK: } + +// Also verify panic behavior. +// CHECK-LABEL: .func (.param .b32 func_retval0) panicking_external_fn +// CHECK: { +// CHECK: %{{p[0-9]+}} bra [[PANIC_LABEL:[a-zA-Z0-9_]+]]; +// CHECK: [[PANIC_LABEL]]: +// CHECK: call.uni +// CHECK: [[PANIC_HANDLER]] +// CHECK: } + +// Verify whether panic handler is present. +// CHECK: .func [[PANIC_HANDLER]]() +// CHECK: { +// CHECK: call.uni +// CHECK: [[PANIC_FMT]] +// CHECK: } + +// And finally, check the dummy panic formatter. +// CHECK: .func [[PANIC_FMT]]() +// CHECK: { +// CHECK: trap; +// CHECK: } + +#[panic_handler] +unsafe fn breakpoint_panic_handler(_: &::core::panic::PanicInfo) -> ! { + core::intrinsics::breakpoint(); + core::hint::unreachable_unchecked(); +} diff --git a/src/test/run-make/nvptx-emit-asm/Makefile b/src/test/run-make/nvptx-emit-asm/Makefile new file mode 100644 index 0000000000000..e03601878bdee --- /dev/null +++ b/src/test/run-make/nvptx-emit-asm/Makefile @@ -0,0 +1,9 @@ +-include ../../run-make-fulldeps/tools.mk + +ifeq ($(TARGET),nvptx64-nvidia-cuda) +all: + $(RUSTC) kernel.rs --crate-type="rlib" --emit asm,llvm-ir -O --target $(TARGET) + FileCheck kernel.rs --input-file $(TMPDIR)/kernel.s +else +all: +endif diff --git a/src/test/run-make/nvptx-emit-asm/kernel.rs b/src/test/run-make/nvptx-emit-asm/kernel.rs new file mode 100644 index 0000000000000..070a6efd2d547 --- /dev/null +++ b/src/test/run-make/nvptx-emit-asm/kernel.rs @@ -0,0 +1,41 @@ +#![no_std] +#![deny(warnings)] +#![feature(abi_ptx)] + +// Verify the default CUDA arch. +// CHECK: .target sm_20 +// CHECK: .address_size 64 + +// Verify function name doesn't contain unacceaptable characters. +// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN:_ZN[a-zA-Z0-9$_]+square[a-zA-Z0-9$_]+]] + +// CHECK-LABEL: .visible .entry top_kernel( +#[no_mangle] +pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) { + // CHECK: call.uni (retval0), + // CHECK-NEXT: [[IMPL_FN]] + *b = deep::private::MyStruct::new(*a).square(); +} + +pub mod deep { + pub mod private { + pub struct MyStruct(T); + + impl MyStruct { + pub fn new(a: u32) -> Self { + MyStruct(a) + } + + #[inline(never)] + pub fn square(&self) -> u32 { + self.0.wrapping_mul(self.0) + } + } + } +} + +// Verify that external function bodies are available. +// CHECK: .func (.param .b32 func_retval0) [[IMPL_FN]] +// CHECK: { +// CHECK: mul.lo.s32 %{{r[0-9]+}}, %{{r[0-9]+}}, %{{r[0-9]+}} +// CHECK: } From 97c8e82fe03c079368e7913c46f6b26fefa30150 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Thu, 24 Jan 2019 01:13:52 +0100 Subject: [PATCH 2/8] Enable CI for `nvptx64-nvidia-cuda` --- .travis.yml | 2 ++ src/ci/docker/dist-various-2/Dockerfile | 1 + src/ci/docker/nvptx-cuda/Dockerfile | 18 ++++++++++++++++++ 3 files changed, 21 insertions(+) create mode 100644 src/ci/docker/nvptx-cuda/Dockerfile diff --git a/.travis.yml b/.travis.yml index c4efa88460325..a8e1bfbbfa93a 100644 --- a/.travis.yml +++ b/.travis.yml @@ -186,6 +186,8 @@ matrix: if: branch = auto - env: IMAGE=mingw-check if: type = pull_request OR branch = auto + - env: IMAGE=nvptx-cuda + if: branch = auto - stage: publish toolstate if: branch = master AND type = push diff --git a/src/ci/docker/dist-various-2/Dockerfile b/src/ci/docker/dist-various-2/Dockerfile index 952c1ba2ccb76..66cbb43196aa2 100644 --- a/src/ci/docker/dist-various-2/Dockerfile +++ b/src/ci/docker/dist-various-2/Dockerfile @@ -70,6 +70,7 @@ ENV TARGETS=$TARGETS,x86_64-sun-solaris ENV TARGETS=$TARGETS,x86_64-unknown-linux-gnux32 ENV TARGETS=$TARGETS,x86_64-unknown-cloudabi ENV TARGETS=$TARGETS,x86_64-fortanix-unknown-sgx +ENV TARGETS=$TARGETS,nvptx64-nvidia-cuda ENV X86_FORTANIX_SGX_LIBS="/x86_64-fortanix-unknown-sgx/lib/" diff --git a/src/ci/docker/nvptx-cuda/Dockerfile b/src/ci/docker/nvptx-cuda/Dockerfile new file mode 100644 index 0000000000000..cdb1f565bd22e --- /dev/null +++ b/src/ci/docker/nvptx-cuda/Dockerfile @@ -0,0 +1,18 @@ +FROM ubuntu:18.04 + +RUN apt-get update +RUN apt-get install -y --no-install-recommends \ + g++ make file curl ca-certificates python git \ + cmake sudo gdb + +# TODO(denzp): setup `ptx-linker` CI for auttomatic binary releases. +RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha/rust-ptx-linker.linux64.tar.gz | \ + tar -xzvC /usr/bin + +COPY scripts/sccache.sh /scripts/ +RUN sh /scripts/sccache.sh + +ENV TARGETS=nvptx64-nvidia-cuda + +ENV SCRIPT python2.7 /checkout/x.py test --target $TARGETS \ + src/test/run-make From ceacde31aeac4732a970ba253a99502b5d9f5280 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Sun, 27 Jan 2019 23:52:22 +0100 Subject: [PATCH 3/8] NVPTX: by-default use target cpu "sm_30" --- src/ci/docker/nvptx-cuda/Dockerfile | 2 +- src/librustc_target/spec/nvptx64_nvidia_cuda.rs | 6 +++--- src/test/run-make/nvptx-emit-asm/kernel.rs | 2 +- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/src/ci/docker/nvptx-cuda/Dockerfile b/src/ci/docker/nvptx-cuda/Dockerfile index cdb1f565bd22e..b52865ced3e38 100644 --- a/src/ci/docker/nvptx-cuda/Dockerfile +++ b/src/ci/docker/nvptx-cuda/Dockerfile @@ -5,7 +5,7 @@ RUN apt-get install -y --no-install-recommends \ g++ make file curl ca-certificates python git \ cmake sudo gdb -# TODO(denzp): setup `ptx-linker` CI for auttomatic binary releases. +# FIXME: setup `ptx-linker` CI for automatic binary releases. RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha/rust-ptx-linker.linux64.tar.gz | \ tar -xzvC /usr/bin diff --git a/src/librustc_target/spec/nvptx64_nvidia_cuda.rs b/src/librustc_target/spec/nvptx64_nvidia_cuda.rs index ed5d0f2450623..e8512415e66a3 100644 --- a/src/librustc_target/spec/nvptx64_nvidia_cuda.rs +++ b/src/librustc_target/spec/nvptx64_nvidia_cuda.rs @@ -22,9 +22,9 @@ pub fn target() -> TargetResult { linker: Some("rust-ptx-linker".to_string()), // With `ptx-linker` approach, it can be later overriden via link flags. - cpu: "sm_20".to_string(), + cpu: "sm_30".to_string(), - // TODO(denzp): create tests for the atomics. + // FIXME: create tests for the atomics. max_atomic_width: Some(64), // Unwinding on CUDA is neither feasible nor useful. @@ -51,7 +51,7 @@ pub fn target() -> TargetResult { // This behavior is not supported by PTX ISA. merge_functions: MergeFunctions::Disabled, - // TODO(denzp): enable compilation tests for the target and + // FIXME: enable compilation tests for the target and // create the tests for this. abi_blacklist: vec![ Abi::Cdecl, diff --git a/src/test/run-make/nvptx-emit-asm/kernel.rs b/src/test/run-make/nvptx-emit-asm/kernel.rs index 070a6efd2d547..b71e18d910391 100644 --- a/src/test/run-make/nvptx-emit-asm/kernel.rs +++ b/src/test/run-make/nvptx-emit-asm/kernel.rs @@ -3,7 +3,7 @@ #![feature(abi_ptx)] // Verify the default CUDA arch. -// CHECK: .target sm_20 +// CHECK: .target sm_30 // CHECK: .address_size 64 // Verify function name doesn't contain unacceaptable characters. From 8d53c9247c23e36219102be68f887b0a39085a32 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Mon, 28 Jan 2019 01:16:59 +0100 Subject: [PATCH 4/8] SymbolPathBuffer shallow refactoring --- src/librustc_codegen_utils/symbol_names.rs | 156 ++++++++++----------- 1 file changed, 76 insertions(+), 80 deletions(-) diff --git a/src/librustc_codegen_utils/symbol_names.rs b/src/librustc_codegen_utils/symbol_names.rs index f2014f7421289..3238a0b10bfd6 100644 --- a/src/librustc_codegen_utils/symbol_names.rs +++ b/src/librustc_codegen_utils/symbol_names.rs @@ -103,7 +103,7 @@ use rustc_mir::monomorphize::Instance; use syntax_pos::symbol::Symbol; -use std::fmt::{self, Write}; +use std::fmt::Write; use std::mem::discriminant; pub fn provide(providers: &mut Providers) { @@ -339,28 +339,29 @@ fn compute_symbol_name<'a, 'tcx>(tcx: TyCtxt<'a, 'tcx, 'tcx>, instance: Instance // // To be able to work on all platforms and get *some* reasonable output, we // use C++ name-mangling. -struct SymbolPathBuffer<'a, 'tcx> { - tcx: TyCtxt<'a, 'tcx, 'tcx>, +#[derive(Debug)] +struct SymbolPathBuffer { result: String, temp_buf: String, + strict_naming: bool, } -impl SymbolPathBuffer<'a, 'tcx> { - fn new(tcx: TyCtxt<'a, 'tcx, 'tcx>) -> Self { +impl SymbolPathBuffer { + fn new(tcx: TyCtxt<'_, '_, '_>) -> Self { let mut result = SymbolPathBuffer { result: String::with_capacity(64), temp_buf: String::with_capacity(16), - tcx, + strict_naming: tcx.has_strict_asm_symbol_naming(), }; result.result.push_str("_ZN"); // _Z == Begin name-sequence, N == nested result } - fn from_interned(symbol: ty::SymbolName, tcx: TyCtxt<'a, 'tcx, 'tcx>) -> Self { + fn from_interned(symbol: ty::SymbolName, tcx: TyCtxt<'_, '_, '_>) -> Self { let mut result = SymbolPathBuffer { result: String::with_capacity(64), temp_buf: String::with_capacity(16), - tcx, + strict_naming: tcx.has_strict_asm_symbol_naming(), }; result.result.push_str(&symbol.as_str()); result @@ -377,93 +378,88 @@ impl SymbolPathBuffer<'a, 'tcx> { let _ = write!(self.result, "17h{:016x}E", hash); self.result } -} -impl ItemPathBuffer for SymbolPathBuffer<'a, 'tcx> { - fn root_mode(&self) -> &RootMode { - const ABSOLUTE: &RootMode = &RootMode::Absolute; - ABSOLUTE - } - - fn push(&mut self, text: &str) { + // Name sanitation. LLVM will happily accept identifiers with weird names, but + // gas doesn't! + // gas accepts the following characters in symbols: a-z, A-Z, 0-9, ., _, $ + // NVPTX assembly has more strict naming rules than gas, so additionally, dots + // are replaced with '$' there. + fn sanitize_and_append(&mut self, s: &str) { self.temp_buf.clear(); - let need_underscore = sanitize(&mut self.temp_buf, text, self.tcx); + + for c in s.chars() { + match c { + // Escape these with $ sequences + '@' => self.temp_buf.push_str("$SP$"), + '*' => self.temp_buf.push_str("$BP$"), + '&' => self.temp_buf.push_str("$RF$"), + '<' => self.temp_buf.push_str("$LT$"), + '>' => self.temp_buf.push_str("$GT$"), + '(' => self.temp_buf.push_str("$LP$"), + ')' => self.temp_buf.push_str("$RP$"), + ',' => self.temp_buf.push_str("$C$"), + + '-' | ':' => if self.strict_naming { + // NVPTX doesn't support these characters in symbol names. + self.temp_buf.push('$') + } + else { + // '.' doesn't occur in types and functions, so reuse it + // for ':' and '-' + self.temp_buf.push('.') + }, + + '.' => if self.strict_naming { + self.temp_buf.push('$') + } + else { + self.temp_buf.push('.') + }, + + // These are legal symbols + 'a'..='z' | 'A'..='Z' | '0'..='9' | '_' | '$' => self.temp_buf.push(c), + + _ => { + self.temp_buf.push('$'); + for c in c.escape_unicode().skip(1) { + match c { + '{' => {} + '}' => self.temp_buf.push('$'), + c => self.temp_buf.push(c), + } + } + } + } + } + + let need_underscore = { + // Underscore-qualify anything that didn't start as an ident. + !self.temp_buf.is_empty() + && self.temp_buf.as_bytes()[0] != '_' as u8 + && !(self.temp_buf.as_bytes()[0] as char).is_xid_start() + }; + let _ = write!( self.result, "{}", self.temp_buf.len() + (need_underscore as usize) ); + if need_underscore { self.result.push('_'); } + self.result.push_str(&self.temp_buf); } } -// Manual Debug implementation to omit non-Debug `tcx` field. -impl fmt::Debug for SymbolPathBuffer<'_, '_> { - fn fmt(&self, fmt: &mut fmt::Formatter) -> fmt::Result { - fmt.debug_struct("SymbolPathBuffer") - .field("result", &self.result) - .field("temp_buf", &self.temp_buf) - .finish() +impl ItemPathBuffer for SymbolPathBuffer { + fn root_mode(&self) -> &RootMode { + const ABSOLUTE: &RootMode = &RootMode::Absolute; + ABSOLUTE } -} - -// Name sanitation. LLVM will happily accept identifiers with weird names, but -// gas doesn't! -// gas accepts the following characters in symbols: a-z, A-Z, 0-9, ., _, $ -// NVPTX assembly has more strict naming rules than gas, so additionally, dots -// are replaced with '$' there. -// -// returns true if an underscore must be added at the start -pub fn sanitize(result: &mut String, s: &str, tcx: TyCtxt<'a, 'tcx, 'tcx>) -> bool { - for c in s.chars() { - match c { - // Escape these with $ sequences - '@' => result.push_str("$SP$"), - '*' => result.push_str("$BP$"), - '&' => result.push_str("$RF$"), - '<' => result.push_str("$LT$"), - '>' => result.push_str("$GT$"), - '(' => result.push_str("$LP$"), - ')' => result.push_str("$RP$"), - ',' => result.push_str("$C$"), - - '-' | ':' => if tcx.has_strict_asm_symbol_naming() { - // NVPTX doesn't support these characters in symbol names. - result.push('$') - } - else { - // '.' doesn't occur in types and functions, so reuse it - // for ':' and '-' - result.push('.') - }, - - '.' => if tcx.has_strict_asm_symbol_naming() { - result.push('$') - } - else { - result.push('.') - }, - - // These are legal symbols - 'a'..='z' | 'A'..='Z' | '0'..='9' | '_' | '$' => result.push(c), - _ => { - result.push('$'); - for c in c.escape_unicode().skip(1) { - match c { - '{' => {} - '}' => result.push('$'), - c => result.push(c), - } - } - } - } + fn push(&mut self, text: &str) { + self.sanitize_and_append(text); } - - // Underscore-qualify anything that didn't start as an ident. - !result.is_empty() && result.as_bytes()[0] != '_' as u8 - && !(result.as_bytes()[0] as char).is_xid_start() } From 6f86a70ea13cbe8b32c6e6ed76b8f57be4d70c68 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Mon, 28 Jan 2019 23:56:37 +0100 Subject: [PATCH 5/8] Adjust PTXLinker LTO logic and CLI --- src/ci/docker/nvptx-cuda/Dockerfile | 2 +- src/librustc_codegen_ssa/back/linker.rs | 17 ++++++++--------- 2 files changed, 9 insertions(+), 10 deletions(-) diff --git a/src/ci/docker/nvptx-cuda/Dockerfile b/src/ci/docker/nvptx-cuda/Dockerfile index b52865ced3e38..c7c3ca6bc5417 100644 --- a/src/ci/docker/nvptx-cuda/Dockerfile +++ b/src/ci/docker/nvptx-cuda/Dockerfile @@ -6,7 +6,7 @@ RUN apt-get install -y --no-install-recommends \ cmake sudo gdb # FIXME: setup `ptx-linker` CI for automatic binary releases. -RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha/rust-ptx-linker.linux64.tar.gz | \ +RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha.1/rust-ptx-linker.linux64.tar.gz | \ tar -xzvC /usr/bin COPY scripts/sccache.sh /scripts/ diff --git a/src/librustc_codegen_ssa/back/linker.rs b/src/librustc_codegen_ssa/back/linker.rs index 5e9aeed7107ac..55b02ebe6c4d5 100644 --- a/src/librustc_codegen_ssa/back/linker.rs +++ b/src/librustc_codegen_ssa/back/linker.rs @@ -13,7 +13,7 @@ use rustc::hir::def_id::{LOCAL_CRATE, CrateNum}; use rustc::middle::dependency_format::Linkage; use rustc::session::Session; use rustc::session::config::{self, CrateType, OptLevel, DebugInfo, - CrossLangLto}; + CrossLangLto, Lto}; use rustc::ty::TyCtxt; use rustc_target::spec::{LinkerFlavor, LldFlavor}; use serialize::{json, Encoder}; @@ -1118,14 +1118,13 @@ impl<'a> Linker for PtxLinker<'a> { } fn optimize(&mut self) { - self.cmd.arg(match self.sess.opts.optimize { - OptLevel::No => "-O0", - OptLevel::Less => "-O1", - OptLevel::Default => "-O2", - OptLevel::Aggressive => "-O3", - OptLevel::Size => "-Os", - OptLevel::SizeMin => "-Os" - }); + match self.sess.lto() { + Lto::Thin | Lto::Fat | Lto::ThinLocal => { + self.cmd.arg("-Olto"); + }, + + Lto::No => { }, + }; } fn output_filename(&mut self, path: &Path) { From 899d936dee555ef2ac67e68a5017ccb0d1e9a697 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Tue, 29 Jan 2019 19:06:42 +0100 Subject: [PATCH 6/8] Merge NVPTX and WASM test images into `test-various` --- .../Dockerfile | 19 ++++++++++++++----- src/test/run-make/nvptx-dylib-crate/kernel.rs | 12 ++---------- 2 files changed, 16 insertions(+), 15 deletions(-) rename src/ci/docker/{wasm32-unknown => test-various}/Dockerfile (63%) diff --git a/src/ci/docker/wasm32-unknown/Dockerfile b/src/ci/docker/test-various/Dockerfile similarity index 63% rename from src/ci/docker/wasm32-unknown/Dockerfile rename to src/ci/docker/test-various/Dockerfile index 161f0c0062fa0..a5ae94262c1ca 100644 --- a/src/ci/docker/wasm32-unknown/Dockerfile +++ b/src/ci/docker/test-various/Dockerfile @@ -13,14 +13,16 @@ RUN apt-get update && apt-get install -y --no-install-recommends \ gdb \ xz-utils +# FIXME: build the `ptx-linker` instead. +RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha.1/rust-ptx-linker.linux64.tar.gz | \ + tar -xzvC /usr/bin + RUN curl -sL https://nodejs.org/dist/v9.2.0/node-v9.2.0-linux-x64.tar.xz | \ - tar -xJ + tar -xJ COPY scripts/sccache.sh /scripts/ RUN sh /scripts/sccache.sh -ENV TARGETS=wasm32-unknown-unknown - ENV RUST_CONFIGURE_ARGS \ --set build.nodejs=/node-v9.2.0-linux-x64/bin/node \ --set rust.lld @@ -31,11 +33,18 @@ ENV RUST_CONFIGURE_ARGS \ # other contexts as well ENV NO_DEBUG_ASSERTIONS=1 -ENV SCRIPT python2.7 /checkout/x.py test --target $TARGETS \ +ENV WASM_TARGETS=wasm32-unknown-unknown +ENV WASM_SCRIPT python2.7 /checkout/x.py test --target $WASM_TARGETS \ src/test/run-make \ src/test/ui \ src/test/run-pass \ src/test/compile-fail \ src/test/mir-opt \ src/test/codegen-units \ - src/libcore \ + src/libcore + +ENV NVPTX_TARGETS=nvptx64-nvidia-cuda +ENV NVPTX_SCRIPT python2.7 /checkout/x.py test --target $NVPTX_TARGETS \ + src/test/run-make + +ENV SCRIPT $WASM_SCRIPT && $NVPTX_SCRIPT diff --git a/src/test/run-make/nvptx-dylib-crate/kernel.rs b/src/test/run-make/nvptx-dylib-crate/kernel.rs index a889e23018d81..5e65cca9140b7 100644 --- a/src/test/run-make/nvptx-dylib-crate/kernel.rs +++ b/src/test/run-make/nvptx-dylib-crate/kernel.rs @@ -12,7 +12,6 @@ extern crate dep; // CHECK: .func (.param .b32 func_retval0) wrapping_external_fn // CHECK: .func (.param .b32 func_retval0) panicking_external_fn // CHECK: .func [[PANIC_HANDLER:_ZN4core9panicking5panic[a-zA-Z0-9]+]] -// CHECK: .func [[PANIC_FMT:_ZN4core9panicking9panic_fmt[a-zA-Z0-9]+]] // CHECK-LABEL: .visible .entry top_kernel( #[no_mangle] @@ -47,15 +46,8 @@ pub unsafe extern "ptx-kernel" fn top_kernel(a: *const u32, b: *mut u32) { // CHECK: [[PANIC_HANDLER]] // CHECK: } -// Verify whether panic handler is present. -// CHECK: .func [[PANIC_HANDLER]]() -// CHECK: { -// CHECK: call.uni -// CHECK: [[PANIC_FMT]] -// CHECK: } - -// And finally, check the dummy panic formatter. -// CHECK: .func [[PANIC_FMT]]() +// Verify whether out dummy panic formatter has a correct body. +// CHECK: .func [[PANIC_FMT:_ZN4core9panicking9panic_fmt[a-zA-Z0-9]+]]() // CHECK: { // CHECK: trap; // CHECK: } From 3f624456e6e284f918d4e189e766e357911b49e2 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Tue, 29 Jan 2019 20:54:23 +0100 Subject: [PATCH 7/8] Provide PTXLinker with fallback to internal `target-cpu` --- src/ci/docker/test-various/Dockerfile | 2 +- src/librustc_codegen_ssa/back/linker.rs | 6 ++++++ src/test/run-make/nvptx-binary-crate/Makefile | 7 +++++-- src/test/run-make/nvptx-dylib-crate/kernel.rs | 2 +- 4 files changed, 13 insertions(+), 4 deletions(-) diff --git a/src/ci/docker/test-various/Dockerfile b/src/ci/docker/test-various/Dockerfile index a5ae94262c1ca..6c419e13c9f05 100644 --- a/src/ci/docker/test-various/Dockerfile +++ b/src/ci/docker/test-various/Dockerfile @@ -14,7 +14,7 @@ RUN apt-get update && apt-get install -y --no-install-recommends \ xz-utils # FIXME: build the `ptx-linker` instead. -RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha.1/rust-ptx-linker.linux64.tar.gz | \ +RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha.2/rust-ptx-linker.linux64.tar.gz | \ tar -xzvC /usr/bin RUN curl -sL https://nodejs.org/dist/v9.2.0/node-v9.2.0-linux-x64.tar.xz | \ diff --git a/src/librustc_codegen_ssa/back/linker.rs b/src/librustc_codegen_ssa/back/linker.rs index 55b02ebe6c4d5..249715a7b6e26 100644 --- a/src/librustc_codegen_ssa/back/linker.rs +++ b/src/librustc_codegen_ssa/back/linker.rs @@ -1132,6 +1132,12 @@ impl<'a> Linker for PtxLinker<'a> { } fn finalize(&mut self) -> Command { + // Provide the linker with fallback to internal `target-cpu`. + self.cmd.arg("--fallback-arch").arg(match self.sess.opts.cg.target_cpu { + Some(ref s) => s, + None => &self.sess.target.target.options.cpu + }); + ::std::mem::replace(&mut self.cmd, Command::new("")) } diff --git a/src/test/run-make/nvptx-binary-crate/Makefile b/src/test/run-make/nvptx-binary-crate/Makefile index 4c22dae265c14..2c211b5c78507 100644 --- a/src/test/run-make/nvptx-binary-crate/Makefile +++ b/src/test/run-make/nvptx-binary-crate/Makefile @@ -2,8 +2,11 @@ ifeq ($(TARGET),nvptx64-nvidia-cuda) all: - $(RUSTC) main.rs -Clink-arg=--arch=sm_60 --crate-type="bin" -O --target $(TARGET) - FileCheck main.rs --input-file $(TMPDIR)/main.ptx + $(RUSTC) main.rs --crate-type="bin" --target $(TARGET) -O -C link-arg=--arch=sm_60 -o $(TMPDIR)/main.link_arg.ptx + $(RUSTC) main.rs --crate-type="bin" --target $(TARGET) -O -C target-cpu=sm_60 -o $(TMPDIR)/main.target_cpu.ptx + + FileCheck main.rs --input-file $(TMPDIR)/main.link_arg.ptx + FileCheck main.rs --input-file $(TMPDIR)/main.target_cpu.ptx else all: endif diff --git a/src/test/run-make/nvptx-dylib-crate/kernel.rs b/src/test/run-make/nvptx-dylib-crate/kernel.rs index 5e65cca9140b7..63fd6b063dd8c 100644 --- a/src/test/run-make/nvptx-dylib-crate/kernel.rs +++ b/src/test/run-make/nvptx-dylib-crate/kernel.rs @@ -5,7 +5,7 @@ extern crate dep; // Verify the default CUDA arch. -// CHECK: .target sm_20 +// CHECK: .target sm_30 // CHECK: .address_size 64 // Make sure declarations are there. From 49931fda56dc6268ba3c104b64768f551cfc4636 Mon Sep 17 00:00:00 2001 From: Denys Zariaiev Date: Tue, 29 Jan 2019 21:00:07 +0100 Subject: [PATCH 8/8] Use correct CI test image for WASM and NVPTX --- .travis.yml | 4 +--- src/ci/docker/nvptx-cuda/Dockerfile | 18 ------------------ 2 files changed, 1 insertion(+), 21 deletions(-) delete mode 100644 src/ci/docker/nvptx-cuda/Dockerfile diff --git a/.travis.yml b/.travis.yml index a8e1bfbbfa93a..7985b6c0e191f 100644 --- a/.travis.yml +++ b/.travis.yml @@ -168,7 +168,7 @@ matrix: if: branch = auto - env: IMAGE=i686-gnu-nopt if: branch = auto - - env: IMAGE=wasm32-unknown + - env: IMAGE=test-various if: branch = auto - env: IMAGE=x86_64-gnu if: branch = auto @@ -186,8 +186,6 @@ matrix: if: branch = auto - env: IMAGE=mingw-check if: type = pull_request OR branch = auto - - env: IMAGE=nvptx-cuda - if: branch = auto - stage: publish toolstate if: branch = master AND type = push diff --git a/src/ci/docker/nvptx-cuda/Dockerfile b/src/ci/docker/nvptx-cuda/Dockerfile deleted file mode 100644 index c7c3ca6bc5417..0000000000000 --- a/src/ci/docker/nvptx-cuda/Dockerfile +++ /dev/null @@ -1,18 +0,0 @@ -FROM ubuntu:18.04 - -RUN apt-get update -RUN apt-get install -y --no-install-recommends \ - g++ make file curl ca-certificates python git \ - cmake sudo gdb - -# FIXME: setup `ptx-linker` CI for automatic binary releases. -RUN curl -sL https://github.com/denzp/rust-ptx-linker/releases/download/v0.9.0-alpha.1/rust-ptx-linker.linux64.tar.gz | \ - tar -xzvC /usr/bin - -COPY scripts/sccache.sh /scripts/ -RUN sh /scripts/sccache.sh - -ENV TARGETS=nvptx64-nvidia-cuda - -ENV SCRIPT python2.7 /checkout/x.py test --target $TARGETS \ - src/test/run-make