diff --git a/StgSharp.Common/Mathematics/Linear/Matrix.cs b/StgSharp.Common/Mathematics/Linear/Matrix.cs
index 62868022..c75d2bf3 100644
--- a/StgSharp.Common/Mathematics/Linear/Matrix.cs
+++ b/StgSharp.Common/Mathematics/Linear/Matrix.cs
@@ -25,38 +25,21 @@
 //     
 // -----------------------------------------------------------------------
 // -----------------------------------------------------------------------
-using StgSharp.HighPerformance;
 using System;
+using System.Collections.Generic;
+using System.Linq;
 using System.Numerics;
 using System.Runtime.CompilerServices;
-using System.Runtime.InteropServices;
+using System.Text;
+using System.Threading.Tasks;
 
 using hlsfAllocator = global::StgSharp.HighPerformance.Memory.HybridLayerSegregatedFitAllocator;
 
-// using hlsfHandle = global::StgSharp.HighPerformance.Memory.HybridLayerSegregatedFitAllocationHandle;
-
 namespace StgSharp.Mathematics.Numeric
 {
-    public abstract unsafe partial class Matrix<T> where T: unmanaged, INumber<T>
+    public static class Matrix
     {
 
-        private protected readonly int t_count = 16 / sizeof(T);
-
-        private protected readonly int t_size = sizeof(T);
-
-        private protected Matrix() { }
-
-        public abstract ref T this[int column, int row]
-        {
-            get;
-        }
-
-        public abstract unsafe Matrix<T> this[ReadOnlySpan<MatrixIndex> column, ReadOnlySpan<MatrixIndex> row]
-        {
-            get;
-            set;
-        }
-
         public static hlsfAllocator? DefaultAllocator
         {
             get
@@ -67,192 +50,29 @@ namespace StgSharp.Mathematics.Numeric
             set;
         } = null;
 
-        /// <summary>
-        ///   Length of the matrix in columns.
-        /// </summary>
-        public int ColumnLength { get; init; }
-
-        /// <summary>
-        ///   Length of the matrix in rows.
-        /// </summary>
-        public int RowLength { get; init; }
-
         public static MatrixAllocation DefaultAllocation { get; set; } = MatrixAllocation.GC;
 
-        internal abstract MatrixKernel<T>* Buffer { get; }
-
-        protected internal int KernelColumnLength { get; init; }
-
-        protected internal int KernelRowLength { get; init; }
-
-        public static Matrix<T> FromDefault(int columnLength, int rowLength)
+        public static Matrix<T> FromDefault<T>(int columnLength, int rowLength) where T : unmanaged, INumber<T>
         {
             return DefaultAllocation switch
             {
-                MatrixAllocation.GC => FromGC(columnLength, rowLength),
-                MatrixAllocation.HLSF => FromHlsf(columnLength, rowLength, DefaultAllocator!),
+                MatrixAllocation.GC => FromGC<T>(columnLength, rowLength),
+                MatrixAllocation.HLSF => FromHlsf<T>(columnLength, rowLength, DefaultAllocator!),
                 _ => throw new NotSupportedException("This allocation is not supported."),
             };
         }
 
-        public static Matrix<T> FromGC(int columnLength, int rowLength)
+        public static Matrix<T> FromGC<T>(int columnLength, int rowLength) where T : unmanaged, INumber<T>
         {
             return new GCMatrix<T>(columnLength, rowLength);
         }
 
         [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public static Matrix<T> FromHlsf(int columnLength, int rowLength, hlsfAllocator allocator)
+        public static Matrix<T> FromHlsf<T>(int columnLength, int rowLength, hlsfAllocator allocator)
+            where T : unmanaged, INumber<T>
         {
             return new HlsfMatrix<T>(columnLength, rowLength, allocator);
         }
 
-        [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public abstract MatrixSegmentEnumeration<T> GetColumnEnumeration();
-
-        public abstract ref MatrixKernel<T> GetMatrixKernel(int column, int row);
-
-        public abstract MatrixSegmentEnumeration<T> GetRowEnumeration();
-
-        [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public abstract MatrixSegmentEnumeration<T> GetSequentialEnumeration();
-
-        [Obsolete("This method does not contain boarder check, so be careful on size or use indexer instead.", false)]
-        public abstract ref T UnsafeIndex(int column, int row);
-
-        private protected void AssertIndex(int column, int row)
-        {
-            if (column < 0 || column >= ColumnLength) {
-                throw new ArgumentOutOfRangeException(nameof(column));
-            }
-
-            if (row < 0 || row >= RowLength) {
-                throw new ArgumentOutOfRangeException(nameof(row));
-            }
-        }
-
-    }
-
-    public readonly unsafe ref struct MatrixSegmentEnumeration<T>(
-                                      MatrixKernel<T>* source,
-                                      int primStride,
-                                      int primSize,
-                                      int secondaryStride,
-                                      int secondarySize,
-                                      int primOffset = 0,
-                                      int secondaryOffset = 0)
-        where T: unmanaged, INumber<T>
-    {
-
-        private readonly MatrixKernel<T>* _source = source;
-        internal readonly int PrimOffset = primOffset;
-
-        internal readonly int PrimSize = primSize;
-        internal readonly int PrimStride = primStride;
-        internal readonly int SecondaryOffset = secondaryOffset;
-        internal readonly int SecondarySize = secondarySize;
-        internal readonly int SecondaryStride = secondaryStride;
-
-        internal readonly MatrixKernel<T>* Source
-        {
-            [MethodImpl(MethodImplOptions.AggressiveInlining)]
-            get => _source;
-        }
-
-        public Enumerator GetEnumerator()
-        {
-            return new Enumerator(in this);
-        }
-
-        public ref struct Enumerator(in MatrixSegmentEnumeration<T> enumeration)
-        {
-
-            private readonly MatrixKernel<T>* _source = enumeration._source;
-            private int _indexCount = -1;
-            private readonly int _stride = enumeration.PrimStride;
-            private readonly long _bound = (long)enumeration.PrimSize * enumeration.PrimStride;
-            private long _index = (long)enumeration.PrimOffset * enumeration.PrimStride;
-
-            private readonly MatrixSegmentEnumeration<T> _enumeration;
-
-            public readonly MatrixSegment<T> Current
-            {
-                [MethodImpl(MethodImplOptions.AggressiveInlining)]
-                get => new(
-                    _source + _index,
-                    _enumeration.SecondarySize,
-                    _enumeration.SecondaryStride,
-                    _enumeration.PrimOffset);
-            }
-
-            public bool MoveNext()
-            {
-                _indexCount++;
-                _index += _stride;
-                if (_index >= _bound) {
-                    _index = 0;
-                }
-                return _indexCount >= _enumeration.PrimSize;
-            }
-
-        }
-
-    }
-
-    public readonly unsafe ref struct MatrixSegment<T>(MatrixKernel<T>* source, int size, int stride, int offset)
-        where T: unmanaged, INumber<T>
-    {
-
-        private readonly MatrixKernel<T>* _source = source;
-        private readonly int _offset = offset;
-        private readonly int _secondarySize = size;
-        private readonly int _size = size;
-        private readonly int _stride = stride;
-
-        public ref MatrixKernel<T> this[int index]
-        {
-            get
-            {
-                if (index < 0 || index >= _size) {
-                    throw new ArgumentOutOfRangeException(nameof(index));
-                }
-                return ref _source[index * _stride];
-            }
-        }
-
-        [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public Enumerator GetEnumerator()
-        {
-            return new Enumerator(in this);
-        }
-
-        public ref struct Enumerator(in MatrixSegment<T> segment)
-        {
-
-            private readonly int _end = segment._size;
-
-            private int _indexCount = -1;
-            private readonly long _bound = (long)segment._size * segment._stride;
-            private long _index = (long)segment._offset * segment._stride;
-            private readonly MatrixSegment<T> _segment = segment;
-
-            public readonly MatrixKernel<T>* Current
-            {
-                [MethodImpl(MethodImplOptions.AggressiveInlining)]
-                get => _segment._source + _index;
-            }
-
-            public bool MoveNext()
-            {
-                _index += _segment._stride;
-                _indexCount++;
-                if (_index >= _bound) {
-                    _index = 0;
-                }
-
-                return _indexCount < _end;
-            }
-
-        }
-
     }
 }
diff --git a/StgSharp.Common/Mathematics/Linear/Matrix.GC.cs b/StgSharp.Common/Mathematics/Linear/Matrix_T.GC.cs
similarity index 91%
rename from StgSharp.Common/Mathematics/Linear/Matrix.GC.cs
rename to StgSharp.Common/Mathematics/Linear/Matrix_T.GC.cs
index 52daa4c7..57eb48e2 100644
--- a/StgSharp.Common/Mathematics/Linear/Matrix.GC.cs
+++ b/StgSharp.Common/Mathematics/Linear/Matrix_T.GC.cs
@@ -1,6 +1,6 @@
 //-----------------------------------------------------------------------
 // -----------------------------------------------------------------------
-// file="Matrix.GC"
+// file="Matrix_T.GC"
 // Project: StgSharp
 // AuthorGroup: Nitload
 // Copyright (c) Nitload. All rights reserved.
@@ -94,10 +94,10 @@ namespace StgSharp.Mathematics.Numeric
 
         #region enumeration
 
-        [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public override unsafe MatrixSegmentEnumeration<T> GetColumnEnumeration()
+        public override unsafe MatrixSegmentEnumeration<T> ColumnEnumeration
         {
-            return new MatrixSegmentEnumeration<T>(_buffer, KernelColumnLength, KernelColumnLength * KernelRowLength, 1,
+            [MethodImpl(MethodImplOptions.AggressiveInlining)]
+            get => new MatrixSegmentEnumeration<T>(_buffer, KernelColumnLength, KernelColumnLength * KernelRowLength, 1,
                                                    KernelColumnLength);
         }
 
@@ -106,10 +106,10 @@ namespace StgSharp.Mathematics.Numeric
             return new MatrixSegmentEnumeration<T>(_buffer, 1, KernelColumnLength, KernelColumnLength, KernelRowLength);
         }
 
-        [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public override unsafe MatrixSegmentEnumeration<T> GetSequentialEnumeration()
+        public override unsafe MatrixSegmentEnumeration<T> SequentialEnumeration
         {
-            return new MatrixSegmentEnumeration<T>(_buffer, KernelRowLength * KernelColumnLength,
+            [MethodImpl(MethodImplOptions.AggressiveInlining)]
+            get => new MatrixSegmentEnumeration<T>(_buffer, KernelRowLength * KernelColumnLength,
                                                    KernelColumnLength * KernelRowLength, 1,
                                                    KernelRowLength * KernelColumnLength);
         }
diff --git a/StgSharp.Common/Mathematics/Linear/Matrix.HLSF.cs b/StgSharp.Common/Mathematics/Linear/Matrix_T.HLSF.cs
similarity index 91%
rename from StgSharp.Common/Mathematics/Linear/Matrix.HLSF.cs
rename to StgSharp.Common/Mathematics/Linear/Matrix_T.HLSF.cs
index 8689a906..e1f1f530 100644
--- a/StgSharp.Common/Mathematics/Linear/Matrix.HLSF.cs
+++ b/StgSharp.Common/Mathematics/Linear/Matrix_T.HLSF.cs
@@ -1,6 +1,6 @@
 //-----------------------------------------------------------------------
 // -----------------------------------------------------------------------
-// file="Matrix.HLSF"
+// file="Matrix_T.HLSF"
 // Project: StgSharp
 // AuthorGroup: Nitload
 // Copyright (c) Nitload. All rights reserved.
@@ -103,10 +103,10 @@ namespace StgSharp.Mathematics.Numeric
 
         #region enumeration
 
-        [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public override unsafe MatrixSegmentEnumeration<T> GetColumnEnumeration()
+        public override unsafe MatrixSegmentEnumeration<T> ColumnEnumeration
         {
-            return new MatrixSegmentEnumeration<T>(k_buffer, KernelColumnLength, KernelColumnLength * KernelRowLength,
+            [MethodImpl(MethodImplOptions.AggressiveInlining)]
+            get => new MatrixSegmentEnumeration<T>(k_buffer, KernelColumnLength, KernelColumnLength * KernelRowLength,
                                                    1, KernelColumnLength);
         }
 
@@ -116,10 +116,11 @@ namespace StgSharp.Mathematics.Numeric
                                                    KernelRowLength);
         }
 
-        [MethodImpl(MethodImplOptions.AggressiveInlining)]
-        public override unsafe MatrixSegmentEnumeration<T> GetSequentialEnumeration()
+        public override unsafe MatrixSegmentEnumeration<T> SequentialEnumeration
         {
-            return new MatrixSegmentEnumeration<T>(k_buffer, KernelRowLength * KernelColumnLength,
+            [MethodImpl(MethodImplOptions.AggressiveInlining)]
+
+            get => new MatrixSegmentEnumeration<T>(k_buffer, KernelRowLength * KernelColumnLength,
                                                    KernelColumnLength * KernelRowLength, 1,
                                                    KernelRowLength * KernelColumnLength);
         }
diff --git a/StgSharp.Common/Mathematics/Linear/Matrix_T.cs b/StgSharp.Common/Mathematics/Linear/Matrix_T.cs
new file mode 100644
index 00000000..3d1e3cff
--- /dev/null
+++ b/StgSharp.Common/Mathematics/Linear/Matrix_T.cs
@@ -0,0 +1,218 @@
+//-----------------------------------------------------------------------
+// -----------------------------------------------------------------------
+// file="Matrix_T"
+// Project: StgSharp
+// AuthorGroup: Nitload
+// Copyright (c) Nitload. All rights reserved.
+//     
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+//     
+// -----------------------------------------------------------------------
+// -----------------------------------------------------------------------
+using StgSharp.HighPerformance;
+using System;
+using System.Numerics;
+using System.Runtime.CompilerServices;
+using System.Runtime.InteropServices;
+
+// using hlsfHandle = global::StgSharp.HighPerformance.Memory.HybridLayerSegregatedFitAllocationHandle;
+
+namespace StgSharp.Mathematics.Numeric
+{
+    public abstract unsafe partial class Matrix<T> where T : unmanaged, INumber<T>
+    {
+
+        private protected readonly int t_count = 16 / sizeof(T);
+
+        private protected readonly int t_size = sizeof(T);
+
+        private protected Matrix() { }
+
+        public abstract ref T this[int column, int row]
+        {
+            get;
+        }
+
+        public abstract unsafe Matrix<T> this[ReadOnlySpan<MatrixIndex> column, ReadOnlySpan<MatrixIndex> row]
+        {
+            get;
+            set;
+        }
+
+        /// <summary>
+        ///   Length of the matrix in columns.
+        /// </summary>
+        public int ColumnLength { get; init; }
+
+        /// <summary>
+        ///   Length of the matrix in rows.
+        /// </summary>
+        public int RowLength { get; init; }
+
+        public abstract MatrixSegmentEnumeration<T> ColumnEnumeration { get; }
+
+        public abstract MatrixSegmentEnumeration<T> SequentialEnumeration { get; }
+
+        internal abstract MatrixKernel<T>* Buffer { get; }
+
+        protected internal int KernelColumnLength { get; init; }
+
+        protected internal int KernelRowLength { get; init; }
+
+        public abstract ref MatrixKernel<T> GetMatrixKernel(int column, int row);
+
+        public abstract MatrixSegmentEnumeration<T> GetRowEnumeration();
+
+        [Obsolete("This method does not contain boarder check, so be careful on size or use indexer instead.", false)]
+        public abstract ref T UnsafeIndex(int column, int row);
+
+        private protected void AssertIndex(int column, int row)
+        {
+            if (column < 0 || column >= ColumnLength) {
+                throw new ArgumentOutOfRangeException(nameof(column));
+            }
+
+            if (row < 0 || row >= RowLength) {
+                throw new ArgumentOutOfRangeException(nameof(row));
+            }
+        }
+
+    }
+
+    public readonly unsafe ref struct MatrixSegmentEnumeration<T>(
+                                      MatrixKernel<T>* source,
+                                      int primStride,
+                                      int primSize,
+                                      int secondaryStride,
+                                      int secondarySize,
+                                      int primOffset = 0,
+                                      int secondaryOffset = 0)
+        where T : unmanaged, INumber<T>
+    {
+
+        private readonly MatrixKernel<T>* _source = source;
+        internal readonly int PrimOffset = primOffset;
+
+        internal readonly int PrimSize = primSize;
+        internal readonly int PrimStride = primStride;
+        internal readonly int SecondaryOffset = secondaryOffset;
+        internal readonly int SecondarySize = secondarySize;
+        internal readonly int SecondaryStride = secondaryStride;
+
+        internal readonly MatrixKernel<T>* Source
+        {
+            [MethodImpl(MethodImplOptions.AggressiveInlining)]
+            get => _source;
+        }
+
+        public Enumerator GetEnumerator()
+        {
+            return new Enumerator(in this);
+        }
+
+        public ref struct Enumerator(in MatrixSegmentEnumeration<T> enumeration)
+        {
+
+            private readonly MatrixKernel<T>* _source = enumeration._source;
+            private int _indexCount = -1;
+            private readonly int _stride = enumeration.PrimStride;
+            private readonly long _bound = (long)enumeration.PrimSize * enumeration.PrimStride;
+            private long _index = (long)enumeration.PrimOffset * enumeration.PrimStride;
+
+            private readonly MatrixSegmentEnumeration<T> _enumeration;
+
+            public readonly MatrixSegment<T> Current
+            {
+                [MethodImpl(MethodImplOptions.AggressiveInlining)]
+                get => new(_source + _index, _enumeration.SecondarySize, _enumeration.SecondaryStride,
+                           _enumeration.PrimOffset);
+            }
+
+            public bool MoveNext()
+            {
+                _indexCount++;
+                _index += _stride;
+                if (_index >= _bound) {
+                    _index = 0;
+                }
+                return _indexCount >= _enumeration.PrimSize;
+            }
+
+        }
+
+    }
+
+    public readonly unsafe ref struct MatrixSegment<T>(MatrixKernel<T>* source, int size, int stride, int offset)
+        where T : unmanaged, INumber<T>
+    {
+
+        private readonly MatrixKernel<T>* _source = source;
+        private readonly int _offset = offset;
+        private readonly int _secondarySize = size;
+        private readonly int _size = size;
+        private readonly int _stride = stride;
+
+        public ref MatrixKernel<T> this[int index]
+        {
+            get
+            {
+                if (index < 0 || index >= _size) {
+                    throw new ArgumentOutOfRangeException(nameof(index));
+                }
+                return ref _source[index * _stride];
+            }
+        }
+
+        [MethodImpl(MethodImplOptions.AggressiveInlining)]
+        public Enumerator GetEnumerator()
+        {
+            return new Enumerator(in this);
+        }
+
+        public ref struct Enumerator(in MatrixSegment<T> segment)
+        {
+
+            private readonly int _end = segment._size;
+
+            private int _indexCount = -1;
+            private readonly long _bound = (long)segment._size * segment._stride;
+            private long _index = (long)segment._offset * segment._stride;
+            private readonly MatrixSegment<T> _segment = segment;
+
+            public readonly MatrixKernel<T>* Current
+            {
+                [MethodImpl(MethodImplOptions.AggressiveInlining)]
+                get => _segment._source + _index;
+            }
+
+            public bool MoveNext()
+            {
+                _index += _segment._stride;
+                _indexCount++;
+                if (_index >= _bound) {
+                    _index = 0;
+                }
+
+                return _indexCount < _end;
+            }
+
+        }
+
+    }
+}
diff --git a/StgSharp.Mathematics/Numeric/MatrixCompute.Basic.cs b/StgSharp.Mathematics/Numeric/MatrixCompute.Basic.cs
index c8e6eed8..ffa2731e 100644
--- a/StgSharp.Mathematics/Numeric/MatrixCompute.Basic.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixCompute.Basic.cs
@@ -47,7 +47,7 @@ namespace StgSharp.Mathematics.Numeric
                 throw new IndexOutOfRangeException("Matrix dimensions must agree.");
             }
             long count = (long)left.KernelColumnLength * left.KernelRowLength;
-            Matrix<float> ans = Matrix<float>.FromDefault(left.ColumnLength, left.RowLength);
+            Matrix<float> ans = Matrix.FromDefault<float>(left.ColumnLength, left.RowLength);
             if (count <= 256)
             {
                 MatrixKernel<float>* leftPtr = left.Buffer;
diff --git a/StgSharp.Mathematics/Numeric/MatrixComputeModel.cs b/StgSharp.Mathematics/Numeric/MatrixComputeModel.cs
index bea0bb03..bb6ad4ad 100644
--- a/StgSharp.Mathematics/Numeric/MatrixComputeModel.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixComputeModel.cs
@@ -47,6 +47,16 @@ namespace StgSharp.Mathematics.Numeric
             MatrixParallelFactory.Release(p);
         }
 
+        internal static unsafe void BufferComputeNoOperatorScalar(MatrixParallelTaskPackage* p)
+        {
+            int size = p->ElementSize * 16;
+            long count = Unsafe.As<int, long>(ref p->PrimCount);
+            long offset = Unsafe.As<int, long>(ref p->PrimTileOffset) * size;
+            ((delegate* unmanaged[Cdecl]<void*, ScalarPacket*,long, void>)p->ComputeHandle)((void*)(p->Result + offset),
+                                                                                                     p->Scalar, count);
+            MatrixParallelFactory.Release(p);
+        }
+
         internal static unsafe void BufferComputeUnaryScalar(MatrixParallelTaskPackage* p)
         {
             int size = p->ElementSize * 16;
diff --git a/StgSharp.Mathematics/Numeric/MatrixParallel.Publish.cs b/StgSharp.Mathematics/Numeric/MatrixParallel.Publish.cs
index 29f35f4e..61bca6fe 100644
--- a/StgSharp.Mathematics/Numeric/MatrixParallel.Publish.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixParallel.Publish.cs
@@ -67,17 +67,21 @@ namespace StgSharp.Mathematics.Numeric
             int primCount = package->PrimCount;
             int secCount = package->SecCount;
             long count = (long)primCount * secCount;
+            Console.WriteLine($"amount of kernels is {count}");
 
             long recommendCount = count / 1024;
             MatrixParallelThread[] pool = LeadParallel((int)recommendCount);
             MatrixParallelHandle handle = new MatrixParallelHandle(pool);
-            if (handle.Wraps.Count() == 1)
+            if (handle.Wraps.Length == 1)
             {
+                // Console.WriteLine($"Scheduling into {pool.Length} threads.");
                 ref long offsetRef = ref Unsafe.As<int, long>(ref package->PrimTileOffset);
                 offsetRef = 0;
                 ref long lengthRef = ref Unsafe.As<int, long>(ref package->PrimCount);
                 lengthRef = count;
-                handle.Wraps.First().WrapTask.Enqueue((nint)package);
+                handle.Wraps[0].WrapTask.Enqueue((nint)package);
+
+                // Console.WriteLine($"Scheduling {count} kernels into thread");
                 return handle;
             }
             long actual = count / pool.Length;
@@ -85,18 +89,25 @@ namespace StgSharp.Mathematics.Numeric
 
             long offset = 0;
             MatrixParallelTaskPackage* current = package;
+
+            // Console.WriteLine($"Scheduling into {pool.Length} threads.");
             foreach (MatrixParallelWrap wrap in handle.Wraps)
             {
                 long wrapTaskLength = actual * wrap.WrapTaskCapacity;
                 long ex = Math.Min(extra, wrap.WrapTaskCapacity);
+                long length = wrapTaskLength + ex;
                 extra -= ex;
                 ref long offsetRef = ref Unsafe.As<int, long>(ref current->PrimTileOffset);
                 offsetRef = offset;
                 ref long lengthRef = ref Unsafe.As<int, long>(ref current->PrimCount);
-                lengthRef = wrapTaskLength + ex;
+                lengthRef = length;
                 wrap.WrapTask.Enqueue((nint) current);
+
+                Console.WriteLine($"Offset is {offset}");
+                Console.WriteLine($"Scheduling {length} kernels into {wrap.WrapTaskCapacity} threads");
+                Console.WriteLine();
                 current = MatrixParallelFactory.FromExistPackage(current);
-                offset += wrapTaskLength;
+                offset += length;
             }
             return handle;
         }
@@ -117,9 +128,9 @@ namespace StgSharp.Mathematics.Numeric
             }
             /**/
             MatrixParallelHandle handle = new MatrixParallelHandle(pool);
-            if (handle.Wraps.Count() == 1)
+            if (handle.Wraps.Length == 1)
             {
-                handle.Wraps.First().WrapTask.Enqueue((nint)package);
+                handle.Wraps[0].WrapTask.Enqueue((nint)package);
                 return handle;
             }
             int threadCount = pool.Length;
@@ -129,7 +140,6 @@ namespace StgSharp.Mathematics.Numeric
             int extraColumnCount = columnAfterBaseTile % threadCount;                  // columns to be sliced
 
             int total = primCount;
-            Stack<MatrixParallelWrap> taskBunch = new Stack<MatrixParallelWrap>(threadCount);
             MatrixParallelTaskPackage* current = package;
             foreach (MatrixParallelWrap wrap in handle.Wraps)
             {
diff --git a/StgSharp.Mathematics/Numeric/MatrixParallel.cs b/StgSharp.Mathematics/Numeric/MatrixParallel.cs
index 7f33c363..ecf98cb2 100644
--- a/StgSharp.Mathematics/Numeric/MatrixParallel.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixParallel.cs
@@ -42,7 +42,6 @@ namespace StgSharp.Mathematics.Numeric
 
         private static CapacityFixedStack<MatrixParallelThread> _threads;
 
-        private static volatile int RemainThreadCount = 0;
         private static readonly
 #if NET10_0_OR_GREATER
             Lock
@@ -68,13 +67,13 @@ namespace StgSharp.Mathematics.Numeric
 
         public static void LaunchParallel(MatrixParallelHandle handle, SleepMode afterWork)
         {
-            RemainThreadCount = _threads.Count;
             _AfterWork = (ThreadPriority)afterWork;
             MatrixParallelThread leader = handle.GetLeader();
             foreach (MatrixParallelWrap w in handle.Wraps) {
                 w.SetScheduler();
             }
             leader.MatrixParallelLeader();
+            handle.WaitAllWraps();
         }
 
         internal static void Init()
diff --git a/StgSharp.Mathematics/Numeric/MatrixParallelHandle.cs b/StgSharp.Mathematics/Numeric/MatrixParallelHandle.cs
index f16124b2..3ab1a077 100644
--- a/StgSharp.Mathematics/Numeric/MatrixParallelHandle.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixParallelHandle.cs
@@ -28,6 +28,7 @@
 using StgSharp.Collections;
 using System;
 using System.Collections.Generic;
+using System.ComponentModel.DataAnnotations;
 using System.Diagnostics.CodeAnalysis;
 using System.Linq;
 using System.Runtime.CompilerServices;
@@ -36,7 +37,7 @@ using System.Threading.Tasks;
 
 namespace StgSharp.Mathematics.Numeric
 {
-    public struct MatrixParallelHandle : IDisposable
+    public class MatrixParallelHandle : IDisposable
     {
 
         private static readonly (int a, int b)[] LookupTable = [
@@ -55,10 +56,11 @@ namespace StgSharp.Mathematics.Numeric
             (1, 4), // 19
             (4, 2), // 20
         ];
-
         private MatrixParallelThread[] _threads;
+        private MatrixParallelWrap[] _wraps;
         private bool _disposed = false;
-        private CapacityFixedStack<MatrixParallelWrap> _wraps;
+        private int _count;
+        private ManualResetEventSlim _leaderEvent = new ManualResetEventSlim(false);
 
         internal MatrixParallelHandle(MatrixParallelThread[] threadArray)
         {
@@ -78,7 +80,7 @@ namespace StgSharp.Mathematics.Numeric
             }
         }
 
-        internal IEnumerable<MatrixParallelWrap> Wraps
+        internal ReadOnlySpan<MatrixParallelWrap> Wraps
         {
             get
             {
@@ -89,21 +91,8 @@ namespace StgSharp.Mathematics.Numeric
             }
         }
 
-        internal ReadOnlySpan<MatrixParallelThread> ThreadsUnsafe
-        {
-            [MethodImpl(MethodImplOptions.AggressiveInlining)]
-            get => _threads;
-        }
-
         private bool IsMultithread { get; init; }
 
-        public MatrixParallelThread DefineLeader()
-        {
-            MatrixParallelThread leader = _threads[1];
-            leader.Role = 2;
-            return leader;
-        }
-
         public void Dispose()
         {
             if (_disposed) {
@@ -115,8 +104,27 @@ namespace StgSharp.Mathematics.Numeric
 
         public MatrixParallelThread GetLeader()
         {
-            Threads[0].Role = 2;
-            return Threads[0];
+            MatrixParallelThread leader = Threads[1];
+            leader.Role = MatrixParallelThread.RoleType.Leader;
+            leader.LeaderEvent = _leaderEvent;
+            return leader;
+        }
+
+        public void WaitAllWraps()
+        {
+            if (_disposed) {
+                throw new ObjectDisposedException(typeof(MatrixParallelHandle).Name);
+            }
+            SpinWait.SpinUntil(() => Volatile.Read(ref _count) == 0);
+        }
+
+        [MethodImpl(MethodImplOptions.AggressiveInlining)]
+        internal void ReportWrapAccomplished()
+        {
+            if (_disposed) {
+                throw new ObjectDisposedException(typeof(MatrixParallelHandle).Name);
+            }
+            Interlocked.Decrement(ref _count);
         }
 
         [MethodImpl(MethodImplOptions.AggressiveInlining)]
@@ -163,23 +171,26 @@ namespace StgSharp.Mathematics.Numeric
             int count = _threads.Length;
             if (count < 6)
             {
-                MatrixParallelWrap wrap = new MatrixParallelWrap(_threads, 0, count);
+                MatrixParallelWrap wrap = new MatrixParallelWrap(this, 0, count);
                 _wraps = [ wrap ];
-                return;
-            }
-            (int c3,int c4) = FastDecompose(_threads.Length);
-            _wraps = new(c3 + c4);
-            int begin = 0;
-            for (int i = 0; i < c3; i++)
+            } else
             {
-                _wraps.Push(new MatrixParallelWrap(_threads, begin, 3));
-                begin += 3;
-            }
-            for (int i = 0; i < c4; i++)
-            {
-                _wraps.Push(new MatrixParallelWrap(_threads, begin, 4));
-                begin += 4;
+                (int c3, int c4) = FastDecompose(_threads.Length);
+                _wraps = new MatrixParallelWrap[c3 + c4];
+                int begin = 0;
+                int i = 0;
+                for (; i < c3; i++)
+                {
+                    _wraps[i] = new MatrixParallelWrap(this, begin, 3);
+                    begin += 3;
+                }
+                for (; i < c3 + c4; i++)
+                {
+                    _wraps[i] = new MatrixParallelWrap(this, begin, 4);
+                    begin += 4;
+                }
             }
+            _count = _wraps.Length;
         }
 
     }
diff --git a/StgSharp.Mathematics/Numeric/MatrixParallelTaskPackage.cs b/StgSharp.Mathematics/Numeric/MatrixParallelTaskPackage.cs
index 538352f4..c6429ee1 100644
--- a/StgSharp.Mathematics/Numeric/MatrixParallelTaskPackage.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixParallelTaskPackage.cs
@@ -119,19 +119,19 @@ namespace StgSharp.Mathematics.Numeric
 
         #region right enum
 
-        [FieldOffset(32)]public int RightPrimOffset;        // 4
-        [FieldOffset(36)]public int RightPrimStride;        // 4
-        [FieldOffset(40)]public int RightSecOffset;         // 4
-        [FieldOffset(44)]public int RightSecStride;         // 4
+        [FieldOffset(48)]public int RightPrimOffset;        // 4
+        [FieldOffset(52)]public int RightPrimStride;        // 4
+        [FieldOffset(56)]public int RightSecOffset;         // 4
+        [FieldOffset(60)]public int RightSecStride;         // 4
 
         #endregion
 
         #region ans enum
 
-        [FieldOffset(48)]public int ResultPrimOffset;       // 4
-        [FieldOffset(52)]public int ResultPrimStride;       // 4
-        [FieldOffset(56)]public int ResultSecOffset;        // 4
-        [FieldOffset(60)] public int ResultSecStride;        // 4
+        [FieldOffset(64)]public int ResultPrimOffset;       // 4
+        [FieldOffset(68)]public int ResultPrimStride;       // 4
+        [FieldOffset(72)]public int ResultSecOffset;        // 4
+        [FieldOffset(76)] public int ResultSecStride;        // 4
 
         #endregion
 
diff --git a/StgSharp.Mathematics/Numeric/MatrixParallelThread.cs b/StgSharp.Mathematics/Numeric/MatrixParallelThread.cs
index 08490395..6b801574 100644
--- a/StgSharp.Mathematics/Numeric/MatrixParallelThread.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixParallelThread.cs
@@ -72,6 +72,8 @@ namespace StgSharp.Mathematics.Numeric
             set => field = value;
         }
 
+        internal ManualResetEventSlim? LeaderEvent { get; set; }
+
         public void Dispose()
         {
             Dispose(disposing:true);
@@ -102,6 +104,7 @@ namespace StgSharp.Mathematics.Numeric
                             case MatrixOperationParam.LEFT_RIGHT_ANS_SCALAR_PARAM:
                                 break;
                             case MatrixOperationParam.ANS_SCALAR_PARAM:
+                                MatrixComputeModel.BufferComputeNoOperatorScalar(p);
                                 break;
                             default:
                                 break;
@@ -135,21 +138,30 @@ namespace StgSharp.Mathematics.Numeric
         [MethodImpl(MethodImplOptions.AggressiveInlining)]
         public void Restart()
         {
-            _resetEvent.Set();
+            if (Role == RoleType.Leader)
+            {
+                LeaderEvent?.Set();
+            } else
+            {
+                _resetEvent.Set();
+            }
         }
 
         internal void MatrixParallelLeader()
         {
             Thread current = Thread.CurrentThread;
             current.Priority = ThreadPriority.Highest;
-            Thread.Sleep(0);
+            LeaderEvent?.Wait();
+            LeaderEvent?.Reset();
 
+            // Console.WriteLine("Leader is computing");
 #if WINDOWS
             if (!UseUmaAffinityPath)
             {
                 // NUMA/affinity path: bind to this thread's reserved mask, then restore
                 nuint previousMask = SetThreadAffinityMask(GetCurrentThread(), _affinity.Mask);
                 bool affinityChanged = previousMask != 0; // nonzero => success
+                Thread.Sleep(0);
                 ExecuteTaskPackage(Tasks);
                 if (affinityChanged) {
                     _ = SetThreadAffinityMask(GetCurrentThread(), previousMask);
@@ -164,7 +176,11 @@ namespace StgSharp.Mathematics.Numeric
             ExecuteTaskPackage(Tasks);
 #endif
 
-            this.Role = 1;
+            BindedWrap.ReportThreadAccomplished();
+
+            // Console.WriteLine("leader's computing is done");
+            Role = RoleType.Worker;
+            LeaderEvent = null;
             current.Priority = ThreadPriority.Normal;
         }
 
@@ -228,16 +244,27 @@ namespace StgSharp.Mathematics.Numeric
                 _resetEvent.Reset();
                 switch (Role)
                 {
-                    case 0://worker
+                    case RoleType.Worker://worker
                         // _ = Interlocked.Decrement(ref RemainThreadCount);
+                        // Console.WriteLine("worker start");
                         ExecuteTaskPackage(Tasks);
+                        BindedWrap.ReportThreadAccomplished();
+
+                        // Console.WriteLine("worker end");
                         break;
-                    case 1://scheduler
-                        BindedWrap.PublishTaskToWorkers(BindedWrap.WrapTask);
+                    case RoleType.Scheduler://scheduler
+                        Console.WriteLine("scheduler start");
+                        BindedWrap.PublishTaskToWorkers();
                         BindedWrap.SetAllWorker();
+
+                        Console.WriteLine("scheduler done");
                         ExecuteTaskPackage(Tasks);
+                        BindedWrap.ReportThreadAccomplished();
+                        BindedWrap.WaitAllWorker();
+                        BindedWrap.ReportWrapAccomplished();
+                        Role = RoleType.Worker;
                         break;
-                    case 2://leader, error because leader should be at outer thread but not here
+                    case RoleType.Leader://leader, error because leader should be at outer thread but not here
                         break;
                     default:
                         break;
@@ -389,21 +416,15 @@ namespace StgSharp.Mathematics.Numeric
 
         internal Queue<IntPtr> Tasks { get; } = [];
 
-        public int Role
+        public RoleType Role { get; internal set; }
+
+        public enum RoleType : int
         {
-            get;
-            internal set
-            {
-                if ((Role & (~3u)) == 0)
-                {
-                    field = value;
-                } else
-                {
-#if DEBUG
-                    throw new ArgumentOutOfRangeException(nameof(Role));
-#endif
-                }
-            }
+
+            Worker = 0,
+            Scheduler = 1,
+            Leader = 2,
+
         }
 
         #endregion
diff --git a/StgSharp.Mathematics/Numeric/MatrixParallelWrap.cs b/StgSharp.Mathematics/Numeric/MatrixParallelWrap.cs
index a8237843..82a10ff0 100644
--- a/StgSharp.Mathematics/Numeric/MatrixParallelWrap.cs
+++ b/StgSharp.Mathematics/Numeric/MatrixParallelWrap.cs
@@ -30,6 +30,7 @@ using System;
 using System.Collections.Generic;
 using System.Diagnostics.CodeAnalysis;
 using System.Linq;
+using System.Net.WebSockets;
 using System.Runtime.CompilerServices;
 using System.Text;
 using System.Threading.Tasks;
@@ -39,23 +40,22 @@ namespace StgSharp.Mathematics.Numeric
     public unsafe class MatrixParallelWrap
     {
 
-        private readonly MatrixParallelThread[] _threads;
         private int _begin;
 
-        internal MatrixParallelWrap(MatrixParallelThread[] threads, int begin, int count)
+        private int _remainThreadCount;
+        private readonly MatrixParallelHandle _handle;
+
+        internal MatrixParallelWrap(MatrixParallelHandle handle, int begin, int count)
         {
             WrapTaskCapacity = count;
-            _threads = threads;
+            _handle = handle;
             _begin = begin;
-            _threads[_begin].Role = 1;
-            for (int i = _begin; i < _begin + WrapTaskCapacity; i++) {
-                _threads[i].BindedWrap = this;
+            for (int i = 0; i < WrapTaskCapacity; i++)
+            {
+                Threads[i].Role = MatrixParallelThread.RoleType.Worker;
+                Threads[i].BindedWrap = this;
             }
-        }
-
-        public MatrixParallelWrap()
-        {
-            _threads = new MatrixParallelThread[5];
+            Threads[0].Role = MatrixParallelThread.RoleType.Scheduler;
         }
 
         public static SleepMode AfterWork
@@ -66,19 +66,41 @@ namespace StgSharp.Mathematics.Numeric
 
         internal int Magic { get; set; }
 
+        private ReadOnlySpan<MatrixParallelThread> Threads
+        {
+            [MethodImpl(MethodImplOptions.AggressiveInlining)]
+            get => _handle.Threads[_begin..(_begin + WrapTaskCapacity)];
+        }
+
+        [MethodImpl(MethodImplOptions.AggressiveInlining)]
+        public void ReportThreadAccomplished()
+        {
+            Interlocked.Decrement(ref _remainThreadCount);
+        }
+
+        [MethodImpl(MethodImplOptions.AggressiveInlining)]
+        public void ReportWrapAccomplished()
+        {
+            _handle.ReportWrapAccomplished();
+        }
+
         [MethodImpl(MethodImplOptions.AggressiveInlining)]
         public void SetAllWorker()
         {
-            for (int i = _begin; i < _begin + WrapTaskCapacity; i++)
-            {
-                if (_threads[i].Role != 3) {
-                    _threads[i].Restart();
-                }
+            for (int i = 1; i < WrapTaskCapacity; i++) {
+                Threads[i].Restart();
             }
         }
 
-        internal void PublishTaskToWorkers(Queue<IntPtr> tasks)
+        [MethodImpl(MethodImplOptions.AggressiveInlining)]
+        public void WaitAllWorker()
         {
+            SpinWait.SpinUntil(() => Volatile.Read(ref _remainThreadCount) == 0);
+        }
+
+        internal void PublishTaskToWorkers()
+        {
+            Queue<nint> tasks = WrapTask;
             while (tasks.TryDequeue(out nint handle))
             {
                 MatrixParallelTaskPackage* task = (MatrixParallelTaskPackage*)handle;
@@ -100,45 +122,54 @@ namespace StgSharp.Mathematics.Numeric
 
         internal void SetScheduler()
         {
-            MatrixParallelThread thread = _threads[_begin];
+            MatrixParallelThread thread = Threads[0];
             thread.Restart();
+            _ = Interlocked.Exchange(ref _remainThreadCount, WrapTaskCapacity);
         }
 
         private void PublishBufferTaskToWorkers(MatrixParallelTaskPackage* package)
         {
-            long count = Unsafe.As<int, long>(ref package->PrimCount);
-            long recommendCount = count / 1024;
-            MatrixParallelThread[] pool = _threads;
+            ReadOnlySpan<MatrixParallelThread> pool = Threads;
             if (pool.Length == 1)
             {
                 pool[0].Tasks.Enqueue((nint)package);
                 return;
             }
+            long count = Unsafe.As<int, long>(ref package->PrimCount);
             long actual = count / pool.Length;
             long extra = count % pool.Length;
 
             long offset = 0;
             MatrixParallelTaskPackage* current = package;
+            string msg = $"""
+                Total: count={count}, length={Threads.Length}{Environment.NewLine}
+                """;
             foreach (MatrixParallelThread thread in pool)
             {
                 long wrapTaskLength = actual;
                 long ex = extra > 0 ? 1 : 0;
+                long length = wrapTaskLength + ex;
                 extra -= ex;
                 ref long offsetRef = ref Unsafe.As<int, long>(ref current->PrimTileOffset);
                 offsetRef += offset;
                 ref long lengthRef = ref Unsafe.As<int, long>(ref current->PrimCount);
-                lengthRef = wrapTaskLength + ex;
-                pool[0].Tasks.Enqueue((nint)current);
+                lengthRef = length;
+                thread.Tasks.Enqueue((nint)current);
+                msg += $"""
+                    Slice to thread: total offset={offsetRef} local offset={offset}, length={lengthRef}{Environment.NewLine}
+                    """;
+
                 current = MatrixParallelFactory.FromExistPackage(current);
-                offset += wrapTaskLength;
+                offset = length;
             }
+            Console.WriteLine(msg);
         }
 
         private void PublishKernelTaskToWorkers(MatrixParallelTaskPackage* package)
         {
             int primCount = package->PrimCount;
             int secCount = package->SecCount;
-            MatrixParallelThread[] pool = _threads;
+            ReadOnlySpan<MatrixParallelThread> pool = Threads;
 
             if (pool.Length == 1)
             {
@@ -201,31 +232,10 @@ namespace StgSharp.Mathematics.Numeric
             return;
         }
 
-        public class Handle(int inWrapIndex, MatrixParallelWrap wrap)
-        {
-
-            public int InWrapIndex { get; } = inWrapIndex;
-
-            /// <summary>
-            ///   Role of the thread in the wrap. 1 is Scheduler, 0 are Workers. 2 is leader, beyond
-            ///   wrap's control.
-            /// </summary>
-            public int Role { get; internal set; }
-
-            public Queue<IntPtr> TaskCache { get; } = [];
-
-            internal ManualResetEventSlim ResetEvent { get; } = new(false);
-
-            internal MatrixParallelWrap Wrap { get; } = wrap;
-
-        }
-
         #region wrap task
 
         internal unsafe Queue<nint> WrapTask { get; set; } = [];
 
-        internal int WrapTaskSliceCount { get; set; }
-
         internal int WrapTaskCapacity { get ; init; }
 
         #endregion
diff --git a/StgSharp.Native/sn-src/k_bufop.c b/StgSharp.Native/sn-src/k_bufop.c
index cfde2dc5..62a881b4 100644
--- a/StgSharp.Native/sn-src/k_bufop.c
+++ b/StgSharp.Native/sn-src/k_bufop.c
@@ -46,6 +46,7 @@ DECLARE_BUF_PROC_LEFT_RIGHT_ANS(float, sse, add)
 DECLARE_BUF_PROC_LEFT_RIGHT_ANS(float, avx, add)
 {
         for (size_t i = 0; i < count; i++) {
+                size_t s = sizeof(MAT_KERNEL(float));
                 register __m256 c0 = _mm256_add_ps(left[i].ymm[0], right[i].ymm[0]);
                 register __m256 c1 = _mm256_add_ps(left[i].ymm[1], right[i].ymm[1]);
                 ans[i].ymm[0] = c0;
diff --git a/StgSharpDebug/Program.cs b/StgSharpDebug/Program.cs
index c7663d88..143c1a94 100644
--- a/StgSharpDebug/Program.cs
+++ b/StgSharpDebug/Program.cs
@@ -67,10 +67,18 @@ namespace StgSharp
 
 
                 // 原有的矩阵测试代码
-                HybridLayerSegregatedFitAllocator hlsf = new HybridLayerSegregatedFitAllocator(64 * 400 * 400);
+                HybridLayerSegregatedFitAllocator hlsf = new HybridLayerSegregatedFitAllocator(64 * 4000 * 4000);
 
-                Matrix<float> m = Matrix<float>.FromHlsf(200, 200, hlsf);
+                Matrix<float> m = Matrix.FromHlsf<float>(2000, 2000, hlsf);
                 m.Fill(1.0f);
+                for (int i = 0; i < 20; i++)
+                {
+                    for (int j = 0; j < 20; j++) {
+                        Console.Write($"{m[i, j]:0.0E+00} \t");
+                    }
+                    Console.WriteLine();
+                }
+                Console.WriteLine();
                 Matrix<float> n = m.Add(m);
                 for (int i = 0; i < 20; i++)
                 {