refactor: 移除 scheme 接口

Signed-off-by: YdrMaster <ydrml@hotmail.com>

Author: YdrMaster

.github/workflows/build.yml

@@ -52,6 +52,7 @@ jobs:
         run:
           cargo clippy
           --all-features
+          --all-targets
           --message-format=json | clippy-sarif | tee rust-clippy-results.sarif | sarif-fmt
         continue-on-error: true
 

operators/src/add/args.rs

@@ -1,6 +1,6 @@
 use crate::{
     get_static, rank_mismatch, shape_mismatch, shape_not_support, utils::type_distinct, ConstPtr,
-    Hardware, MutPtr, SchemeError, TensorLayout,
+    Hardware, LaunchError, MutPtr, TensorLayout,
 };
 use digit_layout::DigitLayout;
 use itertools::izip;
@@ -40,7 +40,7 @@ impl<H: Hardware> Args<H> {
 pub(super) struct Scheme(DigitLayout, Box<[isize]>);
 
 impl Scheme {
-    pub fn new<H: Hardware>(args: &Args<H>) -> Result<Self, SchemeError> {
+    pub fn new<H: Hardware>(args: &Args<H>) -> Result<Self, LaunchError> {
         let Args {
             c_layout: c,
             a_layout: a,

operators/src/add/common_cpu/mod.rs

@@ -1,5 +1,5 @@
 use super::{args::Scheme, Add, Args};
-use crate::{common_cpu::Cpu, ByteOf, LaunchError, QueueAlloc, SchemeError};
+use crate::{common_cpu::Cpu, ByteOf, LaunchError, QueueAlloc};
 use digit_layout::types as ty;
 use half::f16;
 use rayon::iter::{IntoParallelIterator, ParallelIterator};
@@ -17,14 +17,6 @@ impl crate::Operator for Operator {
     fn new(_node: &Self::TopoNode) -> Self {
         Self
     }
-    #[inline]
-    fn scheme(
-        &mut self,
-        _args: &Self::Args,
-        _max_workspace_size: usize,
-    ) -> Result<usize, SchemeError> {
-        Ok(0)
-    }
 
     fn launch<QA>(
         &self,

operators/src/add/cuda/mod.rs

@@ -2,8 +2,8 @@ use super::{args::Scheme, Add, Args};
 use crate::{
     cuda::{dt_name, Gpu, Handle, ModuleBox},
     shape_not_support, strides_not_support,
-    utils::{gcd, type_distinct},
-    ByteOf, LaunchError, QueueAlloc, SchemeDiversity, SchemeError,
+    utils::gcd,
+    ByteOf, LaunchError, QueueAlloc, SchemeDiversity,
 };
 use digit_layout::DigitLayout;
 use lru::LruCache;
@@ -32,20 +32,6 @@ impl crate::Operator for Operator {
         }
     }
 
-    #[inline]
-    fn scheme(
-        &mut self,
-        args: &Self::Args,
-        _max_workspace_size: usize,
-    ) -> Result<usize, SchemeError> {
-        let dt = type_distinct(&[args.c_layout.dt(), args.a_layout.dt(), args.b_layout.dt()])?;
-        self.schemes
-            .lock()
-            .unwrap()
-            .get_or_insert(dt, || compile(&self.handle, dt));
-        Ok(0)
-    }
-
     fn launch<QA>(
         &self,
         args: &Self::Args,
@@ -60,20 +46,20 @@ impl crate::Operator for Operator {
         let count = scheme.count();
 
         let &[1] = scheme.idx_strides() else {
-            return Err(shape_not_support("").into());
+            return Err(shape_not_support(""));
         };
         let &[sc] = scheme.c_strides() else {
-            return Err(shape_not_support("").into());
+            return Err(shape_not_support(""));
         };
         let &[sa] = scheme.a_strides() else {
-            return Err(shape_not_support("").into());
+            return Err(shape_not_support(""));
         };
         let &[sb] = scheme.b_strides() else {
-            return Err(shape_not_support("").into());
+            return Err(shape_not_support(""));
         };
         let unit = dt.nbytes() as isize;
         if sc != unit || sa != unit || sb != unit {
-            return Err(strides_not_support("").into());
+            return Err(strides_not_support(""));
         }
 
         let block_dims = gcd(count, self.max_threads_block);
@@ -124,25 +110,12 @@ extern "C" __global__ void add(
 #[cfg(test)]
 mod test {
     use super::{Args, Gpu, Operator};
-    use crate::{dyn_, Hardware, Operator as _, TensorLayout};
+    use crate::{Hardware, Operator as _, TensorLayout};
     use digit_layout::{
         types::{F16, F64},
         DigitLayout,
     };
-    use std::ptr::null;
 
-    fn dyn_args<H: Hardware>(dt: DigitLayout) -> Args<H> {
-        use std::ptr::null_mut;
-        let layout = TensorLayout::new_dyn(dt, &[dyn_(); 2], &[dyn_(); 2]);
-        Args {
-            c_layout: layout.clone(),
-            c_base: null_mut(),
-            a_layout: layout.clone(),
-            a_base: null(),
-            b_layout: layout.clone(),
-            b_base: null(),
-        }
-    }
     fn args<H: Hardware>(
         dt: DigitLayout,
         n: usize,
@@ -178,10 +151,8 @@ mod test {
             return;
         };
 
-        let mut cpu_op = RefOp::new(&Cpu);
-        let mut gpu_op = Operator::new(&gpu);
-        cpu_op.scheme(&dyn_args(F64), 0).unwrap();
-        gpu_op.scheme(&dyn_args(F16), 0).unwrap();
+        let cpu_op = RefOp::new(&Cpu);
+        let gpu_op = Operator::new(&gpu);
 
         let n = 1;
         let d = 768;

operators/src/add/infini/mod.rs

@@ -4,10 +4,6 @@
 pub mod common_cpu;
 #[cfg(use_cuda)]
 pub mod cuda;
-#[cfg(use_infini)]
-pub mod infini;
-#[cfg(use_cl)]
-pub mod opencl;
 
 mod args;
 pub use args::Args;

operators/src/add/mod.rs

@@ -1,7 +1,7 @@
-use crate::{
+use crate::{
     type_not_support,
     utils::{dim_distinct, rank_error, type_distinct},
-    ConstPtr, Hardware, MaybeDyn, MutPtr, SchemeError, TensorLayout,
+    ConstPtr, Hardware, LaunchError, MaybeDyn, MutPtr, TensorLayout,
 };
 use digit_layout::{DigitLayout, LayoutContent::Unsigned};
 use std::ptr::{null, null_mut};
@@ -44,7 +44,7 @@ pub(super) struct Meta {
 }
 
 impl<H: Hardware> Args<H> {
-    pub(super) fn meta(&self) -> Result<Meta, SchemeError> {
+    pub(super) fn meta(&self) -> Result<Meta, LaunchError> {
         let Self {
             dst_layout: dst,
             src_layout: src,

operators/src/add/opencl/mod.rs

@@ -1,5 +1,5 @@
 use super::{args::Meta, AddRows, Args};
-use crate::{common_cpu::Cpu, get_static, ByteOf, LaunchError, QueueAlloc, SchemeError, Unsigned};
+use crate::{common_cpu::Cpu, get_static, ByteOf, LaunchError, QueueAlloc, Unsigned};
 use digit_layout::types as ty;
 use half::f16;
 use rayon::iter::{IntoParallelIterator, ParallelIterator};
@@ -18,14 +18,6 @@ impl crate::Operator for Operator {
         Self
     }
 
-    fn scheme(
-        &mut self,
-        _args: &Self::Args,
-        _max_workspace_size: usize,
-    ) -> Result<usize, SchemeError> {
-        Ok(0)
-    }
-
     fn launch<QA>(
         &self,
         args: &Self::Args,

operators/src/add_rows/args.rs

@@ -4,7 +4,7 @@ use crate::{
     cuda::{dt_name, Gpu, Handle, ModuleBox},
     get_static, strides_not_support,
     utils::gcd,
-    ByteOf, LaunchError, QueueAlloc, SchemeDiversity, SchemeError,
+    ByteOf, LaunchError, QueueAlloc, SchemeDiversity,
 };
 use digit_layout::DigitLayout;
 use lru::LruCache;
@@ -34,22 +34,6 @@ impl crate::Operator for Operator {
         }
     }
 
-    #[inline]
-    fn scheme(
-        &mut self,
-        args: &Self::Args,
-        _max_workspace_size: usize,
-    ) -> Result<usize, SchemeError> {
-        let Meta { dt, .. } = args.meta()?;
-
-        let key = SchemeKey { dt };
-        self.schemes
-            .lock()
-            .unwrap()
-            .try_get_or_insert(key, || Scheme::new(&self.handle, key))?;
-        Ok(0)
-    }
-
     fn launch<QA>(
         &self,
         args: &Self::Args,
@@ -87,7 +71,7 @@ impl crate::Operator for Operator {
         let unit_dst = dst_layout.dt().nbytes() as isize;
         let unit_idx = idx_layout.dt().nbytes() as isize;
         if nsd != unit_dst || nss != unit_dst || msi != unit_idx {
-            return Err(strides_not_support("").into());
+            return Err(strides_not_support(""));
         };
         fn cast(strides: &[isize], size: usize) -> Vec<isize> {
             strides.iter().map(|x| x / size as isize).collect()
@@ -134,7 +118,7 @@ struct SchemeKey {
 }
 
 impl Scheme {
-    pub fn new(handle: &Arc<Handle>, SchemeKey { dt }: SchemeKey) -> Result<Self, SchemeError> {
+    pub fn new(handle: &Arc<Handle>, SchemeKey { dt }: SchemeKey) -> Result<Self, LaunchError> {
         let device = handle.device();
         let cc = device.compute_capability();
         let type_name = dt_name(dt);
@@ -168,26 +152,14 @@ extern "C" __global__ void {name}(
 #[cfg(test)]
 mod test {
     use super::{Args, Gpu, Operator};
-    use crate::{cuda::cast_load, dyn_, Hardware, Operator as _, TensorLayout};
+    use crate::{cuda::cast_load, Hardware, Operator as _, TensorLayout};
     use cuda::memcpy_d2h;
     use digit_layout::{
         types::{F16, F64, U32},
         DigitLayout,
     };
     use half::f16;
-    use std::ptr::null;
 
-    fn dyn_args<H: Hardware>(dt: DigitLayout) -> Args<H> {
-        use std::ptr::null_mut;
-        Args {
-            dst_layout: TensorLayout::new_dyn(dt, &[dyn_(); 3], &[dyn_(); 3]),
-            dst_base: null_mut(),
-            src_layout: TensorLayout::new_dyn(dt, &[dyn_(); 2], &[dyn_(); 2]),
-            src_base: null(),
-            idx_layout: TensorLayout::new_dyn(U32, &[dyn_(); 2], &[dyn_(); 2]),
-            idx_base: null(),
-        }
-    }
     fn args<H: Hardware>(
         dt: DigitLayout,
         b: usize,
@@ -220,10 +192,8 @@ mod test {
             return;
         };
 
-        let mut cpu_op = RefOp::new(&Cpu);
-        let mut gpu_op = Operator::new(&gpu);
-        cpu_op.scheme(&dyn_args(F64), 0).unwrap();
-        gpu_op.scheme(&dyn_args(F16), 0).unwrap();
+        let cpu_op = RefOp::new(&Cpu);
+        let gpu_op = Operator::new(&gpu);
 
         let b = 1;
         let m = 10;