From c92d77e78ebd4022eac125deeaeb885c27e08acc Mon Sep 17 00:00:00 2001 From: Christian Legnitto Date: Thu, 10 Jul 2025 18:41:30 +0200 Subject: [PATCH] Update Rust toolchain to nightly-2025-06-23 --- crates/cuda_builder/src/lib.rs | 15 ++--- crates/cuda_std/src/thread.rs | 2 +- crates/cust/src/error.rs | 2 +- crates/cust/src/function.rs | 2 +- crates/cust/src/memory/array.rs | 2 +- .../cust/src/memory/device/device_buffer.rs | 2 +- crates/cust/src/memory/device/device_slice.rs | 12 +--- crates/cust_raw/build/callbacks.rs | 7 +-- crates/cust_raw/build/cuda_sdk.rs | 6 +- crates/cust_raw/build/main.rs | 12 +++- crates/cust_raw/src/lib.rs | 7 +++ crates/gpu_rand/src/xoroshiro/common.rs | 2 +- crates/nvvm/src/lib.rs | 9 +-- crates/ptx/src/lexer.rs | 15 +++-- crates/rustc_codegen_nvvm/build.rs | 44 +++++++------- crates/rustc_codegen_nvvm/src/abi.rs | 45 +++++++-------- crates/rustc_codegen_nvvm/src/asm.rs | 2 +- crates/rustc_codegen_nvvm/src/attributes.rs | 6 +- crates/rustc_codegen_nvvm/src/back.rs | 45 ++++++++------- crates/rustc_codegen_nvvm/src/builder.rs | 41 +++++++------ crates/rustc_codegen_nvvm/src/const_ty.rs | 14 ++--- crates/rustc_codegen_nvvm/src/consts.rs | 56 +++++++++++------- crates/rustc_codegen_nvvm/src/context.rs | 22 ++++--- .../src/debug_info/metadata.rs | 57 +++++++++---------- .../src/debug_info/metadata/enums.rs | 4 +- .../rustc_codegen_nvvm/src/debug_info/mod.rs | 6 +- crates/rustc_codegen_nvvm/src/intrinsic.rs | 47 +++++++-------- crates/rustc_codegen_nvvm/src/lib.rs | 20 +++---- crates/rustc_codegen_nvvm/src/link.rs | 33 ++++++----- crates/rustc_codegen_nvvm/src/lto.rs | 1 + crates/rustc_codegen_nvvm/src/mono_item.rs | 6 +- crates/rustc_codegen_nvvm/src/nvvm.rs | 6 +- crates/rustc_codegen_nvvm/src/override_fns.rs | 18 ++++-- crates/rustc_codegen_nvvm/src/ty.rs | 21 ++++--- examples/cuda/gemm/src/main.rs | 8 +-- examples/cuda/vecadd/src/main.rs | 5 +- rust-toolchain.toml | 2 +- xtask/src/extract_llfns.rs | 6 +- 38 files changed, 302 insertions(+), 308 deletions(-) diff --git a/crates/cuda_builder/src/lib.rs b/crates/cuda_builder/src/lib.rs index 1a056d1b..adab4a52 100644 --- a/crates/cuda_builder/src/lib.rs +++ b/crates/cuda_builder/src/lib.rs @@ -25,7 +25,7 @@ impl fmt::Display for CudaBuilderError { } CudaBuilderError::BuildFailed => f.write_str("Build failed"), CudaBuilderError::FailedToCopyPtxFile(err) => { - f.write_str(&format!("Failed to copy PTX file: {:?}", err)) + f.write_str(&format!("Failed to copy PTX file: {err:?}")) } } } @@ -369,19 +369,14 @@ fn find_rustc_codegen_nvvm() -> PathBuf { return path; } } - panic!("Could not find {} in library path", filename); + panic!("Could not find {filename} in library path"); } /// Joins strings together while ensuring none of the strings contain the separator. fn join_checking_for_separators(strings: Vec>, sep: &str) -> String { for s in &strings { let s = s.borrow(); - assert!( - !s.contains(sep), - "{:?} may not contain separator {:?}", - s, - sep - ); + assert!(!s.contains(sep), "{s:?} may not contain separator {sep:?}"); } strings.join(sep) } @@ -404,7 +399,7 @@ fn invoke_rustc(builder: &CudaBuilder) -> Result { EmitOption::LlvmIr => "llvm-ir", EmitOption::Bitcode => "llvm-bc", }; - rustflags.push(format!("--emit={}", string)); + rustflags.push(format!("--emit={string}")); } let mut llvm_args = vec![NvvmOption::Arch(builder.arch).to_string()]; @@ -533,7 +528,7 @@ fn get_last_artifact(out: &str) -> Option { Ok(line) => Some(line), Err(_) => { // Pass through invalid lines - println!("{}", line); + println!("{line}"); None } }) diff --git a/crates/cuda_std/src/thread.rs b/crates/cuda_std/src/thread.rs index 1f70bbb5..3a4d3432 100644 --- a/crates/cuda_std/src/thread.rs +++ b/crates/cuda_std/src/thread.rs @@ -66,7 +66,7 @@ macro_rules! inbounds { }}; ($func_name:ident, $lower_bound:expr, $upper_bound:expr) => {{ let val = unsafe { $func_name() }; - if val < $lower_bound || val > $upper_bound { + if !($lower_bound..=$upper_bound).contains(&val) { // SAFETY: this condition is declared unreachable by compute capability max bound // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities // we do this to potentially allow for better optimizations by LLVM diff --git a/crates/cust/src/error.rs b/crates/cust/src/error.rs index 5b8a881c..d5912014 100644 --- a/crates/cust/src/error.rs +++ b/crates/cust/src/error.rs @@ -105,7 +105,7 @@ impl fmt::Display for CudaError { .to_result() .map_err(|_| fmt::Error)?; let cstr = CStr::from_ptr(ptr); - write!(f, "{:?}", cstr) + write!(f, "{cstr:?}") } } // This shouldn't happen diff --git a/crates/cust/src/function.rs b/crates/cust/src/function.rs index 0bff08f4..a2dffae9 100644 --- a/crates/cust/src/function.rs +++ b/crates/cust/src/function.rs @@ -213,7 +213,7 @@ unsafe impl Send for Function<'_> {} unsafe impl Sync for Function<'_> {} impl Function<'_> { - pub(crate) fn new(inner: CUfunction, _module: &Module) -> Function { + pub(crate) fn new(inner: CUfunction, _module: &Module) -> Function<'_> { Function { inner, module: PhantomData, diff --git a/crates/cust/src/memory/array.rs b/crates/cust/src/memory/array.rs index 27daa49b..40b5ba06 100644 --- a/crates/cust/src/memory/array.rs +++ b/crates/cust/src/memory/array.rs @@ -139,7 +139,7 @@ impl ArrayFormat { // there are literally no docs on what nv12 is??? // it seems to be something with multiplanar arrays, needs some investigation CUarray_format_enum::CU_AD_FORMAT_NV12 => panic!("nv12 is not supported yet"), - _ => panic!("Unsupported array format: {:?}", raw), + _ => panic!("Unsupported array format: {raw:?}"), } } diff --git a/crates/cust/src/memory/device/device_buffer.rs b/crates/cust/src/memory/device/device_buffer.rs index 873a194c..706b058c 100644 --- a/crates/cust/src/memory/device/device_buffer.rs +++ b/crates/cust/src/memory/device/device_buffer.rs @@ -288,7 +288,7 @@ impl DeviceBuffer { #[cfg(feature = "bytemuck")] fn casting_went_wrong(src: &str, err: PodCastError) -> ! { - panic!("{}>{:?}", src, err); + panic!("{src}>{err:?}"); } #[cfg(feature = "bytemuck")] diff --git a/crates/cust/src/memory/device/device_slice.rs b/crates/cust/src/memory/device/device_slice.rs index 2062c06b..ff0db3dd 100644 --- a/crates/cust/src/memory/device/device_slice.rs +++ b/crates/cust/src/memory/device/device_slice.rs @@ -445,27 +445,21 @@ pub trait DeviceSliceIndex { #[cold] #[track_caller] fn slice_start_index_len_fail(index: usize, len: usize) -> ! { - panic!( - "range start index {} out of range for slice of length {}", - index, len - ); + panic!("range start index {index} out of range for slice of length {len}"); } #[inline(never)] #[cold] #[track_caller] fn slice_end_index_len_fail(index: usize, len: usize) -> ! { - panic!( - "range end index {} out of range for slice of length {}", - index, len - ); + panic!("range end index {index} out of range for slice of length {len}"); } #[inline(never)] #[cold] #[track_caller] fn slice_index_order_fail(index: usize, end: usize) -> ! { - panic!("slice index starts at {} but ends at {}", index, end); + panic!("slice index starts at {index} but ends at {end}"); } #[inline(never)] diff --git a/crates/cust_raw/build/callbacks.rs b/crates/cust_raw/build/callbacks.rs index 3ea22f56..d100220b 100644 --- a/crates/cust_raw/build/callbacks.rs +++ b/crates/cust_raw/build/callbacks.rs @@ -42,10 +42,7 @@ impl ParseCallbacks for BindgenCallbacks { match doxygen_bindgen::transform(&cleaned) { Ok(res) => Some(res), Err(err) => { - println!( - "cargo:warning=Problem processing doxygen comment: {}\n{}", - comment, err - ); + println!("cargo:warning=Problem processing doxygen comment: {comment}\n{err}"); None } } @@ -184,7 +181,7 @@ impl FunctionRenames { let expanded = match build.try_expand() { Ok(expanded) => expanded, - Err(e) => panic!("Failed to expand macros: {}", e), + Err(e) => panic!("Failed to expand macros: {e}"), }; let expanded = str::from_utf8(&expanded).unwrap(); diff --git a/crates/cust_raw/build/cuda_sdk.rs b/crates/cust_raw/build/cuda_sdk.rs index d3760027..e49c134f 100644 --- a/crates/cust_raw/build/cuda_sdk.rs +++ b/crates/cust_raw/build/cuda_sdk.rs @@ -164,7 +164,7 @@ impl CudaSdk { cuda_root: &path::Path, ) -> Result, Box> { let (target, triple) = Self::parse_target_triple()?; - assert!(triple.len() >= 3, "Invalid target triple: {:?}", triple); + assert!(triple.len() >= 3, "Invalid target triple: {triple:?}"); let search_dirs = match [triple[0].as_str(), triple[1].as_str(), triple[2].as_str()] { ["x86_64", "pc", "windows"] => { @@ -248,7 +248,7 @@ impl CudaSdk { .ok_or("Cannot find CUDA_VERSION from CUDA header file.")?; let version = version .parse::() - .map_err(|_| format!("Cannot parse CUDA_VERSION as u32: '{}'", version))?; + .map_err(|_| format!("Cannot parse CUDA_VERSION as u32: '{version}'"))?; Ok(version) } @@ -264,7 +264,7 @@ impl CudaSdk { .ok_or("Cannot find CUDART_VERSION from cuda_runtime header file.")?; let version = version .parse::() - .map_err(|_| format!("Cannot parse CUDART_VERSION as u32: '{}'", version))?; + .map_err(|_| format!("Cannot parse CUDART_VERSION as u32: '{version}'"))?; Ok(version) } diff --git a/crates/cust_raw/build/main.rs b/crates/cust_raw/build/main.rs index b70c10a1..d2adf6d3 100644 --- a/crates/cust_raw/build/main.rs +++ b/crates/cust_raw/build/main.rs @@ -59,12 +59,12 @@ fn main() { let metadata_nvvm_include = env::join_paths(sdk.nvvm_include_paths()) .map(|s| s.to_string_lossy().to_string()) .expect("Failed to build metadata for nvvm_include."); - println!("cargo::metadata=includes={}", metadata_cuda_include); - println!("cargo::metadata=nvvm_includes={}", metadata_nvvm_include); + println!("cargo::metadata=includes={metadata_cuda_include}"); + println!("cargo::metadata=nvvm_includes={metadata_nvvm_include}"); // Re-run build script conditions. println!("cargo::rerun-if-changed=build"); for e in sdk.related_cuda_envs() { - println!("cargo::rerun-if-env-changed={}", e); + println!("cargo::rerun-if-env-changed={e}"); } create_cuda_driver_bindings(&sdk, &outdir, &manifest_dir); @@ -138,6 +138,12 @@ fn create_cuda_driver_bindings( .allowlist_type("^cuda.*") .allowlist_var("^CU.*") .allowlist_function("^cu.*") + .no_partialeq("CUDA_HOST_NODE_PARAMS.*") + .no_partialeq("CUDA_KERNEL_NODE_PARAMS.*") + .no_hash("CUDA_HOST_NODE_PARAMS.*") + .no_hash("CUDA_KERNEL_NODE_PARAMS.*") + .no_copy("CUDA_HOST_NODE_PARAMS.*") + .no_copy("CUDA_KERNEL_NODE_PARAMS.*") .default_enum_style(bindgen::EnumVariation::Rust { non_exhaustive: false, }) diff --git a/crates/cust_raw/src/lib.rs b/crates/cust_raw/src/lib.rs index 62bd4b0c..f14e525f 100644 --- a/crates/cust_raw/src/lib.rs +++ b/crates/cust_raw/src/lib.rs @@ -1,16 +1,23 @@ #[cfg(feature = "driver")] +#[allow(clippy::missing_safety_doc)] pub mod driver_sys; #[cfg(feature = "runtime")] +#[allow(clippy::missing_safety_doc)] pub mod runtime_sys; #[cfg(feature = "cublas")] +#[allow(clippy::missing_safety_doc)] pub mod cublas_sys; #[cfg(feature = "cublaslt")] +#[allow(clippy::missing_safety_doc)] pub mod cublaslt_sys; #[cfg(feature = "cublasxt")] +#[allow(clippy::missing_safety_doc)] pub mod cublasxt_sys; #[cfg(feature = "nvptx-compiler")] +#[allow(clippy::missing_safety_doc)] pub mod nvptx_compiler_sys; #[cfg(feature = "nvvm")] +#[allow(clippy::missing_safety_doc)] pub mod nvvm_sys; diff --git a/crates/gpu_rand/src/xoroshiro/common.rs b/crates/gpu_rand/src/xoroshiro/common.rs index 82b321ab..8819ed0f 100644 --- a/crates/gpu_rand/src/xoroshiro/common.rs +++ b/crates/gpu_rand/src/xoroshiro/common.rs @@ -305,7 +305,7 @@ pub struct Seed512(pub [u8; 64]); impl Seed512 { /// Return an iterator over the seed. - pub fn iter(&self) -> core::slice::Iter { + pub fn iter(&self) -> core::slice::Iter<'_, u8> { self.0.iter() } } diff --git a/crates/nvvm/src/lib.rs b/crates/nvvm/src/lib.rs index 188fad6b..f8f4c9f8 100644 --- a/crates/nvvm/src/lib.rs +++ b/crates/nvvm/src/lib.rs @@ -166,7 +166,7 @@ impl Display for NvvmOption { Self::GenDebugInfo => "-g", Self::GenLineInfo => "-generate-line-info", Self::NoOpts => "-opt=0", - Self::Arch(arch) => return f.write_str(&format!("-arch={}", arch)), + Self::Arch(arch) => return f.write_str(&format!("-arch={arch}")), Self::Ftz => "-ftz=1", Self::FastSqrt => "-prec-sqrt=0", Self::FastDiv => "-prec-div=0", @@ -283,7 +283,7 @@ pub enum NvvmArch { impl Display for NvvmArch { fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { - let mut raw = format!("{:?}", self).to_ascii_lowercase(); + let mut raw = format!("{self:?}").to_ascii_lowercase(); raw.insert(7, '_'); f.write_str(&raw) } @@ -325,10 +325,7 @@ impl NvvmProgram { /// pub fn compile(&self, options: &[NvvmOption]) -> Result, NvvmError> { unsafe { - let options = options - .iter() - .map(|x| format!("{}\0", x)) - .collect::>(); + let options = options.iter().map(|x| format!("{x}\0")).collect::>(); let mut options_ptr = options .iter() .map(|x| x.as_ptr().cast()) diff --git a/crates/ptx/src/lexer.rs b/crates/ptx/src/lexer.rs index 08443e25..440e5b5f 100644 --- a/crates/ptx/src/lexer.rs +++ b/crates/ptx/src/lexer.rs @@ -259,7 +259,7 @@ impl<'src> Lexer<'src> { } c => { self.next(); - return Some(Err(format!("Unexpected token `{}`", c))); + return Some(Err(format!("Unexpected token `{c}`"))); } })) } @@ -293,7 +293,7 @@ impl<'src> Lexer<'src> { let val = string .as_str() .parse::() - .map_err(|_| format!("Failed to parse `{}` as f64 literal", string))?; + .map_err(|_| format!("Failed to parse `{string}` as f64 literal"))?; *self.values.last_mut().unwrap() = Some(TokenValue::Double(val)); return Ok(Token { @@ -315,7 +315,7 @@ impl<'src> Lexer<'src> { let val = string .as_str() .parse::() - .map_err(|_| format!("Failed to parse `{}` as f64 literal", string))?; + .map_err(|_| format!("Failed to parse `{string}` as f64 literal"))?; *self.values.last_mut().unwrap() = Some(TokenValue::Double(val)); return Ok(Token { @@ -338,7 +338,7 @@ impl<'src> Lexer<'src> { let val = string .as_str() .parse::() - .map_err(|_| format!("Failed to parse `{}` as f64 literal", string))?; + .map_err(|_| format!("Failed to parse `{string}` as f64 literal"))?; *self.values.last_mut().unwrap() = Some(TokenValue::Double(val)); return Ok(Token { @@ -369,7 +369,7 @@ impl<'src> Lexer<'src> { } let raw = u32::from_str_radix(numbers.as_str(), 16).map_err(|_| { - format!("Failed to parse `{}` as a 32 bit hex integer", numbers) + format!("Failed to parse `{numbers}` as a 32 bit hex integer") })?; *self.values.last_mut().unwrap() = Some(TokenValue::Float(f32::from_bits(raw))); @@ -390,7 +390,7 @@ impl<'src> Lexer<'src> { } let raw = u64::from_str_radix(numbers.as_str(), 16).map_err(|_| { - format!("Failed to parse `{}` as a 64 bit hex integer", numbers) + format!("Failed to parse `{numbers}` as a 64 bit hex integer") })?; *self.values.last_mut().unwrap() = @@ -557,8 +557,7 @@ impl<'src> Lexer<'src> { } Err(format!( - "Expected directive or reserved type, but found `.{}` instead", - ident + "Expected directive or reserved type, but found `.{ident}` instead" )) } } diff --git a/crates/rustc_codegen_nvvm/build.rs b/crates/rustc_codegen_nvvm/build.rs index a1c45b4e..ba453722 100644 --- a/crates/rustc_codegen_nvvm/build.rs +++ b/crates/rustc_codegen_nvvm/build.rs @@ -26,7 +26,7 @@ fn main() { } fn fail(s: &str) -> ! { - println!("\n\n{}\n\n", s); + println!("\n\n{s}\n\n"); std::process::exit(1); } @@ -34,10 +34,7 @@ fn fail(s: &str) -> ! { pub fn output(cmd: &mut Command) -> String { let output = match cmd.stderr(Stdio::inherit()).output() { Ok(status) => status, - Err(e) => fail(&format!( - "failed to execute command: {:?}\nerror: {}", - cmd, e - )), + Err(e) => fail(&format!("failed to execute command: {cmd:?}\nerror: {e}")), }; assert!( output.status.success(), @@ -56,11 +53,10 @@ fn target_to_llvm_prebuilt(target: &str) -> String { // NOTE(RDambrosio016): currently disabled because of weird issues with segfaults and building the C++ shim // "x86_64-unknown-linux-gnu" => "linux-x86_64", _ => panic!( - "Unsupported target with no matching prebuilt LLVM: `{}`, install LLVM and set LLVM_CONFIG", - target + "Unsupported target with no matching prebuilt LLVM: `{target}`, install LLVM and set LLVM_CONFIG" ), }; - format!("{}.tar.xz", base) + format!("{base}.tar.xz") } fn find_llvm_config(target: &str) -> PathBuf { @@ -79,8 +75,7 @@ fn find_llvm_config(target: &str) -> PathBuf { return PathBuf::from(path_to_try); } println!( - "cargo:warning=Prebuilt llvm-config version does not start with {}", - REQUIRED_MAJOR_LLVM_VERSION + "cargo:warning=Prebuilt llvm-config version does not start with {REQUIRED_MAJOR_LLVM_VERSION}" ); } else { println!("cargo:warning=Failed to run prebuilt llvm-config"); @@ -94,7 +89,7 @@ fn find_llvm_config(target: &str) -> PathBuf { .unwrap_or_else(|| PREBUILT_LLVM_URL.to_string()); let prebuilt_name = target_to_llvm_prebuilt(target); - url = format!("{}{}", url, prebuilt_name); + url = format!("{url}{prebuilt_name}"); let out = env::var("OUT_DIR").expect("OUT_DIR was not set"); let mut easy = Easy::new(); @@ -139,7 +134,7 @@ fn detect_llvm_link() -> (&'static str, &'static str) { } pub fn tracked_env_var_os + Display>(key: K) -> Option { - println!("cargo:rerun-if-env-changed={}", key); + println!("cargo:rerun-if-env-changed={key}"); env::var_os(key) } @@ -156,13 +151,12 @@ fn rustc_llvm_build() { for component in required_components { assert!( components.contains(component), - "require llvm component {} but wasn't found", - component + "require llvm component {component} but wasn't found" ); } for component in components.iter() { - println!("cargo:rustc-cfg=llvm_component=\"{}\"", component); + println!("cargo:rustc-cfg=llvm_component=\"{component}\""); } // Link in our own LLVM shims, compiled with the same flags as LLVM @@ -255,7 +249,7 @@ fn rustc_llvm_build() { } else { "dylib" }; - println!("cargo:rustc-link-lib={}={}", kind, name); + println!("cargo:rustc-link-lib={kind}={name}"); } // Link in the system libraries that LLVM depends on @@ -272,11 +266,11 @@ fn rustc_llvm_build() { cmd.arg(llvm_link_arg).arg("--ldflags"); for lib in output(&mut cmd).split_whitespace() { if let Some(stripped) = lib.strip_prefix("-LIBPATH:") { - println!("cargo:rustc-link-search=native={}", stripped); + println!("cargo:rustc-link-search=native={stripped}"); } else if let Some(stripped) = lib.strip_prefix("-l") { - println!("cargo:rustc-link-lib={}", stripped); + println!("cargo:rustc-link-lib={stripped}"); } else if let Some(stripped) = lib.strip_prefix("-L") { - println!("cargo:rustc-link-search=native={}", stripped); + println!("cargo:rustc-link-search=native={stripped}"); } } @@ -288,9 +282,9 @@ fn rustc_llvm_build() { if let Some(s) = llvm_linker_flags { for lib in s.into_string().unwrap().split_whitespace() { if let Some(stripped) = lib.strip_prefix("-l") { - println!("cargo:rustc-link-lib={}", stripped); + println!("cargo:rustc-link-lib={stripped}"); } else if let Some(stripped) = lib.strip_prefix("-L") { - println!("cargo:rustc-link-search=native={}", stripped); + println!("cargo:rustc-link-search=native={stripped}"); } } } @@ -330,14 +324,14 @@ fn rustc_llvm_build() { path.parent().unwrap().display() ); if target.contains("windows") { - println!("cargo:rustc-link-lib=static-nobundle={}", stdcppname); + println!("cargo:rustc-link-lib=static-nobundle={stdcppname}"); } else { - println!("cargo:rustc-link-lib=static={}", stdcppname); + println!("cargo:rustc-link-lib=static={stdcppname}"); } } else if cxxflags.contains("stdlib=libc++") { println!("cargo:rustc-link-lib=c++"); } else { - println!("cargo:rustc-link-lib={}", stdcppname); + println!("cargo:rustc-link-lib={stdcppname}"); } } @@ -365,6 +359,6 @@ fn link_llvm_system_libs(llvm_config: &Path, components: &[&str]) { continue; }; - println!("cargo:rustc-link-lib=dylib={}", name); + println!("cargo:rustc-link-lib=dylib={name}"); } } diff --git a/crates/rustc_codegen_nvvm/src/abi.rs b/crates/rustc_codegen_nvvm/src/abi.rs index 1b305010..eca12b21 100644 --- a/crates/rustc_codegen_nvvm/src/abi.rs +++ b/crates/rustc_codegen_nvvm/src/abi.rs @@ -2,6 +2,7 @@ use std::cmp; use libc::c_uint; use rustc_abi::BackendRepr::Scalar; +use rustc_abi::CanonAbi; use rustc_abi::Size; use rustc_abi::{HasDataLayout, Primitive, Reg, RegKind}; use rustc_codegen_ssa::mir::operand::OperandRef; @@ -13,7 +14,7 @@ use rustc_middle::ty::layout::LayoutOf; pub use rustc_middle::ty::layout::{WIDE_PTR_ADDR, WIDE_PTR_EXTRA}; use rustc_middle::ty::{Ty, TyCtxt, TyKind}; use rustc_target::callconv::{ - ArgAbi, ArgAttribute, ArgAttributes, ArgExtension, CastTarget, Conv, FnAbi, PassMode, + ArgAbi, ArgAttribute, ArgAttributes, ArgExtension, CastTarget, FnAbi, PassMode, }; use tracing::trace; @@ -28,7 +29,7 @@ pub(crate) fn readjust_fn_abi<'tcx>( fn_abi: &'tcx FnAbi<'tcx, Ty<'tcx>>, ) -> &'tcx FnAbi<'tcx, Ty<'tcx>> { // dont override anything in the rust abi for now - if fn_abi.conv == Conv::Rust { + if fn_abi.conv == CanonAbi::Rust { return fn_abi; } let readjust_arg_abi = |arg: &ArgAbi<'tcx, Ty<'tcx>>| { @@ -42,14 +43,14 @@ pub(crate) fn readjust_fn_abi<'tcx>( arg.mode = PassMode::Ignore; } - if let TyKind::Ref(_, ty, _) = arg.layout.ty.kind() { - if matches!(ty.kind(), TyKind::Slice(_)) { - let mut ptr_attrs = ArgAttributes::new(); - if let PassMode::Indirect { attrs, .. } = arg.mode { - ptr_attrs.regular = attrs.regular; - } - arg.mode = PassMode::Pair(ptr_attrs, ArgAttributes::new()); + if let TyKind::Ref(_, ty, _) = arg.layout.ty.kind() + && matches!(ty.kind(), TyKind::Slice(_)) + { + let mut ptr_attrs = ArgAttributes::new(); + if let PassMode::Indirect { attrs, .. } = arg.mode { + ptr_attrs.regular = attrs.regular; } + arg.mode = PassMode::Pair(ptr_attrs, ArgAttributes::new()); } if arg.layout.ty.is_array() && !matches!(arg.mode, PassMode::Direct { .. }) { @@ -260,9 +261,6 @@ impl<'ll, 'tcx> ArgAbiBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { ) { arg_abi.store(self, val, dst) } - fn arg_memory_ty(&self, arg_abi: &ArgAbi<'tcx, Ty<'tcx>>) -> &'ll Type { - arg_abi.memory_ty(self) - } } pub(crate) trait FnAbiLlvmExt<'ll, 'tcx> { @@ -493,11 +491,12 @@ impl<'ll, 'tcx> FnAbiLlvmExt<'ll, 'tcx> for FnAbi<'tcx, Ty<'tcx>> { // If the value is a boolean, the range is 0..2 and that ultimately // become 0..0 when the type becomes i1, which would be rejected // by the LLVM verifier. - if let Primitive::Int(..) = scalar.primitive() { - if !scalar.is_bool() && !scalar.is_always_valid(bx) { - trace!("apply_attrs_callsite -> range_metadata"); - bx.range_metadata(callsite, scalar.valid_range(bx)); - } + if let Primitive::Int(..) = scalar.primitive() + && !scalar.is_bool() + && !scalar.is_always_valid(bx) + { + trace!("apply_attrs_callsite -> range_metadata"); + bx.range_metadata(callsite, scalar.valid_range(bx)); } } for arg in self.args.iter() { @@ -541,7 +540,7 @@ impl<'ll, 'tcx> FnAbiLlvmExt<'ll, 'tcx> for FnAbi<'tcx, Ty<'tcx>> { } } -impl<'tcx> AbiBuilderMethods<'tcx> for Builder<'_, '_, 'tcx> { +impl<'tcx> AbiBuilderMethods for Builder<'_, '_, 'tcx> { fn get_param(&mut self, index: usize) -> Self::Value { let val = llvm::get_param(self.llfn(), index as c_uint); // trace!("Get param `{:?}`", val); @@ -551,11 +550,11 @@ impl<'tcx> AbiBuilderMethods<'tcx> for Builder<'_, '_, 'tcx> { // destructure so rustc doesnt complain in the call to transmute_llval let Self { cx, llbuilder } = self; let map = cx.remapped_integer_args.borrow(); - if let Some((_, key)) = map.get(llfnty) { - if let Some((_, new_ty)) = key.iter().find(|t| t.0 == index) { - trace!("Casting irregular param {:?} to {:?}", val, new_ty); - return transmute_llval(llbuilder, cx, val, new_ty); - } + if let Some((_, key)) = map.get(&llfnty) + && let Some((_, new_ty)) = key.iter().find(|t| t.0 == index) + { + trace!("Casting irregular param {:?} to {:?}", val, new_ty); + return transmute_llval(llbuilder, cx, val, new_ty); } val }; diff --git a/crates/rustc_codegen_nvvm/src/asm.rs b/crates/rustc_codegen_nvvm/src/asm.rs index c34e0e6d..b9d05fbc 100644 --- a/crates/rustc_codegen_nvvm/src/asm.rs +++ b/crates/rustc_codegen_nvvm/src/asm.rs @@ -224,7 +224,7 @@ impl<'tcx> AsmBuilderMethods<'tcx> for Builder<'_, '_, 'tcx> { impl<'tcx> AsmCodegenMethods<'tcx> for CodegenCx<'_, 'tcx> { fn codegen_global_asm( - &self, + &mut self, template: &[InlineAsmTemplatePiece], operands: &[GlobalAsmOperandRef], _options: InlineAsmOptions, diff --git a/crates/rustc_codegen_nvvm/src/attributes.rs b/crates/rustc_codegen_nvvm/src/attributes.rs index 85279724..ab98acf6 100644 --- a/crates/rustc_codegen_nvvm/src/attributes.rs +++ b/crates/rustc_codegen_nvvm/src/attributes.rs @@ -1,8 +1,8 @@ use crate::llvm::{self, AttributePlace::*, Value}; use rustc_ast::{LitKind, MetaItemInner, MetaItemLit}; -use rustc_attr_parsing::{InlineAttr, OptimizeAttr}; +use rustc_attr_data_structures::{InlineAttr, OptimizeAttr}; use rustc_hir::Attribute; -use rustc_middle::{bug, middle::codegen_fn_attrs::CodegenFnAttrFlags, ty}; +use rustc_middle::{middle::codegen_fn_attrs::CodegenFnAttrFlags, ty}; use rustc_session::{Session, config::OptLevel}; use rustc_span::{Symbol, sym}; @@ -16,7 +16,7 @@ fn inline(val: &'_ Value, inline: InlineAttr) { Always => llvm::Attribute::AlwaysInline.apply_llfn(Function, val), Never => llvm::Attribute::NoInline.apply_llfn(Function, val), None => {} - Force { .. } => bug!("Force inline should have been inlined away by now"), // TODO: Verify this + Force { .. } => llvm::Attribute::AlwaysInline.apply_llfn(Function, val), } } diff --git a/crates/rustc_codegen_nvvm/src/back.rs b/crates/rustc_codegen_nvvm/src/back.rs index 98d2861e..f523cf94 100644 --- a/crates/rustc_codegen_nvvm/src/back.rs +++ b/crates/rustc_codegen_nvvm/src/back.rs @@ -29,7 +29,7 @@ use crate::{LlvmMod, NvvmCodegenBackend, builder::Builder, context::CodegenCx, l pub fn llvm_err(handle: DiagCtxtHandle, msg: &str) -> FatalError { match llvm::last_error() { - Some(err) => handle.fatal(format!("{}: {}", msg, err)), + Some(err) => handle.fatal(format!("{msg}: {err}")), None => handle.fatal(msg.to_string()), } } @@ -121,7 +121,7 @@ pub fn target_machine_factory( false, ) }; - tm.ok_or_else(|| format!("Could not create LLVM TargetMachine for triple: {}", triple)) + tm.ok_or_else(|| format!("Could not create LLVM TargetMachine for triple: {triple}")) }) } @@ -146,7 +146,7 @@ pub extern "C" fn demangle_callback( Err(_) => return 0, }; - if write!(cursor, "{:#}", demangled).is_err() { + if write!(cursor, "{demangled:#}").is_err() { // Possible only if provided buffer is not big enough return 0; } @@ -176,11 +176,11 @@ pub(crate) unsafe fn codegen( let llmod = unsafe { module.module_llvm.llmod.as_ref().unwrap() }; let mod_name = module.name.clone(); - let module_name = Some(&mod_name[..]); + let module_name = &mod_name[..]; let out = cgcx .output_filenames - .temp_path(OutputType::Object, module_name); + .temp_path_for_cgu(OutputType::Object, module_name, None); // nvvm ir *is* llvm ir so emit_ir fits the expectation of llvm ir which is why we // implement this. this is copy and pasted straight from rustc_codegen_llvm @@ -189,9 +189,9 @@ pub(crate) unsafe fn codegen( let _timer = cgcx .prof .generic_activity_with_arg("NVVM_module_codegen_emit_ir", &module.name[..]); - let out = cgcx - .output_filenames - .temp_path(OutputType::LlvmAssembly, module_name); + let out = + cgcx.output_filenames + .temp_path_for_cgu(OutputType::LlvmAssembly, module_name, None); let out = out.to_str().unwrap(); let result = unsafe { @@ -199,7 +199,7 @@ pub(crate) unsafe fn codegen( }; result.into_result().map_err(|()| { - let msg = format!("failed to write NVVM IR to {}", out); + let msg = format!("failed to write NVVM IR to {out}"); llvm_err(dcx, &msg) })?; } @@ -229,6 +229,7 @@ pub(crate) unsafe fn codegen( bytecode: None, assembly: None, llvm_ir: None, + links_from_incr_cache: vec![], }) } @@ -254,7 +255,7 @@ pub fn compile_codegen_unit(tcx: TyCtxt<'_>, cgu_name: Symbol) -> (ModuleCodegen // Instantiate monomorphizations without filling out definitions yet... let llvm_module = LlvmMod::new(cgu_name.as_str()); { - let cx = CodegenCx::new(tcx, cgu, &llvm_module); + let mut cx = CodegenCx::new(tcx, cgu, &llvm_module); let mono_items = cx.codegen_unit.items_in_deterministic_order(cx.tcx); @@ -267,22 +268,27 @@ pub fn compile_codegen_unit(tcx: TyCtxt<'_>, cgu_name: Symbol) -> (ModuleCodegen }, ) in &mono_items { - mono_item.predefine::>(&cx, linkage, visibility); + mono_item.predefine::>( + &mut cx, + "mono_item", + linkage, + visibility, + ); } // ... and now that we have everything pre-defined, fill out those definitions. - for &(mono_item, _) in &mono_items { + for &(mono_item, mono_item_data) in &mono_items { if let MonoItem::Fn(func) = mono_item { - define_or_override_fn(func, &cx); + define_or_override_fn(func, &mut cx); } else { - mono_item.define::>(&cx); + mono_item.define::>(&mut cx, "mono_item", mono_item_data); } } // a main function for gpu kernels really makes no sense but // codegen it anyways. // sanitize attrs are not allowed in nvvm so do nothing further. - maybe_create_entry_wrapper::>(&cx); + maybe_create_entry_wrapper::>(&cx, cgu); // Run replace-all-uses-with for statics that need it for &(old_g, new_g) in cx.statics_to_rauw.borrow().iter() { @@ -333,13 +339,8 @@ pub(crate) unsafe fn optimize( let llmod = unsafe { &*module.module_llvm.llmod }; - let module_name = module.name.clone(); - let module_name = Some(&module_name[..]); - if config.emit_no_opt_bc { - let out = cgcx - .output_filenames - .temp_path_ext("no-opt.bc", module_name); + let out = cgcx.output_filenames.with_extension("no-opt.bc"); let out = path_to_c_string(&out); unsafe { llvm::LLVMWriteBitcodeToFile(llmod, out.as_ptr()) }; } @@ -389,7 +390,7 @@ pub(crate) unsafe fn optimize( for pass in &config.passes { if !addpass(pass) { - diag_handler.warn(format!("unknown pass `{}`, ignoring", pass)); + diag_handler.warn(format!("unknown pass `{pass}`, ignoring")); } } diff --git a/crates/rustc_codegen_nvvm/src/builder.rs b/crates/rustc_codegen_nvvm/src/builder.rs index ac2075f2..6bf428b3 100644 --- a/crates/rustc_codegen_nvvm/src/builder.rs +++ b/crates/rustc_codegen_nvvm/src/builder.rs @@ -6,7 +6,7 @@ use libc::{c_char, c_uint}; use rustc_abi as abi; use rustc_abi::{AddressSpace, Align, HasDataLayout, Size, TargetDataLayout, WrappingRange}; use rustc_codegen_ssa::MemFlags; -use rustc_codegen_ssa::common::{AtomicOrdering, IntPredicate, RealPredicate, TypeKind}; +use rustc_codegen_ssa::common::{IntPredicate, RealPredicate, TypeKind}; use rustc_codegen_ssa::mir::operand::{OperandRef, OperandValue}; use rustc_codegen_ssa::mir::place::PlaceRef; use rustc_codegen_ssa::traits::*; @@ -14,6 +14,7 @@ use rustc_data_structures::small_c_str::SmallCStr; use rustc_hir::def_id::DefId; use rustc_middle::bug; use rustc_middle::middle::codegen_fn_attrs::CodegenFnAttrs; +use rustc_middle::ty::AtomicOrdering; use rustc_middle::ty::layout::{ FnAbiError, FnAbiOfHelpers, FnAbiRequest, HasTypingEnv, LayoutError, LayoutOfHelpers, TyAndLayout, @@ -521,10 +522,10 @@ impl<'ll, 'tcx, 'a> BuilderMethods<'a, 'tcx> for Builder<'a, 'll, 'tcx> { bx.nonnull_metadata(load); } - if let Some(pointee) = layout.pointee_info_at(bx, offset) { - if pointee.safe.is_some() { - bx.align_metadata(load, pointee.align); - } + if let Some(pointee) = layout.pointee_info_at(bx, offset) + && pointee.safe.is_some() + { + bx.align_metadata(load, pointee.align); } } abi::Primitive::Float(_) => {} @@ -538,14 +539,12 @@ impl<'ll, 'tcx, 'a> BuilderMethods<'a, 'tcx> for Builder<'a, 'll, 'tcx> { let mut const_llval = None; let llty = place.layout.llvm_type(self); unsafe { - if let Some(global) = llvm::LLVMIsAGlobalVariable(place.val.llval) { - if llvm::LLVMIsGlobalConstant(global) == llvm::True { - if let Some(init) = llvm::LLVMGetInitializer(global) { - if self.val_ty(init) == llty { - const_llval = Some(init); - } - } - } + if let Some(global) = llvm::LLVMIsAGlobalVariable(place.val.llval) + && llvm::LLVMIsGlobalConstant(global) == llvm::True + && let Some(init) = llvm::LLVMGetInitializer(global) + && self.val_ty(init) == llty + { + const_llval = Some(init); } } @@ -704,7 +703,7 @@ impl<'ll, 'tcx, 'a> BuilderMethods<'a, 'tcx> for Builder<'a, 'll, 'tcx> { &mut self, _val: &'ll Value, ptr: &'ll Value, - _order: rustc_codegen_ssa::common::AtomicOrdering, + _order: AtomicOrdering, _size: Size, ) { // see comment in atomic_load @@ -1042,8 +1041,8 @@ impl<'ll, 'tcx, 'a> BuilderMethods<'a, 'tcx> for Builder<'a, 'll, 'tcx> { _dst: &'ll Value, _cmp: &'ll Value, _src: &'ll Value, - _order: rustc_codegen_ssa::common::AtomicOrdering, - _failure_order: rustc_codegen_ssa::common::AtomicOrdering, + _order: AtomicOrdering, + _failure_order: AtomicOrdering, _weak: bool, ) -> (&'ll Value, &'ll Value) { // allowed but only for some things and with restrictions @@ -1055,7 +1054,7 @@ impl<'ll, 'tcx, 'a> BuilderMethods<'a, 'tcx> for Builder<'a, 'll, 'tcx> { _op: rustc_codegen_ssa::common::AtomicRmwBinOp, _dst: &'ll Value, _src: &'ll Value, - _order: rustc_codegen_ssa::common::AtomicOrdering, + _order: AtomicOrdering, ) -> &'ll Value { // see cmpxchg comment self.fatal("atomic rmw is not supported") @@ -1063,7 +1062,7 @@ impl<'ll, 'tcx, 'a> BuilderMethods<'a, 'tcx> for Builder<'a, 'll, 'tcx> { fn atomic_fence( &mut self, - _order: rustc_codegen_ssa::common::AtomicOrdering, + _order: AtomicOrdering, _scope: rustc_codegen_ssa::common::SynchronizationScope, ) { self.fatal("atomic fence is not supported, use cuda_std intrinsics instead") @@ -1120,7 +1119,7 @@ impl<'ll, 'tcx, 'a> BuilderMethods<'a, 'tcx> for Builder<'a, 'll, 'tcx> { while self.cx.type_kind(fn_ty) == TypeKind::Pointer { fn_ty = self.cx.element_type(fn_ty); } - if let Some((Some(ret_ty), _)) = map.get(fn_ty) { + if let Some((Some(ret_ty), _)) = map.get(&fn_ty) { self.cx.last_call_llfn.set(Some(call)); call = transmute_llval(self.llbuilder, self.cx, call, ret_ty); } @@ -1211,9 +1210,7 @@ impl<'a, 'll, 'tcx> Builder<'a, 'll, 'tcx> { ) -> Cow<'b, [&'ll Value]> { assert!( self.cx.type_kind(fn_ty) == TypeKind::Function, - "builder::{} not passed a function, but {:?}", - typ, - fn_ty + "builder::{typ} not passed a function, but {fn_ty:?}" ); let param_tys = self.cx.func_params_types(fn_ty); diff --git a/crates/rustc_codegen_nvvm/src/const_ty.rs b/crates/rustc_codegen_nvvm/src/const_ty.rs index a03de2a2..fe9a4451 100644 --- a/crates/rustc_codegen_nvvm/src/const_ty.rs +++ b/crates/rustc_codegen_nvvm/src/const_ty.rs @@ -14,7 +14,7 @@ use rustc_middle::mir::interpret::{ConstAllocation, GlobalAlloc, Scalar}; use rustc_middle::ty::layout::LayoutOf; use tracing::trace; -impl<'ll, 'tcx> ConstCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { +impl<'ll, 'tcx> ConstCodegenMethods for CodegenCx<'ll, 'tcx> { fn const_data_from_alloc(&self, alloc: ConstAllocation) -> &'ll Value { const_alloc_to_llvm(self, alloc, /*static*/ false) } @@ -82,8 +82,7 @@ impl<'ll, 'tcx> ConstCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { let val = *self .const_cstr_cache .borrow_mut() - .raw_entry_mut() - .from_key(s) + .entry(s.to_string()) .or_insert_with(|| { let sc = self.const_bytes(s.as_bytes()); let sym = self.generate_local_symbol_name("str"); @@ -97,9 +96,8 @@ impl<'ll, 'tcx> ConstCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { llvm::LLVMSetGlobalConstant(g, True); llvm::LLVMRustSetLinkage(g, llvm::Linkage::InternalLinkage); } - (s.to_owned(), g) - }) - .1; + g + }); let len = s.len(); let ty = self.type_ptr_to(self.layout_of(self.tcx.types.str_).llvm_type(self)); let cs = unsafe { llvm::LLVMConstPointerCast(val, ty) }; @@ -247,10 +245,6 @@ impl<'ll, 'tcx> ConstCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { val } - fn is_undef(&self, v: Self::Value) -> bool { - unsafe { llvm::LLVMIsUndef(v) == True } - } - fn const_poison(&self, t: Self::Type) -> Self::Value { // FIXME: Use LLVMGetPoision when possible. self.const_undef(t) diff --git a/crates/rustc_codegen_nvvm/src/consts.rs b/crates/rustc_codegen_nvvm/src/consts.rs index a75cae86..a6b6ca4d 100644 --- a/crates/rustc_codegen_nvvm/src/consts.rs +++ b/crates/rustc_codegen_nvvm/src/consts.rs @@ -186,7 +186,7 @@ fn check_and_apply_linkage<'ll, 'tcx>( _ => cx .sess() .dcx() - .fatal(format!("Unsupported linkage kind: {:?}", linkage)), + .fatal(format!("Unsupported linkage kind: {linkage:?}")), } // If this is a static with a linkage specified, then we need to handle @@ -276,8 +276,7 @@ impl<'ll> CodegenCx<'ll, '_> { assert!( !defined_in_current_codegen_unit, "consts::get_static() should always hit the cache for \ - statics defined in the same CGU, but did not for `{:?}`", - def_id + statics defined in the same CGU, but did not for `{def_id:?}`" ); let ty = instance.ty(self.tcx, self.typing_env()); @@ -286,10 +285,10 @@ impl<'ll> CodegenCx<'ll, '_> { let g = if def_id.is_local() && !self.tcx.is_foreign_item(def_id) { let llty = self.layout_of(ty).llvm_type(self); - if let Some(g) = self.get_declared_value(sym) { - if self.val_ty(g) != self.type_ptr_to(llty) { - span_bug!(self.tcx.def_span(def_id), "Conflicting types for static"); - } + if let Some(g) = self.get_declared_value(sym) + && self.val_ty(g) != self.type_ptr_to(llty) + { + span_bug!(self.tcx.def_span(def_id), "Conflicting types for static"); } let addrspace = self.static_addrspace(instance); @@ -336,7 +335,7 @@ impl<'ll> StaticCodegenMethods for CodegenCx<'ll, '_> { gv } - fn codegen_static(&self, def_id: DefId) { + fn codegen_static(&mut self, def_id: DefId) { unsafe { assert!( llvm::LLVMGetInitializer( @@ -422,23 +421,36 @@ impl<'ll> StaticCodegenMethods for CodegenCx<'ll, '_> { self.unsupported("thread locals"); } - if attrs.flags.contains(CodegenFnAttrFlags::USED) { + if attrs.flags.contains(CodegenFnAttrFlags::USED_COMPILER) { + // `USED` and `USED_LINKER` can't be used together. + assert!(!attrs.flags.contains(CodegenFnAttrFlags::USED_LINKER)); + + // The semantics of #[used] in Rust only require the symbol to make it into the + // object file. It is explicitly allowed for the linker to strip the symbol if it + // is dead, which means we are allowed to use `llvm.compiler.used` instead of + // `llvm.used` here. + // + // Additionally, https://reviews.llvm.org/D97448 in LLVM 13 started emitting unique + // sections with SHF_GNU_RETAIN flag for llvm.used symbols, which may trigger bugs + // in the handling of `.init_array` (the static constructor list) in versions of + // the gold linker (prior to the one released with binutils 2.36). + // + // That said, we only ever emit these when `#[used(compiler)]` is explicitly + // requested. This is to avoid similar breakage on other targets, in particular + // MachO targets have *their* static constructor lists broken if `llvm.compiler.used` + // is emitted rather than `llvm.used`. However, that check happens when assigning + // the `CodegenFnAttrFlags` in the `codegen_fn_attrs` query, so we don't need to + // take care of it here. + self.add_compiler_used_global(g); + } + if attrs.flags.contains(CodegenFnAttrFlags::USED_LINKER) { + // `USED` and `USED_LINKER` can't be used together. + assert!(!attrs.flags.contains(CodegenFnAttrFlags::USED_COMPILER)); + self.add_used_global(g); } + trace!("Codegen static `{:?}`", g); } } - - /// Add a global value to a list to be stored in the `llvm.used` variable, an array of i8*. - fn add_used_global(&self, global: &'ll Value) { - let cast = unsafe { llvm::LLVMConstPointerCast(global, self.type_i8p()) }; - self.used_statics.borrow_mut().push(cast); - } - - /// Add a global value to a list to be stored in the `llvm.compiler.used` variable, - /// an array of i8*. - fn add_compiler_used_global(&self, global: &'ll Value) { - let cast = unsafe { llvm::LLVMConstPointerCast(global, self.type_i8p()) }; - self.compiler_used_statics.borrow_mut().push(cast); - } } diff --git a/crates/rustc_codegen_nvvm/src/context.rs b/crates/rustc_codegen_nvvm/src/context.rs index a185a101..6392705e 100644 --- a/crates/rustc_codegen_nvvm/src/context.rs +++ b/crates/rustc_codegen_nvvm/src/context.rs @@ -184,7 +184,7 @@ impl<'ll, 'tcx> CodegenCx<'ll, 'tcx> { // im lazy i know pub(crate) fn unsupported(&self, thing: &str) -> ! { - self.fatal(format!("{} is unsupported", thing)) + self.fatal(format!("{thing} is unsupported")) } pub(crate) fn create_used_variable_impl(&self, name: &'static CStr, values: &[&'ll Value]) { @@ -232,10 +232,6 @@ impl<'ll, 'tcx> MiscCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { self.tcx.sess } - fn codegen_unit(&self) -> &'tcx CodegenUnit<'tcx> { - self.codegen_unit - } - fn declare_c_main( &self, _fn_type: as rustc_codegen_ssa::traits::BackendTypes>::Type, @@ -286,8 +282,7 @@ impl<'ll, 'tcx> CodegenCx<'ll, 'tcx> { let layout = self.layout_of(ty); if layout.size.bytes() > CONSTANT_MEMORY_SIZE_LIMIT_BYTES { self.tcx.sess.dcx().warn(format!( - "static `{}` exceeds the constant memory limit; placing in global memory (performance may be reduced)", - instance + "static `{instance}` exceeds the constant memory limit; placing in global memory (performance may be reduced)" )); // Place instance in global memory if it is too big for constant memory. AddressSpace(1) @@ -541,6 +536,19 @@ impl<'ll, 'tcx> CodegenCx<'ll, 'tcx> { llfn } + + /// Add a global value to a list to be stored in the `llvm.used` variable, an array of i8*. + pub fn add_used_global(&self, global: &'ll Value) { + let cast = unsafe { llvm::LLVMConstPointerCast(global, self.type_i8p()) }; + self.used_statics.borrow_mut().push(cast); + } + + /// Add a global value to a list to be stored in the `llvm.compiler.used` variable, + /// an array of i8*. + pub fn add_compiler_used_global(&self, global: &'ll Value) { + let cast = unsafe { llvm::LLVMConstPointerCast(global, self.type_i8p()) }; + self.compiler_used_statics.borrow_mut().push(cast); + } } #[derive(Default, Clone)] diff --git a/crates/rustc_codegen_nvvm/src/debug_info/metadata.rs b/crates/rustc_codegen_nvvm/src/debug_info/metadata.rs index 3b9cf93a..dbab74e1 100644 --- a/crates/rustc_codegen_nvvm/src/debug_info/metadata.rs +++ b/crates/rustc_codegen_nvvm/src/debug_info/metadata.rs @@ -800,7 +800,7 @@ pub(crate) fn build_compile_unit_di_node<'ll, 'tcx>( // leave the clang LLVM in there just in case, although it shouldnt be needed because // gpu stuff is different - let producer = format!("clang LLVM ({})", rustc_producer); + let producer = format!("clang LLVM ({rustc_producer})"); let name_in_debuginfo = name_in_debuginfo.to_string_lossy(); let work_dir = tcx @@ -814,7 +814,8 @@ pub(crate) fn build_compile_unit_di_node<'ll, 'tcx>( && let Some(f) = output_filenames.split_dwarf_path( tcx.sess.split_debuginfo(), tcx.sess.opts.unstable_opts.split_dwarf_kind, - Some(codegen_unit_name), + codegen_unit_name, + None, ) { // We get a path relative to the working directory from split_dwarf_path Some(tcx.sess.source_map().path_mapping().to_real_filename(f)) @@ -844,7 +845,7 @@ pub(crate) fn build_compile_unit_di_node<'ll, 'tcx>( 0, ); - let unit_metadata = llvm::LLVMRustDIBuilderCreateCompileUnit( + llvm::LLVMRustDIBuilderCreateCompileUnit( debug_context.builder, dwarf_const::DW_LANG_Rust, compile_unit_file, @@ -861,9 +862,7 @@ pub(crate) fn build_compile_unit_di_node<'ll, 'tcx>( kind, 0, tcx.sess.opts.unstable_opts.split_dwarf_inlining, - ); - - unit_metadata + ) } } @@ -1199,31 +1198,31 @@ fn build_generic_type_param_di_nodes<'ll, 'tcx>( cx: &CodegenCx<'ll, 'tcx>, ty: Ty<'tcx>, ) -> SmallVec<&'ll DIType> { - if let ty::Adt(def, args) = *ty.kind() { - if args.types().next().is_some() { - let generics = cx.tcx.generics_of(def.did()); - let names = get_parameter_names(cx, generics); - let template_params: SmallVec<_> = iter::zip(args, names) - .filter_map(|(kind, name)| { - kind.as_type().map(|ty| { - let actual_type = cx.tcx.normalize_erasing_regions(cx.typing_env(), ty); - let actual_type_di_node = type_di_node(cx, actual_type); - let name = name.as_str(); - unsafe { - llvm::LLVMRustDIBuilderCreateTemplateTypeParameter( - DIB(cx), - None, - name.as_c_char_ptr(), - name.len(), - actual_type_di_node, - ) - } - }) + if let ty::Adt(def, args) = *ty.kind() + && args.types().next().is_some() + { + let generics = cx.tcx.generics_of(def.did()); + let names = get_parameter_names(cx, generics); + let template_params: SmallVec<_> = iter::zip(args, names) + .filter_map(|(kind, name)| { + kind.as_type().map(|ty| { + let actual_type = cx.tcx.normalize_erasing_regions(cx.typing_env(), ty); + let actual_type_di_node = type_di_node(cx, actual_type); + let name = name.as_str(); + unsafe { + llvm::LLVMRustDIBuilderCreateTemplateTypeParameter( + DIB(cx), + None, + name.as_c_char_ptr(), + name.len(), + actual_type_di_node, + ) + } }) - .collect(); + }) + .collect(); - return template_params; - } + return template_params; } return smallvec![]; diff --git a/crates/rustc_codegen_nvvm/src/debug_info/metadata/enums.rs b/crates/rustc_codegen_nvvm/src/debug_info/metadata/enums.rs index b4ba79a0..409094b5 100644 --- a/crates/rustc_codegen_nvvm/src/debug_info/metadata/enums.rs +++ b/crates/rustc_codegen_nvvm/src/debug_info/metadata/enums.rs @@ -189,7 +189,7 @@ pub(super) fn build_coroutine_di_node<'ll, 'tcx>( |cx, coroutine_type_di_node| { let coroutine_layout = cx .tcx - .coroutine_layout(coroutine_def_id, coroutine_args.as_coroutine().kind_ty()) + .coroutine_layout(coroutine_def_id, coroutine_args) .unwrap(); let Variants::Multiple { @@ -393,7 +393,7 @@ fn build_discr_member_di_node<'ll, 'tcx>( align.bits() as u32, enum_or_coroutine_type_and_layout .fields - .offset(tag_field) + .offset(tag_field.into()) .bits(), DIFlags::FlagArtificial, type_di_node(cx, tag_base_type), diff --git a/crates/rustc_codegen_nvvm/src/debug_info/mod.rs b/crates/rustc_codegen_nvvm/src/debug_info/mod.rs index 73ad6886..af8202fb 100644 --- a/crates/rustc_codegen_nvvm/src/debug_info/mod.rs +++ b/crates/rustc_codegen_nvvm/src/debug_info/mod.rs @@ -190,10 +190,6 @@ impl<'ll> DebugInfoBuilderMethods for Builder<'_, 'll, '_> { llvm::LLVMSetCurrentDebugLocation(self.llbuilder, None); } } - - fn get_dbg_loc(&self) -> Option { - None // TODO: implement this - } } /// A source code location used to generate debug information. @@ -366,7 +362,7 @@ impl<'ll, 'tcx> DebugInfoCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { let names = get_parameter_names(cx, generics); iter::zip(args, names) .filter_map(|(kind, name)| { - if let GenericArgKind::Type(ty) = kind.unpack() { + if let GenericArgKind::Type(ty) = kind.kind() { let actual_type = cx.tcx.normalize_erasing_regions(cx.typing_env(), ty); let actual_type_metadata = type_di_node(cx, actual_type); let name = name.as_str(); diff --git a/crates/rustc_codegen_nvvm/src/intrinsic.rs b/crates/rustc_codegen_nvvm/src/intrinsic.rs index 18811266..b6a7c28c 100644 --- a/crates/rustc_codegen_nvvm/src/intrinsic.rs +++ b/crates/rustc_codegen_nvvm/src/intrinsic.rs @@ -1,5 +1,5 @@ use rustc_abi as abi; -use rustc_abi::{self, Float, HasDataLayout, Primitive}; +use rustc_abi::{self, BackendRepr, Float, HasDataLayout, Primitive}; use rustc_codegen_ssa::errors::InvalidMonomorphization; use rustc_codegen_ssa::mir::operand::OperandValue; use rustc_codegen_ssa::mir::place::PlaceValue; @@ -8,18 +8,18 @@ use rustc_codegen_ssa::traits::{ BaseTypeCodegenMethods, BuilderMethods, ConstCodegenMethods, IntrinsicCallBuilderMethods, OverflowOp, }; -use rustc_middle::ty::layout::{HasTypingEnv, LayoutOf}; +use rustc_middle::ty::layout::{FnAbiOf, HasTypingEnv, LayoutOf}; use rustc_middle::ty::{self, Ty}; use rustc_middle::{bug, span_bug}; use rustc_span::symbol::kw; use rustc_span::{Span, Symbol, sym}; -use rustc_target::callconv::{FnAbi, PassMode}; +use rustc_target::callconv::PassMode; use tracing::trace; use crate::abi::LlvmType; use crate::builder::Builder; use crate::context::CodegenCx; -use crate::llvm::{self, Metadata, Type, Value}; +use crate::llvm::{self, Type, Value}; use crate::ty::LayoutLlvmExt; // libnvvm does not support some advanced intrinsics for i128 so we just abort on them for now. In the future @@ -159,9 +159,8 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { fn codegen_intrinsic_call( &mut self, instance: ty::Instance<'tcx>, - fn_abi: &FnAbi<'tcx, Ty<'tcx>>, args: &[OperandRef<'tcx, &'ll Value>], - llresult: &'ll Value, + result: PlaceRef<'tcx, &'ll Value>, span: Span, ) -> Result<(), ty::Instance<'tcx>> { let tcx = self.tcx; @@ -184,7 +183,9 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { ); let llret_ty = self.layout_of(ret_ty).llvm_type(self); - let result = PlaceRef::new_sized(llresult, fn_abi.ret.layout); + + // Compute fn_abi for intrinsics that need it + let fn_abi = self.cx.fn_abi_of_instance(instance, ty::List::empty()); let simple = get_simple_intrinsic(self, name); let llval = match name { @@ -244,7 +245,7 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { self.call(self.type_i1(), None, None, try_func, &[data], None, None); let ret_align = self.data_layout().i32_align.abi; - self.store(self.const_i32(0), llresult, ret_align) + self.store(self.const_i32(0), result.val.llval, ret_align) } sym::breakpoint => { // debugtrap is not supported @@ -254,7 +255,7 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { self.call_intrinsic("llvm.va_copy", &[args[0].immediate(), args[1].immediate()]) } sym::va_arg => { - match fn_abi.ret.layout.backend_repr { + match result.layout.backend_repr { abi::BackendRepr::Scalar(scalar) => { match scalar.primitive() { Primitive::Int(..) => { @@ -295,16 +296,16 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { } } sym::volatile_load | sym::unaligned_volatile_load => { - let tp_ty = fn_args.type_at(0); let mut ptr = args[0].immediate(); + // Handle cast if the ABI requires it if let PassMode::Cast { cast: ty, .. } = &fn_abi.ret.mode { ptr = self.pointercast(ptr, self.type_ptr_to(ty.llvm_type(self))); } - let load = self.volatile_load(self.type_i1(), ptr); + let load = self.volatile_load(result.layout.llvm_type(self), ptr); let align = if name == sym::unaligned_volatile_load { 1 } else { - self.align_of(tp_ty).bytes() as u32 + result.layout.align.abi.bytes() as u32 }; unsafe { llvm::LLVMSetAlignment(load, align); @@ -408,16 +409,16 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { match name { sym::ctlz | sym::cttz => { let y = self.const_bool(false); - let llvm_name = format!("llvm.{}.i{}", name, width); + let llvm_name = format!("llvm.{name}.i{width}"); self.call_intrinsic(&llvm_name, &[args[0].immediate(), y]) } sym::ctlz_nonzero | sym::cttz_nonzero => { let y = self.const_bool(true); - let llvm_name = format!("llvm.{}.i{}", &name_str[..4], width); + let llvm_name = format!("llvm.{}.i{width}", &name_str[..4]); self.call_intrinsic(&llvm_name, &[args[0].immediate(), y]) } sym::ctpop => self.call_intrinsic( - &format!("llvm.ctpop.i{}", width), + &format!("llvm.ctpop.i{width}"), &[args[0].immediate()], ), sym::bswap => { @@ -425,13 +426,13 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { args[0].immediate() // byte swap a u8/i8 is just a no-op } else { self.call_intrinsic( - &format!("llvm.bswap.i{}", width), + &format!("llvm.bswap.i{width}"), &[args[0].immediate()], ) } } sym::bitreverse => self.call_intrinsic( - &format!("llvm.bitreverse.i{}", width), + &format!("llvm.bitreverse.i{width}"), &[args[0].immediate()], ), sym::rotate_left | sym::rotate_right => { @@ -465,14 +466,13 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { } } sym::raw_eq => { - use abi::BackendRepr::*; use rustc_codegen_ssa::common::IntPredicate; let tp_ty = fn_args.type_at(0); let layout = self.layout_of(tp_ty).layout; let use_integer_compare = match layout.backend_repr() { - Scalar(_) | ScalarPair(_, _) => true, - Vector { .. } => false, - Memory { .. } => { + BackendRepr::Scalar(_) | BackendRepr::ScalarPair(_, _) => true, + BackendRepr::SimdVector { .. } => false, + BackendRepr::Memory { .. } => { // For rusty ABIs, small aggregates are actually passed // as `RegKind::Integer` (see `FnAbi::adjust_for_abi`), // so we re-use that same threshold here. @@ -577,11 +577,6 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { self.call_intrinsic("llvm.expect.i1", &[cond, self.const_bool(expected)]) } - fn type_test(&mut self, _pointer: &'ll Value, _typeid: &'ll Metadata) -> &'ll Value { - // LLVM CFI doesnt make sense on the GPU - self.const_i32(0) - } - fn type_checked_load( &mut self, _llvtable: Self::Value, diff --git a/crates/rustc_codegen_nvvm/src/lib.rs b/crates/rustc_codegen_nvvm/src/lib.rs index ba05e279..b0c2da34 100644 --- a/crates/rustc_codegen_nvvm/src/lib.rs +++ b/crates/rustc_codegen_nvvm/src/lib.rs @@ -3,13 +3,13 @@ // make our lives a lot easier for llvm ffi with this. And since rustc's core infra // relies on it its almost guaranteed to not be removed/broken #![feature(extern_types)] -#![feature(hash_raw_entry)] #![feature(let_chains)] #![feature(slice_as_array)] extern crate rustc_abi; extern crate rustc_arena; extern crate rustc_ast; +extern crate rustc_attr_data_structures; extern crate rustc_attr_parsing; extern crate rustc_codegen_ssa; extern crate rustc_data_structures; @@ -70,7 +70,6 @@ use rustc_codegen_ssa::{ }; use rustc_data_structures::fx::FxIndexMap; use rustc_errors::{DiagCtxtHandle, FatalError}; -use rustc_metadata::EncodedMetadata; use rustc_metadata::creader::MetadataLoaderDyn; use rustc_middle::util::Providers; use rustc_middle::{ @@ -147,12 +146,7 @@ impl CodegenBackend for NvvmCodegenBackend { }; } - fn codegen_crate( - &self, - tcx: TyCtxt<'_>, - metadata: EncodedMetadata, - need_metadata_module: bool, - ) -> Box { + fn codegen_crate(&self, tcx: TyCtxt<'_>) -> Box { debug!("Codegen crate"); Box::new(rustc_codegen_ssa::base::codegen_crate( Self, @@ -163,8 +157,6 @@ impl CodegenBackend for NvvmCodegenBackend { .target_cpu .clone() .unwrap_or_else(|| tcx.sess.target.cpu.to_string()), - metadata, - need_metadata_module, )) } @@ -189,6 +181,7 @@ impl CodegenBackend for NvvmCodegenBackend { &self, sess: &rustc_session::Session, codegen_results: rustc_codegen_ssa::CodegenResults, + metadata: rustc_metadata::EncodedMetadata, outputs: &config::OutputFilenames, ) { link::link( @@ -196,6 +189,7 @@ impl CodegenBackend for NvvmCodegenBackend { &codegen_results, outputs, codegen_results.crate_info.local_crate_name.as_str(), + metadata, ); } } @@ -245,7 +239,7 @@ impl WriteBackendMethods for NvvmCodegenBackend { // Not applicable, nvvm doesnt expose pass timing info, maybe we could print llvm pass stuff here. } - unsafe fn optimize( + fn optimize( cgcx: &CodegenContext, diag_handler: DiagCtxtHandle<'_>, module: &mut ModuleCodegen, @@ -254,14 +248,14 @@ impl WriteBackendMethods for NvvmCodegenBackend { unsafe { back::optimize(cgcx, diag_handler, module, config) } } - unsafe fn optimize_thin( + fn optimize_thin( cgcx: &CodegenContext, thin_module: ThinModule, ) -> Result, FatalError> { unsafe { lto::optimize_thin(cgcx, thin_module) } } - unsafe fn codegen( + fn codegen( cgcx: &CodegenContext, diag_handler: DiagCtxtHandle<'_>, module: ModuleCodegen, diff --git a/crates/rustc_codegen_nvvm/src/link.rs b/crates/rustc_codegen_nvvm/src/link.rs index b545d283..9cfe1e5a 100644 --- a/crates/rustc_codegen_nvvm/src/link.rs +++ b/crates/rustc_codegen_nvvm/src/link.rs @@ -77,8 +77,8 @@ fn read_metadata(rlib: &Path) -> Result { match read_meta() { Ok(Some(m)) => Ok(m), - Ok(None) => Err(format!("No .metadata file in rlib: {:?}", rlib)), - Err(io) => Err(format!("Failed to read rlib at {:?}: {}", rlib, io)), + Ok(None) => Err(format!("No .metadata file in rlib: {rlib:?}")), + Err(io) => Err(format!("Failed to read rlib at {rlib:?}: {io}")), } } @@ -110,6 +110,7 @@ pub fn link( codegen_results: &CodegenResults, outputs: &OutputFilenames, crate_name: &str, + metadata: rustc_metadata::EncodedMetadata, ) { debug!("Linking crate `{}`", crate_name); // largely inspired by rust-gpu @@ -133,10 +134,15 @@ pub fn link( if outputs.outputs.should_codegen() { let out_filename = out_filename(sess, crate_type, outputs, Symbol::intern(crate_name)); let out_filename_file_for_writing = - out_filename.file_for_writing(outputs, OutputType::Exe, None); + out_filename.file_for_writing(outputs, OutputType::Exe, "", None); match crate_type { CrateType::Rlib => { - link_rlib(sess, codegen_results, &out_filename_file_for_writing); + link_rlib( + sess, + codegen_results, + &out_filename_file_for_writing, + &metadata, + ); } CrateType::Executable | CrateType::Cdylib | CrateType::Dylib => { let _ = link_exe( @@ -147,13 +153,18 @@ pub fn link( codegen_results, ); } - other => sess.dcx().fatal(format!("Invalid crate type: {:?}", other)), + other => sess.dcx().fatal(format!("Invalid crate type: {other:?}")), } } } } -fn link_rlib(sess: &Session, codegen_results: &CodegenResults, out_filename: &Path) { +fn link_rlib( + sess: &Session, + codegen_results: &CodegenResults, + out_filename: &Path, + metadata: &rustc_metadata::EncodedMetadata, +) { debug!("Linking rlib `{:?}`", out_filename); let mut file_list = Vec::<&Path>::new(); @@ -184,12 +195,7 @@ fn link_rlib(sess: &Session, codegen_results: &CodegenResults, out_filename: &Pa } trace!("Files linked in rlib:\n{:#?}", file_list); - create_archive( - sess, - &file_list, - codegen_results.metadata.raw_data(), - out_filename, - ); + create_archive(sess, &file_list, metadata.stub_or_full(), out_filename); } fn link_exe( @@ -304,8 +310,7 @@ fn codegen_into_ptx_file( fn create_archive(sess: &Session, files: &[&Path], metadata: &[u8], out_filename: &Path) { if let Err(err) = try_create_archive(files, metadata, out_filename) { - sess.dcx() - .fatal(format!("Failed to create archive: {}", err)); + sess.dcx().fatal(format!("Failed to create archive: {err}")); } } diff --git a/crates/rustc_codegen_nvvm/src/lto.rs b/crates/rustc_codegen_nvvm/src/lto.rs index 9e124a65..7d5474ff 100644 --- a/crates/rustc_codegen_nvvm/src/lto.rs +++ b/crates/rustc_codegen_nvvm/src/lto.rs @@ -87,6 +87,7 @@ impl Drop for ThinBuffer { } } +#[allow(dead_code)] pub struct ThinData(&'static mut llvm::ThinLTOData); unsafe impl Send for ThinData {} diff --git a/crates/rustc_codegen_nvvm/src/mono_item.rs b/crates/rustc_codegen_nvvm/src/mono_item.rs index d5e2a3dd..1145668f 100644 --- a/crates/rustc_codegen_nvvm/src/mono_item.rs +++ b/crates/rustc_codegen_nvvm/src/mono_item.rs @@ -23,7 +23,7 @@ pub(crate) fn visibility_to_llvm(linkage: Visibility) -> llvm::Visibility { impl<'tcx> PreDefineCodegenMethods<'tcx> for CodegenCx<'_, 'tcx> { fn predefine_static( - &self, + &mut self, def_id: DefId, linkage: Linkage, visibility: Visibility, @@ -40,7 +40,7 @@ impl<'tcx> PreDefineCodegenMethods<'tcx> for CodegenCx<'_, 'tcx> { .unwrap_or_else(|| { self.sess().dcx().span_fatal( self.tcx.def_span(def_id), - format!("symbol `{}` is already defined", symbol_name), + format!("symbol `{symbol_name}` is already defined"), ) }); @@ -53,7 +53,7 @@ impl<'tcx> PreDefineCodegenMethods<'tcx> for CodegenCx<'_, 'tcx> { } fn predefine_fn( - &self, + &mut self, instance: Instance<'tcx>, linkage: Linkage, visibility: Visibility, diff --git a/crates/rustc_codegen_nvvm/src/nvvm.rs b/crates/rustc_codegen_nvvm/src/nvvm.rs index 5e50db7a..2bee80e3 100644 --- a/crates/rustc_codegen_nvvm/src/nvvm.rs +++ b/crates/rustc_codegen_nvvm/src/nvvm.rs @@ -118,8 +118,7 @@ pub fn codegen_bitcode_modules( let log = prog.compiler_log().unwrap().unwrap_or_default(); let footer = "If you plan to submit a bug report please re-run the codegen with `RUSTFLAGS=\"--emit=llvm-ir\" and include the .ll file corresponding to the .o file mentioned in the log"; panic!( - "Malformed NVVM IR program rejected by libnvvm, dumping verifier log:\n\n{}\n\n{}", - log, footer + "Malformed NVVM IR program rejected by libnvvm, dumping verifier log:\n\n{log}\n\n{footer}" ); } @@ -128,8 +127,7 @@ pub fn codegen_bitcode_modules( Err(error) => { // this should never happen, if it does, something went really bad or its a bug on libnvvm's end panic!( - "libnvvm returned an error that was not previously caught by the verifier: {:?}", - error + "libnvvm returned an error that was not previously caught by the verifier: {error:?}" ); } }; diff --git a/crates/rustc_codegen_nvvm/src/override_fns.rs b/crates/rustc_codegen_nvvm/src/override_fns.rs index 9f44fe4c..e4c0ee23 100644 --- a/crates/rustc_codegen_nvvm/src/override_fns.rs +++ b/crates/rustc_codegen_nvvm/src/override_fns.rs @@ -9,16 +9,26 @@ use crate::llvm; use rustc_codegen_ssa::mono_item::MonoItemExt; use rustc_codegen_ssa::traits::{BaseTypeCodegenMethods, BuilderMethods}; use rustc_hir::def_id::LOCAL_CRATE; -use rustc_middle::mir::mono::MonoItem; +use rustc_middle::mir::mono::{Linkage, MonoItem, MonoItemData, Visibility}; use rustc_middle::ty::layout::FnAbiOf; use rustc_middle::ty::{self, Instance}; /// Either override or define a function. -pub(crate) fn define_or_override_fn<'tcx>(func: Instance<'tcx>, cx: &CodegenCx<'_, 'tcx>) { +pub(crate) fn define_or_override_fn<'tcx>(func: Instance<'tcx>, cx: &mut CodegenCx<'_, 'tcx>) { if should_override(func, cx) { override_libm_function(func, cx); } else { - MonoItem::define::>(&MonoItem::Fn(func), cx); + MonoItem::define::>( + &MonoItem::Fn(func), + cx, + "mono_item", + MonoItemData { + inlined: false, + linkage: Linkage::External, + visibility: Visibility::Default, + size_estimate: 0, + }, + ); } } @@ -40,7 +50,7 @@ fn should_override<'tcx>(func: Instance<'tcx>, cx: &CodegenCx<'_, 'tcx>) -> bool return false; } - let libdevice_name = format!("__nv_{}", name); + let libdevice_name = format!("__nv_{name}"); let ld_fn = if let Some((args, ret)) = cx.intrinsics_map.borrow().get(libdevice_name.as_str()) { cx.type_func(args, ret) } else { diff --git a/crates/rustc_codegen_nvvm/src/ty.rs b/crates/rustc_codegen_nvvm/src/ty.rs index cd7963ed..6bfd0bbd 100644 --- a/crates/rustc_codegen_nvvm/src/ty.rs +++ b/crates/rustc_codegen_nvvm/src/ty.rs @@ -157,7 +157,7 @@ impl<'ll> CodegenCx<'ll, '_> { } } -impl<'ll, 'tcx> BaseTypeCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { +impl<'ll, 'tcx> BaseTypeCodegenMethods for CodegenCx<'ll, 'tcx> { fn type_i8(&self) -> &'ll Type { unsafe { llvm::LLVMInt8TypeInContext(self.llcx) } } @@ -219,8 +219,7 @@ impl<'ll, 'tcx> BaseTypeCodegenMethods<'tcx> for CodegenCx<'ll, 'tcx> { } fn element_type(&self, ty: &'ll Type) -> &'ll Type { - let out = unsafe { llvm::LLVMGetElementType(ty) }; - out + unsafe { llvm::LLVMGetElementType(ty) } } fn vector_length(&self, ty: &'ll Type) -> usize { @@ -300,7 +299,7 @@ pub(crate) trait LayoutLlvmExt<'tcx> { impl<'tcx> LayoutLlvmExt<'tcx> for TyAndLayout<'tcx> { fn is_llvm_immediate(&self) -> bool { match self.backend_repr { - BackendRepr::Scalar(_) | BackendRepr::Vector { .. } => true, + BackendRepr::Scalar(_) | BackendRepr::SimdVector { .. } => true, BackendRepr::ScalarPair(..) => false, BackendRepr::Memory { .. } => self.is_zst(), } @@ -309,9 +308,9 @@ impl<'tcx> LayoutLlvmExt<'tcx> for TyAndLayout<'tcx> { fn is_llvm_scalar_pair(&self) -> bool { match self.backend_repr { BackendRepr::ScalarPair(..) => true, - BackendRepr::Scalar(_) | BackendRepr::Vector { .. } | BackendRepr::Memory { .. } => { - false - } + BackendRepr::Scalar(_) + | BackendRepr::SimdVector { .. } + | BackendRepr::Memory { .. } => false, } } @@ -515,7 +514,7 @@ fn uncached_llvm_type<'a, 'tcx>( trace!("Uncached LLVM type of {:?}", layout); match layout.backend_repr { BackendRepr::Scalar(_) => bug!("handled elsewhere"), - BackendRepr::Vector { element, count } => { + BackendRepr::SimdVector { element, count } => { let element = layout.scalar_llvm_type_at(cx, element); return cx.type_vector(element, count); } @@ -542,10 +541,10 @@ fn uncached_llvm_type<'a, 'tcx>( let mut name = with_no_trimmed_paths!(layout.ty.to_string()); if let (&ty::Adt(def, _), &Variants::Single { index }) = (layout.ty.kind(), &layout.variants) + && def.is_enum() + && !def.variants().is_empty() { - if def.is_enum() && !def.variants().is_empty() { - write!(&mut name, "::{}", def.variant(index).name).unwrap(); - } + write!(&mut name, "::{}", def.variant(index).name).unwrap(); } if let (&ty::Coroutine(_, _), &Variants::Single { index }) = (layout.ty.kind(), &layout.variants) diff --git a/examples/cuda/gemm/src/main.rs b/examples/cuda/gemm/src/main.rs index ff8708c2..8e254176 100644 --- a/examples/cuda/gemm/src/main.rs +++ b/examples/cuda/gemm/src/main.rs @@ -169,7 +169,7 @@ fn run_cublas(stream: &stream::Stream) -> Result<(), Box> { let mut mat_c_actual = Array::::zeros((sz, sz)); mat_c_gpu.copy_to(&mut mat_c_actual.as_slice_mut().unwrap())?; let duration = end.elapsed_time_f32(&beg)? / (NUM_RUNS as f32); - println!("cuBLAS {}x{}: {:.4}ms", sz, sz, duration); + println!("cuBLAS {sz}x{sz}: {duration:.4}ms"); if sz < 1024 { assert_gemm_eq(&mat_a, &mat_b, &mat_c, alpha, beta, &mat_c_actual); } @@ -278,7 +278,7 @@ fn run_gemm_kernel( let mut mat_c_actual = Array::::zeros((sz, sz)); mat_c_gpu.copy_to(&mut mat_c_actual.as_slice_mut().unwrap())?; let duration = end.elapsed_time_f32(&beg)? / (NUM_RUNS as f32); - println!("{} {}x{}: {:.4}ms", kernel_name, sz, sz, duration); + println!("{kernel_name} {sz}x{sz}: {duration:.4}ms"); if sz < 1024 { assert_gemm_eq(&mat_a, &mat_b, &mat_c, alpha, beta, &mat_c_actual); } @@ -315,8 +315,8 @@ fn assert_gemm_eq( let mat_c_expect = alpha * mat_a.dot(&mat_b) + beta * mat_c; let ok = mat_c_expect.relative_eq(&mat_c_actual, EPS.into(), EPS.into()); if !ok { - println!("Actual: {:?}", mat_c_actual); - println!("Expect: {:?}", mat_c_expect); + println!("Actual: {mat_c_actual:?}"); + println!("Expect: {mat_c_expect:?}"); panic!("GEMM result mismatch"); } } diff --git a/examples/cuda/vecadd/src/main.rs b/examples/cuda/vecadd/src/main.rs index c6263f02..cf804d0c 100644 --- a/examples/cuda/vecadd/src/main.rs +++ b/examples/cuda/vecadd/src/main.rs @@ -47,10 +47,7 @@ fn main() -> Result<(), Box> { let grid_size = (NUMBERS_LEN as u32).div_ceil(block_size); - println!( - "using {} blocks and {} threads per block", - grid_size, block_size - ); + println!("using {grid_size} blocks and {block_size} threads per block"); // Actually launch the GPU kernel. This will queue up the launch on the stream, it will // not block the thread until the kernel is finished. diff --git a/rust-toolchain.toml b/rust-toolchain.toml index c8054e8d..858caf73 100644 --- a/rust-toolchain.toml +++ b/rust-toolchain.toml @@ -1,3 +1,3 @@ [toolchain] -channel = "nightly-2025-03-02" +channel = "nightly-2025-06-23" components = ["clippy", "llvm-tools-preview", "rust-src", "rustc-dev", "rustfmt", "rust-analyzer"] \ No newline at end of file diff --git a/xtask/src/extract_llfns.rs b/xtask/src/extract_llfns.rs index c79ad5cb..d6433506 100644 --- a/xtask/src/extract_llfns.rs +++ b/xtask/src/extract_llfns.rs @@ -35,7 +35,7 @@ pub(crate) fn extract_llfns(file: &Path, dir: &Path) { let out_file = format!("{}/{}.ll", dir.display(), name); let _ = Command::new("llvm-extract") .arg(file) - .arg(format!("--func={}", name)) + .arg(format!("--func={name}")) .arg("-S") .arg("--recursive") .arg("-o") @@ -52,13 +52,13 @@ pub(crate) fn extract_llfns(file: &Path, dir: &Path) { for (name, content, failed) in &mut contents { if PRINT_EVERY_EXECUTION { - println!("Running command over `{}.ll`", name); + println!("Running command over `{name}.ll`"); } *failed = !run_command_for_each_fn(name, content); } for (name, _, _) in contents.into_iter().filter(|x| !x.2).take(30) { - println!("Err: {}", name); + println!("Err: {name}"); } }