Merge commit '482e8540a1b757ed7bccc2041c5400f051fdb01e' into subtree-update_cg_gcc_2025-08-04
This commit is contained in:
@@ -31,7 +31,7 @@ pub fn run() -> Result<(), String> {
|
||||
Some("clones/abi-cafe".as_ref()),
|
||||
true,
|
||||
)
|
||||
.map_err(|err| (format!("Git clone failed with message: {err:?}!")))?;
|
||||
.map_err(|err| format!("Git clone failed with message: {err:?}!"))?;
|
||||
// Configure abi-cafe to use the exact same rustc version we use - this is crucial.
|
||||
// Otherwise, the concept of ABI compatibility becomes meanignless.
|
||||
std::fs::copy("rust-toolchain", "clones/abi-cafe/rust-toolchain")
|
||||
|
||||
@@ -43,18 +43,18 @@ pub fn run() -> Result<(), String> {
|
||||
"--start" => {
|
||||
start =
|
||||
str::parse(&args.next().ok_or_else(|| "Fuzz start not provided!".to_string())?)
|
||||
.map_err(|err| (format!("Fuzz start not a number {err:?}!")))?;
|
||||
.map_err(|err| format!("Fuzz start not a number {err:?}!"))?;
|
||||
}
|
||||
"--count" => {
|
||||
count =
|
||||
str::parse(&args.next().ok_or_else(|| "Fuzz count not provided!".to_string())?)
|
||||
.map_err(|err| (format!("Fuzz count not a number {err:?}!")))?;
|
||||
.map_err(|err| format!("Fuzz count not a number {err:?}!"))?;
|
||||
}
|
||||
"-j" | "--jobs" => {
|
||||
threads = str::parse(
|
||||
&args.next().ok_or_else(|| "Fuzz thread count not provided!".to_string())?,
|
||||
)
|
||||
.map_err(|err| (format!("Fuzz thread count not a number {err:?}!")))?;
|
||||
.map_err(|err| format!("Fuzz thread count not a number {err:?}!"))?;
|
||||
}
|
||||
_ => return Err(format!("Unknown option {arg}")),
|
||||
}
|
||||
@@ -66,7 +66,7 @@ pub fn run() -> Result<(), String> {
|
||||
Some("clones/rustlantis".as_ref()),
|
||||
true,
|
||||
)
|
||||
.map_err(|err| (format!("Git clone failed with message: {err:?}!")))?;
|
||||
.map_err(|err| format!("Git clone failed with message: {err:?}!"))?;
|
||||
|
||||
// Ensure that we are on the newest rustlantis commit.
|
||||
let cmd: &[&dyn AsRef<OsStr>] = &[&"git", &"pull", &"origin"];
|
||||
|
||||
@@ -1,29 +1,28 @@
|
||||
From b8f3eed3053c9333b5dfbeaeb2a6a65a4b3156df Mon Sep 17 00:00:00 2001
|
||||
From: Antoni Boucher <bouanto@zoho.com>
|
||||
Date: Tue, 29 Aug 2023 13:06:34 -0400
|
||||
From 190e26c9274b3c93a9ee3516b395590e6bd9213b Mon Sep 17 00:00:00 2001
|
||||
From: None <none@example.com>
|
||||
Date: Sun, 3 Aug 2025 19:54:56 -0400
|
||||
Subject: [PATCH] Patch 0001-Add-stdarch-Cargo.toml-for-testing.patch
|
||||
|
||||
---
|
||||
library/stdarch/Cargo.toml | 23 +++++++++++++++++++++++
|
||||
1 file changed, 23 insertions(+)
|
||||
library/stdarch/Cargo.toml | 20 ++++++++++++++++++++
|
||||
1 file changed, 20 insertions(+)
|
||||
create mode 100644 library/stdarch/Cargo.toml
|
||||
|
||||
diff --git a/library/stdarch/Cargo.toml b/library/stdarch/Cargo.toml
|
||||
new file mode 100644
|
||||
index 0000000..4c63700
|
||||
index 0000000..bd6725c
|
||||
--- /dev/null
|
||||
+++ b/library/stdarch/Cargo.toml
|
||||
@@ -0,0 +1,21 @@
|
||||
@@ -0,0 +1,20 @@
|
||||
+[workspace]
|
||||
+resolver = "1"
|
||||
+members = [
|
||||
+ "crates/core_arch",
|
||||
+ "crates/std_detect",
|
||||
+ "crates/stdarch-gen-arm",
|
||||
+ "crates/*",
|
||||
+ #"examples/"
|
||||
+]
|
||||
+exclude = [
|
||||
+ "crates/wasm-assert-instr-tests"
|
||||
+ "crates/wasm-assert-instr-tests",
|
||||
+ "rust_programs",
|
||||
+]
|
||||
+
|
||||
+[profile.release]
|
||||
@@ -36,5 +35,5 @@ index 0000000..4c63700
|
||||
+opt-level = 3
|
||||
+incremental = true
|
||||
--
|
||||
2.42.0
|
||||
2.50.1
|
||||
|
||||
|
||||
@@ -1,3 +1,3 @@
|
||||
[toolchain]
|
||||
channel = "nightly-2025-07-04"
|
||||
channel = "nightly-2025-08-03"
|
||||
components = ["rust-src", "rustc-dev", "llvm-tools-preview"]
|
||||
|
||||
@@ -99,11 +99,15 @@ fn create_const_value_function(
|
||||
let func = context.new_function(None, FunctionType::Exported, output, &[], name, false);
|
||||
|
||||
#[cfg(feature = "master")]
|
||||
func.add_attribute(FnAttribute::Visibility(symbol_visibility_to_gcc(
|
||||
tcx.sess.default_visibility(),
|
||||
)));
|
||||
{
|
||||
func.add_attribute(FnAttribute::Visibility(symbol_visibility_to_gcc(
|
||||
tcx.sess.default_visibility(),
|
||||
)));
|
||||
|
||||
func.add_attribute(FnAttribute::AlwaysInline);
|
||||
// FIXME(antoyo): cg_llvm sets AlwaysInline, but AlwaysInline is different in GCC and using
|
||||
// it here will causes linking errors when using LTO.
|
||||
func.add_attribute(FnAttribute::Inline);
|
||||
}
|
||||
|
||||
if tcx.sess.must_emit_unwind_tables() {
|
||||
// TODO(antoyo): emit unwind tables.
|
||||
|
||||
@@ -540,9 +540,15 @@ impl<'a, 'gcc, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'gcc, 'tcx> {
|
||||
|
||||
fn ret(&mut self, mut value: RValue<'gcc>) {
|
||||
let expected_return_type = self.current_func().get_return_type();
|
||||
if !expected_return_type.is_compatible_with(value.get_type()) {
|
||||
// NOTE: due to opaque pointers now being used, we need to cast here.
|
||||
value = self.context.new_cast(self.location, value, expected_return_type);
|
||||
let value_type = value.get_type();
|
||||
if !expected_return_type.is_compatible_with(value_type) {
|
||||
// NOTE: due to opaque pointers now being used, we need to (bit)cast here.
|
||||
if self.is_native_int_type(value_type) && self.is_native_int_type(expected_return_type)
|
||||
{
|
||||
value = self.context.new_cast(self.location, value, expected_return_type);
|
||||
} else {
|
||||
value = self.context.new_bitcast(self.location, value, expected_return_type);
|
||||
}
|
||||
}
|
||||
self.llbb().end_with_return(self.location, value);
|
||||
}
|
||||
@@ -1279,11 +1285,19 @@ impl<'a, 'gcc, 'tcx> BuilderMethods<'a, 'tcx> for Builder<'a, 'gcc, 'tcx> {
|
||||
|
||||
fn intcast(
|
||||
&mut self,
|
||||
value: RValue<'gcc>,
|
||||
mut value: RValue<'gcc>,
|
||||
dest_typ: Type<'gcc>,
|
||||
_is_signed: bool,
|
||||
is_signed: bool,
|
||||
) -> RValue<'gcc> {
|
||||
// NOTE: is_signed is for value, not dest_typ.
|
||||
let value_type = value.get_type();
|
||||
if is_signed && !value_type.is_signed(self.cx) {
|
||||
let signed_type = value_type.to_signed(self.cx);
|
||||
value = self.gcc_int_cast(value, signed_type);
|
||||
} else if !is_signed && value_type.is_signed(self.cx) {
|
||||
let unsigned_type = value_type.to_unsigned(self.cx);
|
||||
value = self.gcc_int_cast(value, unsigned_type);
|
||||
}
|
||||
|
||||
self.gcc_int_cast(value, dest_typ)
|
||||
}
|
||||
|
||||
|
||||
@@ -4,12 +4,15 @@
|
||||
|
||||
// cSpell:words cmpti divti modti mulodi muloti udivti umodti
|
||||
|
||||
use gccjit::{BinaryOp, ComparisonOp, FunctionType, Location, RValue, ToRValue, Type, UnaryOp};
|
||||
use gccjit::{
|
||||
BinaryOp, CType, ComparisonOp, FunctionType, Location, RValue, ToRValue, Type, UnaryOp,
|
||||
};
|
||||
use rustc_abi::{CanonAbi, Endian, ExternAbi};
|
||||
use rustc_codegen_ssa::common::{IntPredicate, TypeKind};
|
||||
use rustc_codegen_ssa::traits::{BackendTypes, BaseTypeCodegenMethods, BuilderMethods, OverflowOp};
|
||||
use rustc_middle::ty::{self, Ty};
|
||||
use rustc_target::callconv::{ArgAbi, ArgAttributes, FnAbi, PassMode};
|
||||
use rustc_type_ir::{Interner, TyKind};
|
||||
|
||||
use crate::builder::{Builder, ToGccComp};
|
||||
use crate::common::{SignType, TypeReflection};
|
||||
@@ -167,9 +170,9 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
if a_type.is_vector() {
|
||||
// Vector types need to be bitcast.
|
||||
// TODO(antoyo): perhaps use __builtin_convertvector for vector casting.
|
||||
b = self.context.new_bitcast(self.location, b, a.get_type());
|
||||
b = self.context.new_bitcast(self.location, b, a_type);
|
||||
} else {
|
||||
b = self.context.new_cast(self.location, b, a.get_type());
|
||||
b = self.context.new_cast(self.location, b, a_type);
|
||||
}
|
||||
}
|
||||
self.context.new_binary_op(self.location, operation, a_type, a, b)
|
||||
@@ -216,13 +219,22 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
operation_name: &str,
|
||||
signed: bool,
|
||||
a: RValue<'gcc>,
|
||||
b: RValue<'gcc>,
|
||||
mut b: RValue<'gcc>,
|
||||
) -> RValue<'gcc> {
|
||||
let a_type = a.get_type();
|
||||
let b_type = b.get_type();
|
||||
if (self.is_native_int_type_or_bool(a_type) && self.is_native_int_type_or_bool(b_type))
|
||||
|| (a_type.is_vector() && b_type.is_vector())
|
||||
{
|
||||
if !a_type.is_compatible_with(b_type) {
|
||||
if a_type.is_vector() {
|
||||
// Vector types need to be bitcast.
|
||||
// TODO(antoyo): perhaps use __builtin_convertvector for vector casting.
|
||||
b = self.context.new_bitcast(self.location, b, a_type);
|
||||
} else {
|
||||
b = self.context.new_cast(self.location, b, a_type);
|
||||
}
|
||||
}
|
||||
self.context.new_binary_op(self.location, operation, a_type, a, b)
|
||||
} else {
|
||||
debug_assert!(a_type.dyncast_array().is_some());
|
||||
@@ -351,6 +363,11 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
// TODO(antoyo): is it correct to use rhs type instead of the parameter typ?
|
||||
.new_local(self.location, rhs.get_type(), "binopResult")
|
||||
.get_address(self.location);
|
||||
let new_type = type_kind_to_gcc_type(new_kind);
|
||||
let new_type = self.context.new_c_type(new_type);
|
||||
let lhs = self.context.new_cast(self.location, lhs, new_type);
|
||||
let rhs = self.context.new_cast(self.location, rhs, new_type);
|
||||
let res = self.context.new_cast(self.location, res, new_type.make_pointer());
|
||||
let overflow = self.overflow_call(intrinsic, &[lhs, rhs, res], None);
|
||||
(res.dereference(self.location).to_rvalue(), overflow)
|
||||
}
|
||||
@@ -477,11 +494,27 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
let lhs_low = self.context.new_cast(self.location, self.low(lhs), unsigned_type);
|
||||
let rhs_low = self.context.new_cast(self.location, self.low(rhs), unsigned_type);
|
||||
|
||||
let mut lhs_high = self.high(lhs);
|
||||
let mut rhs_high = self.high(rhs);
|
||||
|
||||
match op {
|
||||
IntPredicate::IntUGT
|
||||
| IntPredicate::IntUGE
|
||||
| IntPredicate::IntULT
|
||||
| IntPredicate::IntULE => {
|
||||
lhs_high = self.context.new_cast(self.location, lhs_high, unsigned_type);
|
||||
rhs_high = self.context.new_cast(self.location, rhs_high, unsigned_type);
|
||||
}
|
||||
// TODO(antoyo): we probably need to handle signed comparison for unsigned
|
||||
// integers.
|
||||
_ => (),
|
||||
}
|
||||
|
||||
let condition = self.context.new_comparison(
|
||||
self.location,
|
||||
ComparisonOp::LessThan,
|
||||
self.high(lhs),
|
||||
self.high(rhs),
|
||||
lhs_high,
|
||||
rhs_high,
|
||||
);
|
||||
self.llbb().end_with_conditional(self.location, condition, block1, block2);
|
||||
|
||||
@@ -495,8 +528,8 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
let condition = self.context.new_comparison(
|
||||
self.location,
|
||||
ComparisonOp::GreaterThan,
|
||||
self.high(lhs),
|
||||
self.high(rhs),
|
||||
lhs_high,
|
||||
rhs_high,
|
||||
);
|
||||
block2.end_with_conditional(self.location, condition, block3, block4);
|
||||
|
||||
@@ -620,7 +653,7 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
}
|
||||
}
|
||||
|
||||
pub fn gcc_xor(&self, a: RValue<'gcc>, b: RValue<'gcc>) -> RValue<'gcc> {
|
||||
pub fn gcc_xor(&self, a: RValue<'gcc>, mut b: RValue<'gcc>) -> RValue<'gcc> {
|
||||
let a_type = a.get_type();
|
||||
let b_type = b.get_type();
|
||||
if a_type.is_vector() && b_type.is_vector() {
|
||||
@@ -628,6 +661,9 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
a ^ b
|
||||
} else if self.is_native_int_type_or_bool(a_type) && self.is_native_int_type_or_bool(b_type)
|
||||
{
|
||||
if !a_type.is_compatible_with(b_type) {
|
||||
b = self.context.new_cast(self.location, b, a_type);
|
||||
}
|
||||
a ^ b
|
||||
} else {
|
||||
self.concat_low_high_rvalues(
|
||||
@@ -1042,3 +1078,25 @@ impl<'gcc, 'tcx> CodegenCx<'gcc, 'tcx> {
|
||||
self.context.new_array_constructor(None, typ, &values)
|
||||
}
|
||||
}
|
||||
|
||||
fn type_kind_to_gcc_type<I: Interner>(kind: TyKind<I>) -> CType {
|
||||
use rustc_middle::ty::IntTy::*;
|
||||
use rustc_middle::ty::UintTy::*;
|
||||
use rustc_middle::ty::{Int, Uint};
|
||||
|
||||
match kind {
|
||||
Int(I8) => CType::Int8t,
|
||||
Int(I16) => CType::Int16t,
|
||||
Int(I32) => CType::Int32t,
|
||||
Int(I64) => CType::Int64t,
|
||||
Int(I128) => CType::Int128t,
|
||||
|
||||
Uint(U8) => CType::UInt8t,
|
||||
Uint(U16) => CType::UInt16t,
|
||||
Uint(U32) => CType::UInt32t,
|
||||
Uint(U64) => CType::UInt64t,
|
||||
Uint(U128) => CType::UInt128t,
|
||||
|
||||
_ => unimplemented!("Kind: {:?}", kind),
|
||||
}
|
||||
}
|
||||
|
||||
@@ -95,8 +95,11 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
|
||||
"cubema" => "__builtin_amdgcn_cubema",
|
||||
"cubesc" => "__builtin_amdgcn_cubesc",
|
||||
"cubetc" => "__builtin_amdgcn_cubetc",
|
||||
"cvt.f16.bf8" => "__builtin_amdgcn_cvt_f16_bf8",
|
||||
"cvt.f16.fp8" => "__builtin_amdgcn_cvt_f16_fp8",
|
||||
"cvt.f32.bf8" => "__builtin_amdgcn_cvt_f32_bf8",
|
||||
"cvt.f32.fp8" => "__builtin_amdgcn_cvt_f32_fp8",
|
||||
"cvt.f32.fp8.e5m3" => "__builtin_amdgcn_cvt_f32_fp8_e5m3",
|
||||
"cvt.off.f32.i4" => "__builtin_amdgcn_cvt_off_f32_i4",
|
||||
"cvt.pk.bf8.f32" => "__builtin_amdgcn_cvt_pk_bf8_f32",
|
||||
"cvt.pk.f16.bf8" => "__builtin_amdgcn_cvt_pk_f16_bf8",
|
||||
@@ -181,6 +184,12 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
|
||||
"dot4.f32.fp8.bf8" => "__builtin_amdgcn_dot4_f32_fp8_bf8",
|
||||
"dot4.f32.fp8.fp8" => "__builtin_amdgcn_dot4_f32_fp8_fp8",
|
||||
"ds.add.gs.reg.rtn" => "__builtin_amdgcn_ds_add_gs_reg_rtn",
|
||||
"ds.atomic.async.barrier.arrive.b64" => {
|
||||
"__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64"
|
||||
}
|
||||
"ds.atomic.barrier.arrive.rtn.b64" => {
|
||||
"__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64"
|
||||
}
|
||||
"ds.bpermute" => "__builtin_amdgcn_ds_bpermute",
|
||||
"ds.bpermute.fi.b32" => "__builtin_amdgcn_ds_bpermute_fi_b32",
|
||||
"ds.gws.barrier" => "__builtin_amdgcn_ds_gws_barrier",
|
||||
@@ -198,8 +207,32 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
|
||||
"fdot2.f16.f16" => "__builtin_amdgcn_fdot2_f16_f16",
|
||||
"fdot2.f32.bf16" => "__builtin_amdgcn_fdot2_f32_bf16",
|
||||
"fdot2c.f32.bf16" => "__builtin_amdgcn_fdot2c_f32_bf16",
|
||||
"flat.prefetch" => "__builtin_amdgcn_flat_prefetch",
|
||||
"fmul.legacy" => "__builtin_amdgcn_fmul_legacy",
|
||||
"global.load.async.to.lds.b128" => {
|
||||
"__builtin_amdgcn_global_load_async_to_lds_b128"
|
||||
}
|
||||
"global.load.async.to.lds.b32" => {
|
||||
"__builtin_amdgcn_global_load_async_to_lds_b32"
|
||||
}
|
||||
"global.load.async.to.lds.b64" => {
|
||||
"__builtin_amdgcn_global_load_async_to_lds_b64"
|
||||
}
|
||||
"global.load.async.to.lds.b8" => "__builtin_amdgcn_global_load_async_to_lds_b8",
|
||||
"global.load.lds" => "__builtin_amdgcn_global_load_lds",
|
||||
"global.prefetch" => "__builtin_amdgcn_global_prefetch",
|
||||
"global.store.async.from.lds.b128" => {
|
||||
"__builtin_amdgcn_global_store_async_from_lds_b128"
|
||||
}
|
||||
"global.store.async.from.lds.b32" => {
|
||||
"__builtin_amdgcn_global_store_async_from_lds_b32"
|
||||
}
|
||||
"global.store.async.from.lds.b64" => {
|
||||
"__builtin_amdgcn_global_store_async_from_lds_b64"
|
||||
}
|
||||
"global.store.async.from.lds.b8" => {
|
||||
"__builtin_amdgcn_global_store_async_from_lds_b8"
|
||||
}
|
||||
"groupstaticsize" => "__builtin_amdgcn_groupstaticsize",
|
||||
"iglp.opt" => "__builtin_amdgcn_iglp_opt",
|
||||
"implicit.buffer.ptr" => "__builtin_amdgcn_implicit_buffer_ptr",
|
||||
@@ -291,6 +324,7 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
|
||||
"s.incperflevel" => "__builtin_amdgcn_s_incperflevel",
|
||||
"s.memrealtime" => "__builtin_amdgcn_s_memrealtime",
|
||||
"s.memtime" => "__builtin_amdgcn_s_memtime",
|
||||
"s.monitor.sleep" => "__builtin_amdgcn_s_monitor_sleep",
|
||||
"s.sendmsg" => "__builtin_amdgcn_s_sendmsg",
|
||||
"s.sendmsghalt" => "__builtin_amdgcn_s_sendmsghalt",
|
||||
"s.setprio" => "__builtin_amdgcn_s_setprio",
|
||||
@@ -300,11 +334,15 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
|
||||
"s.sleep.var" => "__builtin_amdgcn_s_sleep_var",
|
||||
"s.ttracedata" => "__builtin_amdgcn_s_ttracedata",
|
||||
"s.ttracedata.imm" => "__builtin_amdgcn_s_ttracedata_imm",
|
||||
"s.wait.asynccnt" => "__builtin_amdgcn_s_wait_asynccnt",
|
||||
"s.wait.event.export.ready" => "__builtin_amdgcn_s_wait_event_export_ready",
|
||||
"s.wait.tensorcnt" => "__builtin_amdgcn_s_wait_tensorcnt",
|
||||
"s.waitcnt" => "__builtin_amdgcn_s_waitcnt",
|
||||
"sad.hi.u8" => "__builtin_amdgcn_sad_hi_u8",
|
||||
"sad.u16" => "__builtin_amdgcn_sad_u16",
|
||||
"sad.u8" => "__builtin_amdgcn_sad_u8",
|
||||
"sat.pk4.i4.i8" => "__builtin_amdgcn_sat_pk4_i4_i8",
|
||||
"sat.pk4.u4.u8" => "__builtin_amdgcn_sat_pk4_u4_u8",
|
||||
"sched.barrier" => "__builtin_amdgcn_sched_barrier",
|
||||
"sched.group.barrier" => "__builtin_amdgcn_sched_group_barrier",
|
||||
"sdot2" => "__builtin_amdgcn_sdot2",
|
||||
@@ -346,8 +384,13 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
|
||||
"smfmac.i32.16x16x64.i8" => "__builtin_amdgcn_smfmac_i32_16x16x64_i8",
|
||||
"smfmac.i32.32x32x32.i8" => "__builtin_amdgcn_smfmac_i32_32x32x32_i8",
|
||||
"smfmac.i32.32x32x64.i8" => "__builtin_amdgcn_smfmac_i32_32x32x64_i8",
|
||||
"struct.ptr.buffer.load.lds" => "__builtin_amdgcn_struct_ptr_buffer_load_lds",
|
||||
"sudot4" => "__builtin_amdgcn_sudot4",
|
||||
"sudot8" => "__builtin_amdgcn_sudot8",
|
||||
"tensor.load.to.lds" => "__builtin_amdgcn_tensor_load_to_lds",
|
||||
"tensor.load.to.lds.d2" => "__builtin_amdgcn_tensor_load_to_lds_d2",
|
||||
"tensor.store.from.lds" => "__builtin_amdgcn_tensor_store_from_lds",
|
||||
"tensor.store.from.lds.d2" => "__builtin_amdgcn_tensor_store_from_lds_d2",
|
||||
"udot2" => "__builtin_amdgcn_udot2",
|
||||
"udot4" => "__builtin_amdgcn_udot4",
|
||||
"udot8" => "__builtin_amdgcn_udot8",
|
||||
@@ -6326,6 +6369,23 @@ fn map_arch_intrinsic(full_name: &str) -> &'static str {
|
||||
}
|
||||
s390(name, full_name)
|
||||
}
|
||||
"spv" => {
|
||||
#[allow(non_snake_case)]
|
||||
fn spv(name: &str, full_name: &str) -> &'static str {
|
||||
match name {
|
||||
// spv
|
||||
"num.subgroups" => "__builtin_spirv_num_subgroups",
|
||||
"subgroup.id" => "__builtin_spirv_subgroup_id",
|
||||
"subgroup.local.invocation.id" => {
|
||||
"__builtin_spirv_subgroup_local_invocation_id"
|
||||
}
|
||||
"subgroup.max.size" => "__builtin_spirv_subgroup_max_size",
|
||||
"subgroup.size" => "__builtin_spirv_subgroup_size",
|
||||
_ => unimplemented!("***** unsupported LLVM intrinsic {full_name}"),
|
||||
}
|
||||
}
|
||||
spv(name, full_name)
|
||||
}
|
||||
"ve" => {
|
||||
#[allow(non_snake_case)]
|
||||
fn ve(name: &str, full_name: &str) -> &'static str {
|
||||
|
||||
@@ -925,10 +925,17 @@ impl<'a, 'gcc, 'tcx> Builder<'a, 'gcc, 'tcx> {
|
||||
// TODO(antoyo): use width?
|
||||
let arg_type = arg.get_type();
|
||||
let result_type = self.u32_type;
|
||||
let arg = if arg_type.is_signed(self.cx) {
|
||||
let new_type = arg_type.to_unsigned(self.cx);
|
||||
self.gcc_int_cast(arg, new_type)
|
||||
} else {
|
||||
arg
|
||||
};
|
||||
let arg_type = arg.get_type();
|
||||
let count_leading_zeroes =
|
||||
// TODO(antoyo): write a new function Type::is_compatible_with(&Type) and use it here
|
||||
// instead of using is_uint().
|
||||
if arg_type.is_uint(self.cx) {
|
||||
if arg_type.is_uchar(self.cx) || arg_type.is_ushort(self.cx) || arg_type.is_uint(self.cx) {
|
||||
"__builtin_clz"
|
||||
}
|
||||
else if arg_type.is_ulong(self.cx) {
|
||||
|
||||
@@ -206,6 +206,28 @@ pub fn generic_simd_intrinsic<'a, 'gcc, 'tcx>(
|
||||
);
|
||||
}
|
||||
|
||||
#[cfg(feature = "master")]
|
||||
if name == sym::simd_funnel_shl {
|
||||
return Ok(simd_funnel_shift(
|
||||
bx,
|
||||
args[0].immediate(),
|
||||
args[1].immediate(),
|
||||
args[2].immediate(),
|
||||
true,
|
||||
));
|
||||
}
|
||||
|
||||
#[cfg(feature = "master")]
|
||||
if name == sym::simd_funnel_shr {
|
||||
return Ok(simd_funnel_shift(
|
||||
bx,
|
||||
args[0].immediate(),
|
||||
args[1].immediate(),
|
||||
args[2].immediate(),
|
||||
false,
|
||||
));
|
||||
}
|
||||
|
||||
if name == sym::simd_bswap {
|
||||
return Ok(simd_bswap(bx, args[0].immediate()));
|
||||
}
|
||||
@@ -1434,3 +1456,62 @@ pub fn generic_simd_intrinsic<'a, 'gcc, 'tcx>(
|
||||
|
||||
unimplemented!("simd {}", name);
|
||||
}
|
||||
|
||||
#[cfg(feature = "master")]
|
||||
fn simd_funnel_shift<'a, 'gcc, 'tcx>(
|
||||
bx: &mut Builder<'a, 'gcc, 'tcx>,
|
||||
a: RValue<'gcc>,
|
||||
b: RValue<'gcc>,
|
||||
shift: RValue<'gcc>,
|
||||
shift_left: bool,
|
||||
) -> RValue<'gcc> {
|
||||
use crate::common::SignType;
|
||||
|
||||
let a_type = a.get_type();
|
||||
let vector_type = a_type.unqualified().dyncast_vector().expect("vector type");
|
||||
let num_units = vector_type.get_num_units();
|
||||
let elem_type = vector_type.get_element_type();
|
||||
|
||||
let (new_int_type, int_shift_val, int_mask) = if elem_type.is_compatible_with(bx.u8_type)
|
||||
|| elem_type.is_compatible_with(bx.i8_type)
|
||||
{
|
||||
(bx.u16_type, 8, u8::MAX as u64)
|
||||
} else if elem_type.is_compatible_with(bx.u16_type) || elem_type.is_compatible_with(bx.i16_type)
|
||||
{
|
||||
(bx.u32_type, 16, u16::MAX as u64)
|
||||
} else if elem_type.is_compatible_with(bx.u32_type) || elem_type.is_compatible_with(bx.i32_type)
|
||||
{
|
||||
(bx.u64_type, 32, u32::MAX as u64)
|
||||
} else if elem_type.is_compatible_with(bx.u64_type) || elem_type.is_compatible_with(bx.i64_type)
|
||||
{
|
||||
(bx.u128_type, 64, u64::MAX)
|
||||
} else {
|
||||
unimplemented!("funnel shift on {:?}", elem_type);
|
||||
};
|
||||
|
||||
let int_mask = bx.context.new_rvalue_from_long(new_int_type, int_mask as i64);
|
||||
let int_shift_val = bx.context.new_rvalue_from_int(new_int_type, int_shift_val);
|
||||
let mut elements = vec![];
|
||||
let unsigned_type = elem_type.to_unsigned(bx);
|
||||
for i in 0..num_units {
|
||||
let index = bx.context.new_rvalue_from_int(bx.int_type, i as i32);
|
||||
let a_val = bx.context.new_vector_access(None, a, index).to_rvalue();
|
||||
let a_val = bx.context.new_bitcast(None, a_val, unsigned_type);
|
||||
// TODO: we probably need to use gcc_int_cast instead.
|
||||
let a_val = bx.gcc_int_cast(a_val, new_int_type);
|
||||
let b_val = bx.context.new_vector_access(None, b, index).to_rvalue();
|
||||
let b_val = bx.context.new_bitcast(None, b_val, unsigned_type);
|
||||
let b_val = bx.gcc_int_cast(b_val, new_int_type);
|
||||
let shift_val = bx.context.new_vector_access(None, shift, index).to_rvalue();
|
||||
let shift_val = bx.gcc_int_cast(shift_val, new_int_type);
|
||||
let mut val = a_val << int_shift_val | b_val;
|
||||
if shift_left {
|
||||
val = (val << shift_val) >> int_shift_val;
|
||||
} else {
|
||||
val = (val >> shift_val) & int_mask;
|
||||
}
|
||||
let val = bx.gcc_int_cast(val, elem_type);
|
||||
elements.push(val);
|
||||
}
|
||||
bx.context.new_rvalue_from_vector(None, a_type, &elements)
|
||||
}
|
||||
|
||||
@@ -50,6 +50,7 @@ extern crate rustc_session;
|
||||
extern crate rustc_span;
|
||||
extern crate rustc_symbol_mangling;
|
||||
extern crate rustc_target;
|
||||
extern crate rustc_type_ir;
|
||||
|
||||
// This prevents duplicating functions and statics that are already part of the host rustc process.
|
||||
#[allow(unused_extern_crates)]
|
||||
@@ -362,9 +363,9 @@ impl WriteBackendMethods for GccCodegenBackend {
|
||||
_exported_symbols_for_lto: &[String],
|
||||
each_linked_rlib_for_lto: &[PathBuf],
|
||||
modules: Vec<FatLtoInput<Self>>,
|
||||
diff_fncs: Vec<AutoDiffItem>,
|
||||
diff_functions: Vec<AutoDiffItem>,
|
||||
) -> Result<ModuleCodegen<Self::Module>, FatalError> {
|
||||
if !diff_fncs.is_empty() {
|
||||
if !diff_functions.is_empty() {
|
||||
unimplemented!();
|
||||
}
|
||||
|
||||
|
||||
@@ -28,6 +28,6 @@ tests/ui/macros/macro-comma-behavior-rpass.rs
|
||||
tests/ui/macros/rfc-2011-nicer-assert-messages/assert-with-custom-errors-does-not-create-unnecessary-code.rs
|
||||
tests/ui/macros/rfc-2011-nicer-assert-messages/feature-gate-generic_assert.rs
|
||||
tests/ui/macros/stringify.rs
|
||||
tests/ui/reexport-test-harness-main.rs
|
||||
tests/ui/rfcs/rfc-1937-termination-trait/termination-trait-in-test.rs
|
||||
tests/ui/binding/fn-arg-incomplete-pattern-drop-order.rs
|
||||
tests/ui/lto/debuginfo-lto-alloc.rs
|
||||
|
||||
@@ -6,7 +6,6 @@ tests/run-make/doctests-keep-binaries/
|
||||
tests/run-make/doctests-runtool/
|
||||
tests/run-make/emit-shared-files/
|
||||
tests/run-make/exit-code/
|
||||
tests/run-make/issue-22131/
|
||||
tests/run-make/issue-64153/
|
||||
tests/run-make/llvm-ident/
|
||||
tests/run-make/native-link-modifier-bundle/
|
||||
|
||||
@@ -10,11 +10,10 @@ tests/ui/iterators/iter-sum-overflow-overflow-checks.rs
|
||||
tests/ui/mir/mir_drop_order.rs
|
||||
tests/ui/mir/mir_let_chains_drop_order.rs
|
||||
tests/ui/mir/mir_match_guard_let_chains_drop_order.rs
|
||||
tests/ui/oom_unwind.rs
|
||||
tests/ui/panics/oom-panic-unwind.rs
|
||||
tests/ui/panic-runtime/abort-link-to-unwinding-crates.rs
|
||||
tests/ui/panic-runtime/abort.rs
|
||||
tests/ui/panic-runtime/link-to-abort.rs
|
||||
tests/ui/unwind-no-uwtable.rs
|
||||
tests/ui/parser/unclosed-delimiter-in-dep.rs
|
||||
tests/ui/consts/missing_span_in_backtrace.rs
|
||||
tests/ui/drop/dynamic-drop.rs
|
||||
@@ -82,3 +81,8 @@ tests/ui/coroutine/panic-drops.rs
|
||||
tests/ui/coroutine/panic-safe.rs
|
||||
tests/ui/process/nofile-limit.rs
|
||||
tests/ui/simd/intrinsic/generic-arithmetic-pass.rs
|
||||
tests/ui/linking/no-gc-encapsulation-symbols.rs
|
||||
tests/ui/panics/unwind-force-no-unwind-tables.rs
|
||||
tests/ui/attributes/fn-align-dyn.rs
|
||||
tests/ui/linkage-attr/raw-dylib/elf/glibc-x86_64.rs
|
||||
tests/ui/explicit-tail-calls/recursion-etc.rs
|
||||
|
||||
@@ -8,6 +8,7 @@ clzll
|
||||
cmse
|
||||
codegened
|
||||
csky
|
||||
ctfe
|
||||
ctlz
|
||||
ctpop
|
||||
cttz
|
||||
@@ -25,6 +26,7 @@ fwrapv
|
||||
gimple
|
||||
hrtb
|
||||
immediates
|
||||
interner
|
||||
liblto
|
||||
llbb
|
||||
llcx
|
||||
@@ -47,6 +49,7 @@ mavx
|
||||
mcmodel
|
||||
minimumf
|
||||
minnumf
|
||||
miri
|
||||
monomorphization
|
||||
monomorphizations
|
||||
monomorphized
|
||||
|
||||
Reference in New Issue
Block a user