diff --git a/compiler/rustc_codegen_llvm/src/intrinsic.rs b/compiler/rustc_codegen_llvm/src/intrinsic.rs index 481f75f337d63..c6aae89f1e51e 100644 --- a/compiler/rustc_codegen_llvm/src/intrinsic.rs +++ b/compiler/rustc_codegen_llvm/src/intrinsic.rs @@ -560,6 +560,12 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> { return Ok(()); } + sym::amdgpu_dispatch_ptr => { + let val = self.call_intrinsic("llvm.amdgcn.dispatch.ptr", &[], &[]); + // Relying on `LLVMBuildPointerCast` to produce an addrspacecast + self.pointercast(val, self.type_ptr()) + } + _ if name.as_str().starts_with("simd_") => { // Unpack non-power-of-2 #[repr(packed, simd)] arguments. // This gives them the expected layout of a regular #[repr(simd)] vector. diff --git a/compiler/rustc_codegen_ssa/src/mir/intrinsic.rs b/compiler/rustc_codegen_ssa/src/mir/intrinsic.rs index f4fae40d8828f..f5ee9406f4bf1 100644 --- a/compiler/rustc_codegen_ssa/src/mir/intrinsic.rs +++ b/compiler/rustc_codegen_ssa/src/mir/intrinsic.rs @@ -112,6 +112,7 @@ impl<'a, 'tcx, Bx: BuilderMethods<'a, 'tcx>> FunctionCx<'a, 'tcx, Bx> { | sym::unreachable | sym::cold_path | sym::breakpoint + | sym::amdgpu_dispatch_ptr | sym::assert_zero_valid | sym::assert_mem_uninitialized_valid | sym::assert_inhabited diff --git a/compiler/rustc_hir_analysis/src/check/intrinsic.rs b/compiler/rustc_hir_analysis/src/check/intrinsic.rs index c84c1a8ca16d8..d3d167f6e2544 100644 --- a/compiler/rustc_hir_analysis/src/check/intrinsic.rs +++ b/compiler/rustc_hir_analysis/src/check/intrinsic.rs @@ -70,6 +70,7 @@ fn intrinsic_operation_unsafety(tcx: TyCtxt<'_>, intrinsic_id: LocalDefId) -> hi | sym::add_with_overflow | sym::aggregate_raw_ptr | sym::align_of + | sym::amdgpu_dispatch_ptr | sym::assert_inhabited | sym::assert_mem_uninitialized_valid | sym::assert_zero_valid @@ -286,6 +287,7 @@ pub(crate) fn check_intrinsic_type( let (n_tps, n_cts, inputs, output) = match intrinsic_name { sym::autodiff => (4, 0, vec![param(0), param(1), param(2)], param(3)), sym::abort => (0, 0, vec![], tcx.types.never), + sym::amdgpu_dispatch_ptr => (0, 0, vec![], Ty::new_imm_ptr(tcx, tcx.types.unit)), sym::unreachable => (0, 0, vec![], tcx.types.never), sym::breakpoint => (0, 0, vec![], tcx.types.unit), sym::size_of | sym::align_of | sym::variant_count => (1, 0, vec![], tcx.types.usize), diff --git a/compiler/rustc_mir_build/src/builder/matches/match_pair.rs b/compiler/rustc_mir_build/src/builder/matches/match_pair.rs index 3edd0234b0ad7..e80e29415e6f0 100644 --- a/compiler/rustc_mir_build/src/builder/matches/match_pair.rs +++ b/compiler/rustc_mir_build/src/builder/matches/match_pair.rs @@ -2,7 +2,6 @@ use std::sync::Arc; use rustc_abi::FieldIdx; use rustc_middle::mir::*; -use rustc_middle::span_bug; use rustc_middle::thir::*; use rustc_middle::ty::{self, Ty, TypeVisitableExt}; @@ -160,10 +159,7 @@ impl<'tcx> MatchPairTree<'tcx> { } PatKind::Constant { value } => { - // CAUTION: The type of the pattern node (`pattern.ty`) is - // _often_ the same as the type of the const value (`value.ty`), - // but there are some cases where those types differ - // (e.g. when `deref!(..)` patterns interact with `String`). + assert_eq!(pattern.ty, value.ty); // Classify the constant-pattern into further kinds, to // reduce the number of ad-hoc type tests needed later on. @@ -175,16 +171,6 @@ impl<'tcx> MatchPairTree<'tcx> { } else if pat_ty.is_floating_point() { PatConstKind::Float } else if pat_ty.is_str() { - // Deref-patterns can cause string-literal patterns to have - // type `str` instead of the usual `&str`. - if !cx.tcx.features().deref_patterns() { - span_bug!( - pattern.span, - "const pattern has type `str` but deref_patterns is not enabled" - ); - } - PatConstKind::String - } else if pat_ty.is_imm_ref_str() { PatConstKind::String } else { // FIXME(Zalathar): This still covers several different diff --git a/compiler/rustc_mir_build/src/builder/matches/mod.rs b/compiler/rustc_mir_build/src/builder/matches/mod.rs index 0463f7c914a4a..11a181cfa8cea 100644 --- a/compiler/rustc_mir_build/src/builder/matches/mod.rs +++ b/compiler/rustc_mir_build/src/builder/matches/mod.rs @@ -1339,19 +1339,13 @@ enum TestKind<'tcx> { /// Tests the place against a string constant using string equality. StringEq { - /// Constant `&str` value to test against. + /// Constant string value to test against. + /// Note that this value has type `str` (not `&str`). value: ty::Value<'tcx>, - /// Type of the corresponding pattern node. Usually `&str`, but could - /// be `str` for patterns like `deref!("..."): String`. - pat_ty: Ty<'tcx>, }, /// Tests the place against a constant using scalar equality. - ScalarEq { - value: ty::Value<'tcx>, - /// Type of the corresponding pattern node. - pat_ty: Ty<'tcx>, - }, + ScalarEq { value: ty::Value<'tcx> }, /// Test whether the value falls within an inclusive or exclusive range. Range(Arc>), diff --git a/compiler/rustc_mir_build/src/builder/matches/test.rs b/compiler/rustc_mir_build/src/builder/matches/test.rs index c2e39d47a92ca..5c3173a7b1488 100644 --- a/compiler/rustc_mir_build/src/builder/matches/test.rs +++ b/compiler/rustc_mir_build/src/builder/matches/test.rs @@ -9,10 +9,10 @@ use std::sync::Arc; use rustc_data_structures::fx::FxIndexMap; use rustc_hir::{LangItem, RangeEnd}; +use rustc_middle::bug; use rustc_middle::mir::*; use rustc_middle::ty::util::IntTypeExt; use rustc_middle::ty::{self, GenericArg, Ty, TyCtxt}; -use rustc_middle::{bug, span_bug}; use rustc_span::def_id::DefId; use rustc_span::source_map::Spanned; use rustc_span::{DUMMY_SP, Span, Symbol, sym}; @@ -39,10 +39,10 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { TestKind::SwitchInt } TestableCase::Constant { value, kind: PatConstKind::String } => { - TestKind::StringEq { value, pat_ty: match_pair.pattern_ty } + TestKind::StringEq { value } } TestableCase::Constant { value, kind: PatConstKind::Float | PatConstKind::Other } => { - TestKind::ScalarEq { value, pat_ty: match_pair.pattern_ty } + TestKind::ScalarEq { value } } TestableCase::Range(ref range) => { @@ -141,47 +141,33 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { self.cfg.terminate(block, self.source_info(match_start_span), terminator); } - TestKind::StringEq { value, pat_ty } => { + TestKind::StringEq { value } => { let tcx = self.tcx; let success_block = target_block(TestBranch::Success); let fail_block = target_block(TestBranch::Failure); - let expected_value_ty = value.ty; - let expected_value_operand = - self.literal_operand(test.span, Const::from_ty_value(tcx, value)); + let ref_str_ty = Ty::new_imm_ref(tcx, tcx.lifetimes.re_erased, tcx.types.str_); + assert!(ref_str_ty.is_imm_ref_str(), "{ref_str_ty:?}"); - let mut actual_value_ty = pat_ty; - let mut actual_value_place = place; - - match pat_ty.kind() { - ty::Str => { - // String literal patterns may have type `str` if `deref_patterns` is - // enabled, in order to allow `deref!("..."): String`. In this case, `value` - // is of type `&str`, so we compare it to `&place`. - if !tcx.features().deref_patterns() { - span_bug!( - test.span, - "matching on `str` went through without enabling deref_patterns" - ); - } - let re_erased = tcx.lifetimes.re_erased; - let ref_str_ty = Ty::new_imm_ref(tcx, re_erased, tcx.types.str_); - let ref_place = self.temp(ref_str_ty, test.span); - // `let ref_place: &str = &place;` - self.cfg.push_assign( - block, - self.source_info(test.span), - ref_place, - Rvalue::Ref(re_erased, BorrowKind::Shared, place), - ); - actual_value_place = ref_place; - actual_value_ty = ref_str_ty; - } - _ => {} - } + // The string constant we're testing against has type `str`, but + // calling `::eq` requires `&str` operands. + // + // Because `str` and `&str` have the same valtree representation, + // we can "cast" to the desired type by just replacing the type. + assert!(value.ty.is_str(), "unexpected value type for StringEq test: {value:?}"); + let expected_value = ty::Value { ty: ref_str_ty, valtree: value.valtree }; + let expected_value_operand = + self.literal_operand(test.span, Const::from_ty_value(tcx, expected_value)); - assert_eq!(expected_value_ty, actual_value_ty); - assert!(actual_value_ty.is_imm_ref_str()); + // Similarly, the scrutinized place has type `str`, but we need `&str`. + // Get a reference by doing `let actual_value_ref_place: &str = &place`. + let actual_value_ref_place = self.temp(ref_str_ty, test.span); + self.cfg.push_assign( + block, + self.source_info(test.span), + actual_value_ref_place, + Rvalue::Ref(tcx.lifetimes.re_erased, BorrowKind::Shared, place), + ); // Compare two strings using `::eq`. // (Interestingly this means that exhaustiveness analysis relies, for soundness, @@ -192,11 +178,11 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { fail_block, source_info, expected_value_operand, - Operand::Copy(actual_value_place), + Operand::Copy(actual_value_ref_place), ); } - TestKind::ScalarEq { value, pat_ty } => { + TestKind::ScalarEq { value } => { let tcx = self.tcx; let success_block = target_block(TestBranch::Success); let fail_block = target_block(TestBranch::Failure); @@ -205,12 +191,10 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { let mut expected_value_operand = self.literal_operand(test.span, Const::from_ty_value(tcx, value)); - let mut actual_value_ty = pat_ty; let mut actual_value_place = place; - match pat_ty.kind() { + match value.ty.kind() { &ty::Pat(base, _) => { - assert_eq!(pat_ty, value.ty); assert!(base.is_trivially_pure_clone_copy()); let transmuted_place = self.temp(base, test.span); @@ -234,15 +218,13 @@ impl<'a, 'tcx> Builder<'a, 'tcx> { ); actual_value_place = transmuted_place; - actual_value_ty = base; expected_value_operand = Operand::Copy(transmuted_expect); expected_value_ty = base; } _ => {} } - assert_eq!(expected_value_ty, actual_value_ty); - assert!(actual_value_ty.is_scalar()); + assert!(expected_value_ty.is_scalar()); self.compare( block, diff --git a/compiler/rustc_mir_build/src/thir/pattern/const_to_pat.rs b/compiler/rustc_mir_build/src/thir/pattern/const_to_pat.rs index 02409d2bae9fa..70bc142131e4f 100644 --- a/compiler/rustc_mir_build/src/thir/pattern/const_to_pat.rs +++ b/compiler/rustc_mir_build/src/thir/pattern/const_to_pat.rs @@ -289,32 +289,29 @@ impl<'tcx> ConstToPat<'tcx> { suffix: Box::new([]), }, ty::Str => { - // String literal patterns may have type `str` if `deref_patterns` is enabled, in - // order to allow `deref!("..."): String`. Since we need a `&str` for the comparison - // when lowering to MIR in `Builder::perform_test`, treat the constant as a `&str`. - // This works because `str` and `&str` have the same valtree representation. - let ref_str_ty = Ty::new_imm_ref(tcx, tcx.lifetimes.re_erased, ty); - PatKind::Constant { value: ty::Value { ty: ref_str_ty, valtree: cv } } + // Constant/literal patterns of type `&str` are lowered to a + // `PatKind::Deref` wrapping a `PatKind::Constant` of type `str`. + // This pattern node is the `str` constant part. + // + // Under `feature(deref_patterns)`, string literal patterns can also + // have type `str` directly, without the `&`, in order to allow things + // like `deref!("...")` to work when the scrutinee is `String`. + PatKind::Constant { value: ty::Value { ty, valtree: cv } } } - ty::Ref(_, pointee_ty, ..) => match *pointee_ty.kind() { - // `&str` is represented as a valtree, let's keep using this - // optimization for now. - ty::Str => PatKind::Constant { value: ty::Value { ty, valtree: cv } }, - // All other references are converted into deref patterns and then recursively - // convert the dereferenced constant to a pattern that is the sub-pattern of the - // deref pattern. - _ => { - if !pointee_ty.is_sized(tcx, self.typing_env) && !pointee_ty.is_slice() { - return self.mk_err( - tcx.dcx().create_err(UnsizedPattern { span, non_sm_ty: *pointee_ty }), - ty, - ); - } else { - // References have the same valtree representation as their pointee. - PatKind::Deref { subpattern: self.valtree_to_pat(cv, *pointee_ty) } - } + ty::Ref(_, pointee_ty, ..) => { + if pointee_ty.is_str() + || pointee_ty.is_slice() + || pointee_ty.is_sized(tcx, self.typing_env) + { + // References have the same valtree representation as their pointee. + PatKind::Deref { subpattern: self.valtree_to_pat(cv, *pointee_ty) } + } else { + return self.mk_err( + tcx.dcx().create_err(UnsizedPattern { span, non_sm_ty: *pointee_ty }), + ty, + ); } - }, + } ty::Float(flt) => { let v = cv.to_leaf(); let is_nan = match flt { diff --git a/compiler/rustc_pattern_analysis/src/rustc.rs b/compiler/rustc_pattern_analysis/src/rustc.rs index 721635ed48ff5..5e75192ff3092 100644 --- a/compiler/rustc_pattern_analysis/src/rustc.rs +++ b/compiler/rustc_pattern_analysis/src/rustc.rs @@ -583,19 +583,13 @@ impl<'p, 'tcx: 'p> RustcPatCtxt<'p, 'tcx> { fields = vec![]; arity = 0; } - ty::Ref(_, t, _) if t.is_str() => { - // We want a `&str` constant to behave like a `Deref` pattern, to be compatible - // with other `Deref` patterns. This could have been done in `const_to_pat`, - // but that causes issues with the rest of the matching code. - // So here, the constructor for a `"foo"` pattern is `&` (represented by - // `Ref`), and has one field. That field has constructor `Str(value)` and no - // subfields. - // Note: `t` is `str`, not `&str`. - let ty = self.reveal_opaque_ty(*t); - let subpattern = DeconstructedPat::new(Str(*value), Vec::new(), 0, ty, pat); - ctor = Ref; - fields = vec![subpattern.at_index(0)]; - arity = 1; + ty::Str => { + // For constant/literal patterns of type `&str`, the THIR + // pattern is a `PatKind::Deref` of type `&str` wrapping a + // `PatKind::Const` of type `str`. + ctor = Str(*value); + fields = vec![]; + arity = 0; } // All constants that can be structurally matched have already been expanded // into the corresponding `Pat`s by `const_to_pat`. Constants that remain are diff --git a/compiler/rustc_span/src/symbol.rs b/compiler/rustc_span/src/symbol.rs index 4080c1cd59ec7..98db0df593bec 100644 --- a/compiler/rustc_span/src/symbol.rs +++ b/compiler/rustc_span/src/symbol.rs @@ -458,6 +458,7 @@ symbols! { alu32, always, amdgpu, + amdgpu_dispatch_ptr, analysis, and, and_then, diff --git a/library/core/src/intrinsics/gpu.rs b/library/core/src/intrinsics/gpu.rs new file mode 100644 index 0000000000000..9e7624841d0c6 --- /dev/null +++ b/library/core/src/intrinsics/gpu.rs @@ -0,0 +1,23 @@ +//! Intrinsics for GPU targets. +//! +//! Intrinsics in this module are intended for use on GPU targets. +//! They can be target specific but in general GPU targets are similar. + +#![unstable(feature = "gpu_intrinsics", issue = "none")] + +/// Returns a pointer to the HSA kernel dispatch packet. +/// +/// A `gpu-kernel` on amdgpu is always launched through a kernel dispatch packet. +/// The dispatch packet contains the workgroup size, launch size and other data. +/// The content is defined by the [HSA Platform System Architecture Specification], +/// which is implemented e.g. in AMD's [hsa.h]. +/// The intrinsic returns a unit pointer so that rustc does not need to know the packet struct. +/// The pointer is valid for the whole lifetime of the program. +/// +/// [HSA Platform System Architecture Specification]: https://hsafoundation.com/wp-content/uploads/2021/02/HSA-SysArch-1.2.pdf +/// [hsa.h]: https://github.com/ROCm/rocm-systems/blob/rocm-7.1.0/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa.h#L2959 +#[rustc_nounwind] +#[rustc_intrinsic] +#[cfg(target_arch = "amdgpu")] +#[must_use = "returns a pointer that does nothing unless used"] +pub fn amdgpu_dispatch_ptr() -> *const (); diff --git a/library/core/src/intrinsics/mod.rs b/library/core/src/intrinsics/mod.rs index 20f34036b25c9..8d112b4c5d187 100644 --- a/library/core/src/intrinsics/mod.rs +++ b/library/core/src/intrinsics/mod.rs @@ -60,6 +60,7 @@ use crate::{mem, ptr}; mod bounds; pub mod fallback; +pub mod gpu; pub mod mir; pub mod simd; diff --git a/library/coretests/tests/floats/f128.rs b/library/coretests/tests/floats/f128.rs deleted file mode 100644 index 8e4f0c9899e1c..0000000000000 --- a/library/coretests/tests/floats/f128.rs +++ /dev/null @@ -1,50 +0,0 @@ -// FIXME(f16_f128): only tested on platforms that have symbols and aren't buggy -#![cfg(target_has_reliable_f128)] - -use super::assert_biteq; - -// Note these tolerances make sense around zero, but not for more extreme exponents. - -/// Default tolerances. Works for values that should be near precise but not exact. Roughly -/// the precision carried by `100 * 100`. -#[allow(unused)] -const TOL: f128 = 1e-12; - -/// For operations that are near exact, usually not involving math of different -/// signs. -#[allow(unused)] -const TOL_PRECISE: f128 = 1e-28; - -// FIXME(f16_f128,miri): many of these have to be disabled since miri does not yet support -// the intrinsics. - -#[test] -fn test_from() { - assert_biteq!(f128::from(false), 0.0); - assert_biteq!(f128::from(true), 1.0); - assert_biteq!(f128::from(u8::MIN), 0.0); - assert_biteq!(f128::from(42_u8), 42.0); - assert_biteq!(f128::from(u8::MAX), 255.0); - assert_biteq!(f128::from(i8::MIN), -128.0); - assert_biteq!(f128::from(42_i8), 42.0); - assert_biteq!(f128::from(i8::MAX), 127.0); - assert_biteq!(f128::from(u16::MIN), 0.0); - assert_biteq!(f128::from(42_u16), 42.0); - assert_biteq!(f128::from(u16::MAX), 65535.0); - assert_biteq!(f128::from(i16::MIN), -32768.0); - assert_biteq!(f128::from(42_i16), 42.0); - assert_biteq!(f128::from(i16::MAX), 32767.0); - assert_biteq!(f128::from(u32::MIN), 0.0); - assert_biteq!(f128::from(42_u32), 42.0); - assert_biteq!(f128::from(u32::MAX), 4294967295.0); - assert_biteq!(f128::from(i32::MIN), -2147483648.0); - assert_biteq!(f128::from(42_i32), 42.0); - assert_biteq!(f128::from(i32::MAX), 2147483647.0); - // FIXME(f16_f128): Uncomment these tests once the From<{u64,i64}> impls are added. - // assert_eq!(f128::from(u64::MIN), 0.0); - // assert_eq!(f128::from(42_u64), 42.0); - // assert_eq!(f128::from(u64::MAX), 18446744073709551615.0); - // assert_eq!(f128::from(i64::MIN), -9223372036854775808.0); - // assert_eq!(f128::from(42_i64), 42.0); - // assert_eq!(f128::from(i64::MAX), 9223372036854775807.0); -} diff --git a/library/coretests/tests/floats/f16.rs b/library/coretests/tests/floats/f16.rs deleted file mode 100644 index 3cff4259de54f..0000000000000 --- a/library/coretests/tests/floats/f16.rs +++ /dev/null @@ -1,35 +0,0 @@ -// FIXME(f16_f128): only tested on platforms that have symbols and aren't buggy -#![cfg(target_has_reliable_f16)] - -use super::assert_biteq; - -/// Tolerance for results on the order of 10.0e-2 -#[allow(unused)] -const TOL_N2: f16 = 0.0001; - -/// Tolerance for results on the order of 10.0e+0 -#[allow(unused)] -const TOL_0: f16 = 0.01; - -/// Tolerance for results on the order of 10.0e+2 -#[allow(unused)] -const TOL_P2: f16 = 0.5; - -/// Tolerance for results on the order of 10.0e+4 -#[allow(unused)] -const TOL_P4: f16 = 10.0; - -// FIXME(f16_f128,miri): many of these have to be disabled since miri does not yet support -// the intrinsics. - -#[test] -fn test_from() { - assert_biteq!(f16::from(false), 0.0); - assert_biteq!(f16::from(true), 1.0); - assert_biteq!(f16::from(u8::MIN), 0.0); - assert_biteq!(f16::from(42_u8), 42.0); - assert_biteq!(f16::from(u8::MAX), 255.0); - assert_biteq!(f16::from(i8::MIN), -128.0); - assert_biteq!(f16::from(42_i8), 42.0); - assert_biteq!(f16::from(i8::MAX), 127.0); -} diff --git a/library/coretests/tests/floats/mod.rs b/library/coretests/tests/floats/mod.rs index 63d5b8fb2c6e9..87e21b21f310d 100644 --- a/library/coretests/tests/floats/mod.rs +++ b/library/coretests/tests/floats/mod.rs @@ -375,9 +375,6 @@ macro_rules! float_test { }; } -mod f128; -mod f16; - float_test! { name: num, attrs: { @@ -1582,3 +1579,78 @@ float_test! { assert_biteq!((flt(-3.2)).mul_add(2.4, neg_inf), neg_inf); } } + +float_test! { + name: from, + attrs: { + f16: #[cfg(any(miri, target_has_reliable_f16))], + f128: #[cfg(any(miri, target_has_reliable_f128))], + }, + test { + assert_biteq!(Float::from(false), Float::ZERO); + assert_biteq!(Float::from(true), Float::ONE); + + assert_biteq!(Float::from(u8::MIN), Float::ZERO); + assert_biteq!(Float::from(42_u8), 42.0); + assert_biteq!(Float::from(u8::MAX), 255.0); + + assert_biteq!(Float::from(i8::MIN), -128.0); + assert_biteq!(Float::from(42_i8), 42.0); + assert_biteq!(Float::from(i8::MAX), 127.0); + } +} + +float_test! { + name: from_u16_i16, + attrs: { + f16: #[cfg(false)], + const f16: #[cfg(false)], + f128: #[cfg(any(miri, target_has_reliable_f128))], + }, + test { + assert_biteq!(Float::from(u16::MIN), Float::ZERO); + assert_biteq!(Float::from(42_u16), 42.0); + assert_biteq!(Float::from(u16::MAX), 65535.0); + assert_biteq!(Float::from(i16::MIN), -32768.0); + assert_biteq!(Float::from(42_i16), 42.0); + assert_biteq!(Float::from(i16::MAX), 32767.0); + } +} + +float_test! { + name: from_u32_i32, + attrs: { + f16: #[cfg(false)], + const f16: #[cfg(false)], + f32: #[cfg(false)], + const f32: #[cfg(false)], + f128: #[cfg(any(miri, target_has_reliable_f128))], + }, + test { + assert_biteq!(Float::from(u32::MIN), Float::ZERO); + assert_biteq!(Float::from(42_u32), 42.0); + assert_biteq!(Float::from(u32::MAX), 4294967295.0); + assert_biteq!(Float::from(i32::MIN), -2147483648.0); + assert_biteq!(Float::from(42_i32), 42.0); + assert_biteq!(Float::from(i32::MAX), 2147483647.0); + } +} + +// FIXME(f16_f128): Uncomment and adapt these tests once the From<{u64,i64}> impls are added. +// float_test! { +// name: from_u64_i64, +// attrs: { +// f16: #[cfg(false)], +// f32: #[cfg(false)], +// f64: #[cfg(false)], +// f128: #[cfg(any(miri, target_has_reliable_f128))], +// }, +// test { +// assert_biteq!(Float::from(u64::MIN), Float::ZERO); +// assert_biteq!(Float::from(42_u64), 42.0); +// assert_biteq!(Float::from(u64::MAX), 18446744073709551615.0); +// assert_biteq!(Float::from(i64::MIN), -9223372036854775808.0); +// assert_biteq!(Float::from(42_i64), 42.0); +// assert_biteq!(Float::from(i64::MAX), 9223372036854775807.0); +// } +// } diff --git a/tests/codegen-llvm/amdgpu-dispatch-ptr.rs b/tests/codegen-llvm/amdgpu-dispatch-ptr.rs new file mode 100644 index 0000000000000..00bde96c3d596 --- /dev/null +++ b/tests/codegen-llvm/amdgpu-dispatch-ptr.rs @@ -0,0 +1,27 @@ +// Tests the amdgpu_dispatch_ptr intrinsic. + +//@ compile-flags: --crate-type=rlib --target amdgcn-amd-amdhsa -Ctarget-cpu=gfx900 +//@ needs-llvm-components: amdgpu +//@ add-minicore +#![feature(intrinsics, no_core, rustc_attrs)] +#![no_core] + +extern crate minicore; + +pub struct DispatchPacket { + pub header: u16, + pub setup: u16, + pub workgroup_size_x: u16, // and more +} + +#[rustc_intrinsic] +#[rustc_nounwind] +fn amdgpu_dispatch_ptr() -> *const (); + +// CHECK-LABEL: @get_dispatch_data +// CHECK: %[[ORIG_PTR:[^ ]+]] = {{(tail )?}}call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +// CHECK-NEXT: %[[PTR:[^ ]+]] = addrspacecast ptr addrspace(4) %[[ORIG_PTR]] to ptr +#[unsafe(no_mangle)] +pub fn get_dispatch_data() -> &'static DispatchPacket { + unsafe { &*(amdgpu_dispatch_ptr() as *const _) } +} diff --git a/tests/mir-opt/building/match/sort_candidates.constant_eq.SimplifyCfg-initial.after.mir b/tests/mir-opt/building/match/sort_candidates.constant_eq.SimplifyCfg-initial.after.mir index b8f54fef0fafc..4d13d087586e7 100644 --- a/tests/mir-opt/building/match/sort_candidates.constant_eq.SimplifyCfg-initial.after.mir +++ b/tests/mir-opt/building/match/sort_candidates.constant_eq.SimplifyCfg-initial.after.mir @@ -7,11 +7,14 @@ fn constant_eq(_1: &str, _2: bool) -> u32 { let mut _3: (&str, bool); let mut _4: &str; let mut _5: bool; - let mut _6: &&str; - let mut _7: &bool; - let mut _8: bool; - let mut _9: bool; + let mut _6: &str; + let mut _7: &&str; + let mut _8: &bool; + let mut _9: &str; let mut _10: bool; + let mut _11: &str; + let mut _12: bool; + let mut _13: bool; bb0: { StorageLive(_3); @@ -23,7 +26,8 @@ fn constant_eq(_1: &str, _2: bool) -> u32 { StorageDead(_5); StorageDead(_4); PlaceMention(_3); - _9 = ::eq(copy (_3.0: &str), const "a") -> [return: bb9, unwind: bb19]; + _11 = &(*(_3.0: &str)); + _12 = ::eq(copy _11, const "a") -> [return: bb9, unwind: bb19]; } bb1: { @@ -43,7 +47,8 @@ fn constant_eq(_1: &str, _2: bool) -> u32 { } bb5: { - _8 = ::eq(copy (_3.0: &str), const "b") -> [return: bb8, unwind: bb19]; + _9 = &(*(_3.0: &str)); + _10 = ::eq(copy _9, const "b") -> [return: bb8, unwind: bb19]; } bb6: { @@ -55,11 +60,11 @@ fn constant_eq(_1: &str, _2: bool) -> u32 { } bb8: { - switchInt(move _8) -> [0: bb1, otherwise: bb6]; + switchInt(move _10) -> [0: bb1, otherwise: bb6]; } bb9: { - switchInt(move _9) -> [0: bb5, otherwise: bb2]; + switchInt(move _12) -> [0: bb5, otherwise: bb2]; } bb10: { @@ -87,23 +92,25 @@ fn constant_eq(_1: &str, _2: bool) -> u32 { } bb15: { - _6 = &fake shallow (_3.0: &str); - _7 = &fake shallow (_3.1: bool); - StorageLive(_10); - _10 = const true; - switchInt(move _10) -> [0: bb17, otherwise: bb16]; + _6 = &fake shallow (*(_3.0: &str)); + _7 = &fake shallow (_3.0: &str); + _8 = &fake shallow (_3.1: bool); + StorageLive(_13); + _13 = const true; + switchInt(move _13) -> [0: bb17, otherwise: bb16]; } bb16: { - StorageDead(_10); + StorageDead(_13); FakeRead(ForMatchGuard, _6); FakeRead(ForMatchGuard, _7); + FakeRead(ForMatchGuard, _8); _0 = const 1_u32; goto -> bb18; } bb17: { - StorageDead(_10); + StorageDead(_13); falseEdge -> [real: bb3, imaginary: bb5]; } diff --git a/tests/ui/thir-print/str-patterns.rs b/tests/ui/thir-print/str-patterns.rs new file mode 100644 index 0000000000000..4ad782f63a505 --- /dev/null +++ b/tests/ui/thir-print/str-patterns.rs @@ -0,0 +1,17 @@ +#![crate_type = "rlib"] +//@ edition: 2024 +//@ compile-flags: -Zunpretty=thir-flat +//@ check-pass + +// Snapshot test capturing the THIR pattern structure produced by +// string-literal and string-constant patterns. + +pub fn hello_world(x: &str) { + match x { + "hello" => {} + CONSTANT => {} + _ => {} + } +} + +const CONSTANT: &str = "constant"; diff --git a/tests/ui/thir-print/str-patterns.stdout b/tests/ui/thir-print/str-patterns.stdout new file mode 100644 index 0000000000000..bc2549f813c06 --- /dev/null +++ b/tests/ui/thir-print/str-patterns.stdout @@ -0,0 +1,344 @@ +DefId(0:3 ~ str_patterns[fc71]::hello_world): +Thir { + body_type: Fn( + fn(&'{erased} str), + ), + arms: [ + Arm { + pattern: Pat { + ty: &'{erased} str, + span: $DIR/str-patterns.rs:11:9: 11:16 (#0), + extra: None, + kind: Deref { + subpattern: Pat { + ty: str, + span: $DIR/str-patterns.rs:11:9: 11:16 (#0), + extra: None, + kind: Constant { + value: Value { + ty: str, + valtree: Branch( + [ + 104_u8, + 101_u8, + 108_u8, + 108_u8, + 111_u8, + ], + ), + }, + }, + }, + }, + }, + guard: None, + body: e3, + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).9), + ), + scope: Node(9), + span: $DIR/str-patterns.rs:11:9: 11:22 (#0), + }, + Arm { + pattern: Pat { + ty: &'{erased} str, + span: $DIR/str-patterns.rs:12:9: 12:17 (#0), + extra: Some( + PatExtra { + expanded_const: Some( + DefId(0:4 ~ str_patterns[fc71]::CONSTANT), + ), + ascriptions: [], + }, + ), + kind: Deref { + subpattern: Pat { + ty: str, + span: $DIR/str-patterns.rs:12:9: 12:17 (#0), + extra: None, + kind: Constant { + value: Value { + ty: str, + valtree: Branch( + [ + 99_u8, + 111_u8, + 110_u8, + 115_u8, + 116_u8, + 97_u8, + 110_u8, + 116_u8, + ], + ), + }, + }, + }, + }, + }, + guard: None, + body: e5, + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).15), + ), + scope: Node(15), + span: $DIR/str-patterns.rs:12:9: 12:23 (#0), + }, + Arm { + pattern: Pat { + ty: &'{erased} str, + span: $DIR/str-patterns.rs:13:9: 13:10 (#0), + extra: None, + kind: Wild, + }, + guard: None, + body: e7, + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).19), + ), + scope: Node(19), + span: $DIR/str-patterns.rs:13:9: 13:16 (#0), + }, + ], + blocks: [ + Block { + targeted_by_break: false, + region_scope: Node(11), + span: $DIR/str-patterns.rs:11:20: 11:22 (#0), + stmts: [], + expr: None, + safety_mode: Safe, + }, + Block { + targeted_by_break: false, + region_scope: Node(17), + span: $DIR/str-patterns.rs:12:21: 12:23 (#0), + stmts: [], + expr: None, + safety_mode: Safe, + }, + Block { + targeted_by_break: false, + region_scope: Node(21), + span: $DIR/str-patterns.rs:13:14: 13:16 (#0), + stmts: [], + expr: None, + safety_mode: Safe, + }, + Block { + targeted_by_break: false, + region_scope: Node(3), + span: $DIR/str-patterns.rs:9:29: 15:2 (#0), + stmts: [], + expr: Some( + e9, + ), + safety_mode: Safe, + }, + ], + exprs: [ + Expr { + kind: VarRef { + id: LocalVarId( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).2), + ), + }, + ty: &'{erased} str, + temp_scope_id: 5, + span: $DIR/str-patterns.rs:10:11: 10:12 (#0), + }, + Expr { + kind: Scope { + region_scope: Node(5), + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).5), + ), + value: e0, + }, + ty: &'{erased} str, + temp_scope_id: 5, + span: $DIR/str-patterns.rs:10:11: 10:12 (#0), + }, + Expr { + kind: Block { + block: b0, + }, + ty: (), + temp_scope_id: 10, + span: $DIR/str-patterns.rs:11:20: 11:22 (#0), + }, + Expr { + kind: Scope { + region_scope: Node(10), + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).10), + ), + value: e2, + }, + ty: (), + temp_scope_id: 10, + span: $DIR/str-patterns.rs:11:20: 11:22 (#0), + }, + Expr { + kind: Block { + block: b1, + }, + ty: (), + temp_scope_id: 16, + span: $DIR/str-patterns.rs:12:21: 12:23 (#0), + }, + Expr { + kind: Scope { + region_scope: Node(16), + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).16), + ), + value: e4, + }, + ty: (), + temp_scope_id: 16, + span: $DIR/str-patterns.rs:12:21: 12:23 (#0), + }, + Expr { + kind: Block { + block: b2, + }, + ty: (), + temp_scope_id: 20, + span: $DIR/str-patterns.rs:13:14: 13:16 (#0), + }, + Expr { + kind: Scope { + region_scope: Node(20), + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).20), + ), + value: e6, + }, + ty: (), + temp_scope_id: 20, + span: $DIR/str-patterns.rs:13:14: 13:16 (#0), + }, + Expr { + kind: Match { + scrutinee: e1, + arms: [ + a0, + a1, + a2, + ], + match_source: Normal, + }, + ty: (), + temp_scope_id: 4, + span: $DIR/str-patterns.rs:10:5: 14:6 (#0), + }, + Expr { + kind: Scope { + region_scope: Node(4), + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).4), + ), + value: e8, + }, + ty: (), + temp_scope_id: 4, + span: $DIR/str-patterns.rs:10:5: 14:6 (#0), + }, + Expr { + kind: Block { + block: b3, + }, + ty: (), + temp_scope_id: 22, + span: $DIR/str-patterns.rs:9:29: 15:2 (#0), + }, + Expr { + kind: Scope { + region_scope: Node(22), + lint_level: Explicit( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).22), + ), + value: e10, + }, + ty: (), + temp_scope_id: 22, + span: $DIR/str-patterns.rs:9:29: 15:2 (#0), + }, + ], + stmts: [], + params: [ + Param { + pat: Some( + Pat { + ty: &'{erased} str, + span: $DIR/str-patterns.rs:9:20: 9:21 (#0), + extra: None, + kind: Binding { + name: "x", + mode: BindingMode( + No, + Not, + ), + var: LocalVarId( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).2), + ), + ty: &'{erased} str, + subpattern: None, + is_primary: true, + is_shorthand: false, + }, + }, + ), + ty: &'{erased} str, + ty_span: Some( + $DIR/str-patterns.rs:9:23: 9:27 (#0), + ), + self_kind: None, + hir_id: Some( + HirId(DefId(0:3 ~ str_patterns[fc71]::hello_world).1), + ), + }, + ], +} + +DefId(0:4 ~ str_patterns[fc71]::CONSTANT): +Thir { + body_type: Const( + &'{erased} str, + ), + arms: [], + blocks: [], + exprs: [ + Expr { + kind: Literal { + lit: Spanned { + node: Str( + "constant", + Cooked, + ), + span: $DIR/str-patterns.rs:17:24: 17:34 (#0), + }, + neg: false, + }, + ty: &'{erased} str, + temp_scope_id: 5, + span: $DIR/str-patterns.rs:17:24: 17:34 (#0), + }, + Expr { + kind: Scope { + region_scope: Node(5), + lint_level: Explicit( + HirId(DefId(0:4 ~ str_patterns[fc71]::CONSTANT).5), + ), + value: e0, + }, + ty: &'{erased} str, + temp_scope_id: 5, + span: $DIR/str-patterns.rs:17:24: 17:34 (#0), + }, + ], + stmts: [], + params: [], +} +