From 2d14dd68dc9ea2823db17231395f6c191cbc4607 Mon Sep 17 00:00:00 2001 From: craig-b <3112503+craig-b@users.noreply.github.com> Date: Sat, 15 Aug 2026 14:31:40 +0100 Subject: [PATCH 1/5] Migrate TensorSharp.Core P/Invokes to LibraryImport Mechanical DllImport -> source-generated LibraryImport conversion. TS_Std/TS_Var normByN now marshals explicitly as I1 (1-byte C++ bool); the previous default was a 4-byte BOOL, equivalent for a by-value register argument on the supported ABIs. --- TensorSharp.Core/Cpu/CpuOpsNative.cs | 377 +++++++++++++------ TensorSharp.Core/Cpu/MatrixMultiplication.cs | 10 +- TensorSharp.Core/Cpu/OpenBlasNative.cs | 36 +- 3 files changed, 287 insertions(+), 136 deletions(-) diff --git a/TensorSharp.Core/Cpu/CpuOpsNative.cs b/TensorSharp.Core/Cpu/CpuOpsNative.cs index 600dd98d..491cdad8 100644 --- a/TensorSharp.Core/Cpu/CpuOpsNative.cs +++ b/TensorSharp.Core/Cpu/CpuOpsNative.cs @@ -8,6 +8,7 @@ // TensorSharp is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even the implied warranty of // MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the BSD-3-Clause License for more details. using System; +using System.Runtime.CompilerServices; using System.Runtime.InteropServices; namespace TensorSharp.Cpu @@ -32,131 +33,275 @@ public struct TensorRef64 } - public static class CpuOpsNative + public static partial class CpuOpsNative { private const string dll = "CpuOps.dll"; private const CallingConvention cc = CallingConvention.Cdecl; - [DllImport(dll, CallingConvention = cc)] - public static extern IntPtr TS_GetLastError(); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Copy(IntPtr result, IntPtr src); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Abs(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Neg(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Sign(IntPtr result, IntPtr src); - - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Sqrt(IntPtr result, IntPtr src); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Log1p(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Floor(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Ceil(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Round(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Trunc(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Frac(IntPtr result, IntPtr src); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Sin(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Cos(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Tan(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Asin(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Acos(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Atan(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Sinh(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Cosh(IntPtr result, IntPtr src); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Add3(IntPtr result, IntPtr x, IntPtr y, IntPtr z); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Add4(IntPtr result, IntPtr x, IntPtr y, IntPtr z, IntPtr w); - - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_MaskFill(IntPtr result, IntPtr t, IntPtr mask, float defValue); - - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Atan2(IntPtr result, IntPtr srcY, IntPtr srcX); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Tpow(IntPtr result, float value, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Lerp(IntPtr result, IntPtr srcA, IntPtr srcB, float weight); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial IntPtr TS_GetLastError(); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Copy(IntPtr result, IntPtr src); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Abs(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Neg(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Sign(IntPtr result, IntPtr src); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Sqrt(IntPtr result, IntPtr src); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Log1p(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Floor(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Ceil(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Round(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Trunc(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Frac(IntPtr result, IntPtr src); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Sin(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Cos(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Tan(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Asin(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Acos(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Atan(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Sinh(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Cosh(IntPtr result, IntPtr src); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Add3(IntPtr result, IntPtr x, IntPtr y, IntPtr z); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Add4(IntPtr result, IntPtr x, IntPtr y, IntPtr z, IntPtr w); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_MaskFill(IntPtr result, IntPtr t, IntPtr mask, float defValue); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Atan2(IntPtr result, IntPtr srcY, IntPtr srcX); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Tpow(IntPtr result, float value, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Lerp(IntPtr result, IntPtr srcA, IntPtr srcB, float weight); //[DllImport(dll, CallingConvention = cc)] public static extern int TS_Clamp(IntPtr result, IntPtr src, float min, float max); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_AddTanh3(IntPtr result, IntPtr srcX, IntPtr srcY, IntPtr srcZ); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Add(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Sub(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_AddTanh3(IntPtr result, IntPtr srcX, IntPtr srcY, IntPtr srcZ); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Add(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Sub(IntPtr result, IntPtr lhs, float rhs); //[DllImport(dll, CallingConvention = cc)] public static extern int TS_Div(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Rdiv(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Mod(IntPtr result, IntPtr lhs, float rhs); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_gtValue(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_ltValue(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_geValue(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_leValue(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_eqValue(IntPtr result, IntPtr lhs, float rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_neValue(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Rdiv(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Mod(IntPtr result, IntPtr lhs, float rhs); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_gtValue(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_ltValue(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_geValue(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_leValue(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_eqValue(IntPtr result, IntPtr lhs, float rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_neValue(IntPtr result, IntPtr lhs, float rhs); //[DllImport(dll, CallingConvention = cc)] public static extern int TS_CDiv(IntPtr result, IntPtr lhs, IntPtr rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_CMod(IntPtr result, IntPtr lhs, IntPtr rhs); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_gtTensor(IntPtr result, IntPtr lhs, IntPtr rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_ltTensor(IntPtr result, IntPtr lhs, IntPtr rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_geTensor(IntPtr result, IntPtr lhs, IntPtr rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_leTensor(IntPtr result, IntPtr lhs, IntPtr rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_eqTensor(IntPtr result, IntPtr lhs, IntPtr rhs); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_neTensor(IntPtr result, IntPtr lhs, IntPtr rhs); - - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Sum(IntPtr result, IntPtr src, int dimension); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Prod(IntPtr result, IntPtr src, int dimension); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Min(IntPtr result, IntPtr src, int dimension); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Argmin(IntPtr result, IntPtr src, int dimension); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Norm(IntPtr result, IntPtr src, int dimension, float value); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Std(IntPtr result, IntPtr src, int dimension, bool normByN); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_Var(IntPtr result, IntPtr src, int dimension, bool normByN); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_SumAll(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_ProdAll(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_MinAll(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_MaxAll(IntPtr result, IntPtr src); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_MeanAll(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_VarAll(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_StdAll(IntPtr result, IntPtr src); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_NormAll(IntPtr result, IntPtr src, float value); - - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_NewRNG(out IntPtr rng); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_DeleteRNG(IntPtr rng); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_SetRNGSeed(IntPtr rng, int newSeed); - - [DllImport(dll, CallingConvention = cc)] public static extern int TS_RandomUniform(IntPtr rng, IntPtr result, float min, float max); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_RandomNormal(IntPtr rng, IntPtr result, float mean, float stdv); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_RandomExponential(IntPtr rng, IntPtr result, float lambda); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_RandomCauchy(IntPtr rng, IntPtr result, float median, float sigma); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_RandomLogNormal(IntPtr rng, IntPtr result, float mean, float stdv); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_RandomGeometric(IntPtr rng, IntPtr result, float p); - [DllImport(dll, CallingConvention = cc)] public static extern int TS_RandomBernoulli(IntPtr rng, IntPtr result, float p); - - - [DllImport(dll, CallingConvention = cc)] - public static extern int TS_Unfolded_Acc(IntPtr finput, IntPtr input, int kW, int kH, int dW, int dH, int padW, int padH, int nInputPlane, int inputWidth, int inputHeight, int outputWidth, int outputHeight); - [DllImport(dll, CallingConvention = cc)] - public static extern int TS_Unfolded_Copy(IntPtr finput, IntPtr input, int kW, int kH, int dW, int dH, int padW, int padH, int nInputPlane, int inputWidth, int inputHeight, int outputWidth, int outputHeight); - - - [DllImport(dll, CallingConvention = cc)] - public static extern int TS_AddLayerNorm(IntPtr out_, IntPtr in1_, IntPtr in2_, IntPtr gamma_, IntPtr beta_, float eps, int rows, int cols); - - [DllImport(dll, CallingConvention = cc)] - public static extern int TS_AddLayerNormGrad(IntPtr result1, IntPtr result2, IntPtr gradGamma_, IntPtr gradBeta_, IntPtr adj_, IntPtr y_, IntPtr x1_, IntPtr x2_, IntPtr gamma_, IntPtr beta_, int rows, int cols, float eps); - - - [DllImport(dll, CallingConvention = cc)] - public static extern int TS_RMSProp(IntPtr tw, IntPtr tg, IntPtr tc, int rows, int cols, int batchSize, float step_size, float clipval, float regc, float decay_rate, float eps); - - [DllImport(dll, CallingConvention = cc)] - public static extern int TS_SpatialMaxPooling_updateOutput_frame(IntPtr input_p, IntPtr output_p, IntPtr ind_p, long nslices, long iwidth, long iheight, long owidth, long oheight, int kW, int kH, int dW, int dH, int padW, int padH); - - [DllImport(dll, CallingConvention = cc)] - public static extern int TS_SpatialMaxPooling_updateGradInput_frame(IntPtr gradInput, IntPtr gradOutput, IntPtr ind, long nslices, long iwidth, long iheight, long owidth, long oheight, int dW, int dH); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_CMod(IntPtr result, IntPtr lhs, IntPtr rhs); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_gtTensor(IntPtr result, IntPtr lhs, IntPtr rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_ltTensor(IntPtr result, IntPtr lhs, IntPtr rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_geTensor(IntPtr result, IntPtr lhs, IntPtr rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_leTensor(IntPtr result, IntPtr lhs, IntPtr rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_eqTensor(IntPtr result, IntPtr lhs, IntPtr rhs); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_neTensor(IntPtr result, IntPtr lhs, IntPtr rhs); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Sum(IntPtr result, IntPtr src, int dimension); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Prod(IntPtr result, IntPtr src, int dimension); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Min(IntPtr result, IntPtr src, int dimension); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Argmin(IntPtr result, IntPtr src, int dimension); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Norm(IntPtr result, IntPtr src, int dimension, float value); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Std(IntPtr result, IntPtr src, int dimension, [MarshalAs(UnmanagedType.I1)] bool normByN); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Var(IntPtr result, IntPtr src, int dimension, [MarshalAs(UnmanagedType.I1)] bool normByN); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_SumAll(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_ProdAll(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_MinAll(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_MaxAll(IntPtr result, IntPtr src); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_MeanAll(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_VarAll(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_StdAll(IntPtr result, IntPtr src); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_NormAll(IntPtr result, IntPtr src, float value); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_NewRNG(out IntPtr rng); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_DeleteRNG(IntPtr rng); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_SetRNGSeed(IntPtr rng, int newSeed); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RandomUniform(IntPtr rng, IntPtr result, float min, float max); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RandomNormal(IntPtr rng, IntPtr result, float mean, float stdv); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RandomExponential(IntPtr rng, IntPtr result, float lambda); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RandomCauchy(IntPtr rng, IntPtr result, float median, float sigma); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RandomLogNormal(IntPtr rng, IntPtr result, float mean, float stdv); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RandomGeometric(IntPtr rng, IntPtr result, float p); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RandomBernoulli(IntPtr rng, IntPtr result, float p); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Unfolded_Acc(IntPtr finput, IntPtr input, int kW, int kH, int dW, int dH, int padW, int padH, int nInputPlane, int inputWidth, int inputHeight, int outputWidth, int outputHeight); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_Unfolded_Copy(IntPtr finput, IntPtr input, int kW, int kH, int dW, int dH, int padW, int padH, int nInputPlane, int inputWidth, int inputHeight, int outputWidth, int outputHeight); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_AddLayerNorm(IntPtr out_, IntPtr in1_, IntPtr in2_, IntPtr gamma_, IntPtr beta_, float eps, int rows, int cols); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_AddLayerNormGrad(IntPtr result1, IntPtr result2, IntPtr gradGamma_, IntPtr gradBeta_, IntPtr adj_, IntPtr y_, IntPtr x1_, IntPtr x2_, IntPtr gamma_, IntPtr beta_, int rows, int cols, float eps); + + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_RMSProp(IntPtr tw, IntPtr tg, IntPtr tc, int rows, int cols, int batchSize, float step_size, float clipval, float regc, float decay_rate, float eps); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_SpatialMaxPooling_updateOutput_frame(IntPtr input_p, IntPtr output_p, IntPtr ind_p, long nslices, long iwidth, long iheight, long owidth, long oheight, int kW, int kH, int dW, int dH, int padW, int padH); + + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static partial int TS_SpatialMaxPooling_updateGradInput_frame(IntPtr gradInput, IntPtr gradOutput, IntPtr ind, long nslices, long iwidth, long iheight, long owidth, long oheight, int dW, int dH); // [DllImport(dll, CallingConvention = cc)] public static extern int TS_ScatterFill(IntPtr result, float value, int dim, IntPtr indices); } diff --git a/TensorSharp.Core/Cpu/MatrixMultiplication.cs b/TensorSharp.Core/Cpu/MatrixMultiplication.cs index 79dd8acf..96f0acd2 100644 --- a/TensorSharp.Core/Cpu/MatrixMultiplication.cs +++ b/TensorSharp.Core/Cpu/MatrixMultiplication.cs @@ -75,7 +75,7 @@ public enum Transpose ConjTrans = 113 } - unsafe public static class MatrixMultiplication + unsafe public static partial class MatrixMultiplication { const string mklDllName = "mkl_rt.2.dll"; internal const string mklDllNameLinux = "mkl_rt"; @@ -1374,8 +1374,8 @@ void ComputeBatch(int batchIndex) } } - [DllImport(mklDllName)] - public static extern unsafe void cblas_sgemm(Order order, byte transa, byte transb, int m, int n, int k, float alpha, float* a, int lda, float* b, int ldb, float beta, float* c, int ldc); + [LibraryImport(mklDllName)] + public static unsafe partial void cblas_sgemm(Order order, byte transa, byte transb, int m, int n, int k, float alpha, float* a, int lda, float* b, int ldb, float beta, float* c, int ldc); private static void GemmOp(BlasOp transA, BlasOp transB, float alpha, Tensor a, Tensor b, float beta, Tensor c) { @@ -1450,8 +1450,8 @@ private static void GemmOp(BlasOp transA, BlasOp transB, float alpha, Tensor a, } } - [DllImport(mklDllName)] - public static extern unsafe void cblas_sgemm_batch_strided(Order order, byte transa, byte transb, int m, int n, int k, float alpha, float* a, int lda, int stra, float* b, int ldb, int strb, float beta, float* c, int ldc, int stridec, int batch_size); + [LibraryImport(mklDllName)] + public static unsafe partial void cblas_sgemm_batch_strided(Order order, byte transa, byte transb, int m, int n, int k, float alpha, float* a, int lda, int stra, float* b, int ldb, int strb, float beta, float* c, int ldc, int stridec, int batch_size); private static void GemmOpBatch(BlasOp transA, BlasOp transB, float alpha, Tensor a, Tensor b, float beta, Tensor c) { diff --git a/TensorSharp.Core/Cpu/OpenBlasNative.cs b/TensorSharp.Core/Cpu/OpenBlasNative.cs index 157e20ab..481a8306 100644 --- a/TensorSharp.Core/Cpu/OpenBlasNative.cs +++ b/TensorSharp.Core/Cpu/OpenBlasNative.cs @@ -7,38 +7,44 @@ // // TensorSharp is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even the implied warranty of // MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the BSD-3-Clause License for more details. -using System.Runtime.InteropServices; +using System.Runtime.CompilerServices; +using System.Runtime.InteropServices; namespace TensorSharp.Cpu { // When used with 64bit openblas, this interface requires that it is compiled with 32-bit ints - public static class OpenBlasNative + public static partial class OpenBlasNative { private const string dll = "libopenblas.dll"; - private const CallingConvention cc = CallingConvention.Cdecl; - [DllImport(dll, CallingConvention = cc)] - public static extern unsafe void sgemm_(byte* transa, byte* transb, int* m, int* n, int* k, + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static unsafe partial void sgemm_(byte* transa, byte* transb, int* m, int* n, int* k, float* alpha, float* a, int* lda, float* b, int* ldb, float* beta, float* c, int* ldc); - [DllImport(dll, CallingConvention = cc)] - public static extern unsafe void dgemm_(byte* transa, byte* transb, int* m, int* n, int* k, + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static unsafe partial void dgemm_(byte* transa, byte* transb, int* m, int* n, int* k, double* alpha, double* a, int* lda, double* b, int* ldb, double* beta, double* c, int* ldc); - [DllImport(dll, CallingConvention = cc)] - public static extern unsafe void sgemv_(byte* trans, int* m, int* n, + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static unsafe partial void sgemv_(byte* trans, int* m, int* n, float* alpha, float* a, int* lda, float* x, int* incx, float* beta, float* y, int* incy); - [DllImport(dll, CallingConvention = cc)] - public static extern unsafe void dgemv_(byte* trans, int* m, int* n, + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static unsafe partial void dgemv_(byte* trans, int* m, int* n, double* alpha, double* a, int* lda, double* x, int* incx, double* beta, double* y, int* incy); - [DllImport(dll, CallingConvention = cc)] - public static extern unsafe float sdot_(int* n, float* x, int* incx, float* y, int* incy); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static unsafe partial float sdot_(int* n, float* x, int* incx, float* y, int* incy); - [DllImport(dll, CallingConvention = cc)] - public static extern unsafe double ddot_(int* n, double* x, int* incx, double* y, int* incy); + [LibraryImport(dll)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + public static unsafe partial double ddot_(int* n, double* x, int* incx, double* y, int* incy); } } From af405e7582f8bda27528ead0f107119ad4eb2935 Mon Sep 17 00:00:00 2001 From: craig-b <3112503+craig-b@users.noreply.github.com> Date: Sat, 15 Aug 2026 14:31:40 +0100 Subject: [PATCH 2/5] Migrate GGML backend P/Invokes to LibraryImport GgmlOps imports keep Cdecl via UnmanagedCallConv; the kernel32/libc memory-pool imports keep platform-default conventions. The UTF-8 string parameters move from per-parameter LPUTF8Str to StringMarshalling.Utf8. --- .../GgmlDeepSeek4Native.cs | 90 +- TensorSharp.Backends.GGML/GgmlMemoryPool.cs | 18 +- TensorSharp.Backends.GGML/GgmlNative.cs | 801 +++++++++++------- .../GgmlTensorParallel.cs | 61 +- 4 files changed, 579 insertions(+), 391 deletions(-) diff --git a/TensorSharp.Backends.GGML/GgmlDeepSeek4Native.cs b/TensorSharp.Backends.GGML/GgmlDeepSeek4Native.cs index 3aa5cd4e..e9c71b66 100644 --- a/TensorSharp.Backends.GGML/GgmlDeepSeek4Native.cs +++ b/TensorSharp.Backends.GGML/GgmlDeepSeek4Native.cs @@ -10,14 +10,14 @@ // (split) GGUF itself, places the layers across every visible GPU, owns the // DSV4 KV caches, and runs prefill/decode ubatches through ggml_backend_sched. using System; +using System.Runtime.CompilerServices; using System.Runtime.InteropServices; namespace TensorSharp.GGML { - public static class GgmlDeepSeek4Native + public static partial class GgmlDeepSeek4Native { private const string DllName = "GgmlOps"; - private const CallingConvention Conv = CallingConvention.Cdecl; static GgmlDeepSeek4Native() { @@ -26,58 +26,74 @@ static GgmlDeepSeek4Native() // Paths must cross as UTF-8: ggml_fopen decodes them as UTF-8 on Windows, // while CharSet.Ansi would marshal the active code page. - [DllImport(DllName, CallingConvention = Conv)] - private static extern IntPtr TSGgml_Dsv4LoadModel([MarshalAs(UnmanagedType.LPUTF8Str)] string ggufPath, + [LibraryImport(DllName, StringMarshalling = StringMarshalling.Utf8)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr TSGgml_Dsv4LoadModel(string ggufPath, int nGpu, int nCtx, int nUbatch, int nThreads, - int nCpuMoe, [MarshalAs(UnmanagedType.LPUTF8Str)] string backendName); + int nCpuMoe, string backendName); - [DllImport(DllName, CallingConvention = Conv)] - private static extern IntPtr TSGgml_Dsv4LoadModelDspark([MarshalAs(UnmanagedType.LPUTF8Str)] string ggufPath, + [LibraryImport(DllName, StringMarshalling = StringMarshalling.Utf8)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr TSGgml_Dsv4LoadModelDspark(string ggufPath, int nGpu, int nCtx, int nUbatch, - int nThreads, [MarshalAs(UnmanagedType.LPUTF8Str)] string dsparkPath, int nCpuMoe, - [MarshalAs(UnmanagedType.LPUTF8Str)] string backendName); + int nThreads, string dsparkPath, int nCpuMoe, + string backendName); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4DsparkBlockSize(IntPtr handle); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4DsparkBlockSize(IntPtr handle); - [DllImport(DllName, CallingConvention = Conv)] - private static extern unsafe int TSGgml_Dsv4ForwardSpec(IntPtr handle, int* tokens, int nTokens, float* logitsOut); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static unsafe partial int TSGgml_Dsv4ForwardSpec(IntPtr handle, int* tokens, int nTokens, float* logitsOut); - [DllImport(DllName, CallingConvention = Conv)] - private static extern unsafe int TSGgml_Dsv4DsparkDraft(IntPtr handle, int anchorToken, int* toksOut, float* confOut); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static unsafe partial int TSGgml_Dsv4DsparkDraft(IntPtr handle, int anchorToken, int* toksOut, float* confOut); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4Rewind(IntPtr handle, int nPast); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4Rewind(IntPtr handle, int nPast); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4VocabSize(IntPtr handle); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4VocabSize(IntPtr handle); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4CtxSize(IntPtr handle); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4CtxSize(IntPtr handle); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4NPast(IntPtr handle); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4NPast(IntPtr handle); - [DllImport(DllName, CallingConvention = Conv)] - private static extern unsafe int TSGgml_Dsv4Forward(IntPtr handle, int* tokens, int nTokens, float* logitsOut); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static unsafe partial int TSGgml_Dsv4Forward(IntPtr handle, int* tokens, int nTokens, float* logitsOut); - [DllImport(DllName, CallingConvention = Conv)] - private static extern void TSGgml_Dsv4Reset(IntPtr handle); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Dsv4Reset(IntPtr handle); - [DllImport(DllName, CallingConvention = Conv)] - private static extern void TSGgml_Dsv4Free(IntPtr handle); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Dsv4Free(IntPtr handle); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4SlotAlloc(IntPtr handle); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4SlotAlloc(IntPtr handle); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4SetActiveSlot(IntPtr handle, int slotId); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4SetActiveSlot(IntPtr handle, int slotId); - [DllImport(DllName, CallingConvention = Conv)] - private static extern int TSGgml_Dsv4SlotFree(IntPtr handle, int slotId); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Dsv4SlotFree(IntPtr handle, int slotId); - [DllImport(DllName, CallingConvention = Conv)] - private static extern unsafe int TSGgml_Dsv4ForwardBatchedDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static unsafe partial int TSGgml_Dsv4ForwardBatchedDecode( IntPtr handle, int n, int* slotIds, int* tokens, int* positions, float* logitsOut); /// diff --git a/TensorSharp.Backends.GGML/GgmlMemoryPool.cs b/TensorSharp.Backends.GGML/GgmlMemoryPool.cs index 2d19e4cb..5ec515b2 100644 --- a/TensorSharp.Backends.GGML/GgmlMemoryPool.cs +++ b/TensorSharp.Backends.GGML/GgmlMemoryPool.cs @@ -19,7 +19,7 @@ namespace TensorSharp.GGML /// GGML host-ptr buffers require aligned addresses, so use aligned allocations on every /// platform: 16KB on macOS for Metal shared memory, 32 bytes elsewhere for GGML CPU. /// - internal sealed class GgmlMemoryPool + internal sealed partial class GgmlMemoryPool { /// 16KB - Apple Silicon page size; required for Metal newBufferWithBytesNoCopy. private const int MetalPageSize = 16 * 1024; @@ -247,17 +247,17 @@ private static bool FreeVirtual(IntPtr ptr, nuint size) private const uint WindowsMemRelease = 0x8000; private const uint WindowsPageReadWrite = 0x04; - [DllImport("kernel32.dll", EntryPoint = "VirtualAlloc", ExactSpelling = true, SetLastError = true)] - private static extern IntPtr WindowsVirtualAlloc(IntPtr lpAddress, nuint dwSize, uint flAllocationType, uint flProtect); + [LibraryImport("kernel32.dll", EntryPoint = "VirtualAlloc", SetLastError = true)] + private static partial IntPtr WindowsVirtualAlloc(IntPtr lpAddress, nuint dwSize, uint flAllocationType, uint flProtect); - [DllImport("kernel32.dll", EntryPoint = "VirtualFree", ExactSpelling = true, SetLastError = true)] + [LibraryImport("kernel32.dll", EntryPoint = "VirtualFree", SetLastError = true)] [return: MarshalAs(UnmanagedType.Bool)] - private static extern bool WindowsVirtualFree(IntPtr lpAddress, UIntPtr dwSize, uint dwFreeType); + private static partial bool WindowsVirtualFree(IntPtr lpAddress, UIntPtr dwSize, uint dwFreeType); - [DllImport("libc", EntryPoint = "mmap", SetLastError = true)] - private static extern IntPtr UnixMmap(IntPtr addr, nuint length, int prot, int flags, int fd, IntPtr offset); + [LibraryImport("libc", EntryPoint = "mmap", SetLastError = true)] + private static partial IntPtr UnixMmap(IntPtr addr, nuint length, int prot, int flags, int fd, IntPtr offset); - [DllImport("libc", EntryPoint = "munmap", SetLastError = true)] - private static extern int UnixMunmap(IntPtr addr, nuint length); + [LibraryImport("libc", EntryPoint = "munmap", SetLastError = true)] + private static partial int UnixMunmap(IntPtr addr, nuint length); } } diff --git a/TensorSharp.Backends.GGML/GgmlNative.cs b/TensorSharp.Backends.GGML/GgmlNative.cs index 1def415f..cc71c179 100644 --- a/TensorSharp.Backends.GGML/GgmlNative.cs +++ b/TensorSharp.Backends.GGML/GgmlNative.cs @@ -12,6 +12,7 @@ using System.IO; using System.Linq; using System.Reflection; +using System.Runtime.CompilerServices; using System.Runtime.InteropServices; using System.Threading; @@ -924,7 +925,6 @@ internal enum GgmlIndexReductionOp internal static partial class GgmlNative { private const string DllName = "GgmlOps"; - private const CallingConvention CallingConventionType = CallingConvention.Cdecl; private static int s_windowsDependencySearchPathsInitialized; static GgmlNative() @@ -1064,34 +1064,43 @@ internal static bool SetNativeEnvironmentVariable(string name, string value, boo return TSGgml_SetNativeEnvironmentVariable(name, value, overwrite ? 1 : 0) != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern IntPtr TSGgml_GetLastError(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr TSGgml_GetLastError(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_IsMetalAvailable(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_IsMetalAvailable(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_SetNativeEnvironmentVariable( - [MarshalAs(UnmanagedType.LPUTF8Str)] string name, - [MarshalAs(UnmanagedType.LPUTF8Str)] string value, int overwrite); + [LibraryImport(DllName, StringMarshalling = StringMarshalling.Utf8)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_SetNativeEnvironmentVariable( + string name, + string value, int overwrite); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_CanInitializeBackend(int backendType); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_CanInitializeBackend(int backendType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_IsBackendAvailable(int backendType); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_IsBackendAvailable(int backendType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_SetVulkanDeviceIndex(int deviceIndex); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_SetVulkanDeviceIndex(int deviceIndex); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetVulkanDeviceCount(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetVulkanDeviceCount(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetVulkanDeviceDescription(int deviceIndex, byte[] description, int descriptionSize); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetVulkanDeviceDescription(int deviceIndex, byte[] description, int descriptionSize); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_AddmmF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_AddmmF32( GgmlTensorView2D result, GgmlTensorView2D src, GgmlTensorView2D m1, @@ -1099,8 +1108,9 @@ private static extern int TSGgml_AddmmF32( float beta, float alpha); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_AddmmQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_AddmmQuantF32( GgmlTensorView2D result, GgmlTensorView2D m1, IntPtr m2Data, @@ -1109,8 +1119,9 @@ private static extern int TSGgml_AddmmQuantF32( long m2Ne1, long m2RawBytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedRmsNormMatMulQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedRmsNormMatMulQuantF32( GgmlTensorView2D result, GgmlTensorView2D input, IntPtr normWeightData, @@ -1122,8 +1133,9 @@ private static extern int TSGgml_FusedRmsNormMatMulQuantF32( long m2Ne1, long m2RawBytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedMatMulQuantAddF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedMatMulQuantAddF32( GgmlTensorView2D residual, GgmlTensorView2D input, IntPtr m2Data, @@ -1133,11 +1145,13 @@ private static extern int TSGgml_FusedMatMulQuantAddF32( long m2RawBytes, int tpDegree, out IntPtr tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_ReleaseFusedMatmulAddTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_ReleaseFusedMatmulAddTpGraphs(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedFFNSwiGLUQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedFFNSwiGLUQuantF32( GgmlTensorView2D residual, GgmlTensorView2D input, IntPtr normWeightData, @@ -1156,8 +1170,9 @@ private static extern int TSGgml_FusedFFNSwiGLUQuantF32( int halfDim, int tpDegree, out IntPtr tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_ReleaseFusedFfnTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_ReleaseFusedFfnTpGraphs(); public static void ReleaseFusedFfnTpGraphs() { @@ -1165,8 +1180,9 @@ public static void ReleaseFusedFfnTpGraphs() catch (EntryPointNotFoundException) { } } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedFFNActProjectQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedFFNActProjectQuantF32( GgmlTensorView2D output, GgmlTensorView2D input, IntPtr normWeightData, @@ -1185,24 +1201,27 @@ private static extern int TSGgml_FusedFFNActProjectQuantF32( int halfDim, int actType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedRmsNormResidualAddF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedRmsNormResidualAddF32( GgmlTensorView2D residual, GgmlTensorView2D input, IntPtr normWeightData, int normWeightCount, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedPleBlockQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedPleBlockQuantF32( GgmlTensorView2D residual, GgmlTensorView2D perLayerInput, IntPtr inpGateData, int inpGateGgmlType, long inpGateNe0, long inpGateNe1, long inpGateRawBytes, IntPtr projData, int projGgmlType, long projNe0, long projNe1, long projRawBytes, IntPtr postNormData, int postNormCount, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedOutProjNormRouterQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedOutProjNormRouterQuantF32( GgmlTensorView2D residual, GgmlTensorView2D input, IntPtr outProjData, int outProjType, long outNe0, long outNe1, long outBytes, IntPtr normData, int normCount, float eps, @@ -1210,8 +1229,9 @@ private static extern int TSGgml_FusedOutProjNormRouterQuantF32( IntPtr routerData, int routerType, long routerNe0, long routerNe1, long routerBytes, GgmlTensorView2D routerOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedVisionMLPF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedVisionMLPF32( GgmlTensorView2D hidden, IntPtr lnW, IntPtr lnB, int lnDim, float eps, IntPtr upW, int upNe0, int upNe1, long upBytes, @@ -1219,12 +1239,14 @@ private static extern int TSGgml_FusedVisionMLPF32( IntPtr downW, int downNe0, int downNe1, long downBytes, IntPtr downB, int downBDim); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_MuseGlimmerVisionBlockQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_MuseGlimmerVisionBlockQuantF32( in GgmlMuseGlimmerVisionBlockArgs args); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedOutProjFFNQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedOutProjFFNQuantF32( GgmlTensorView2D residual, GgmlTensorView2D input, IntPtr outProjData, int outProjType, long outNe0, long outNe1, long outRawBytes, IntPtr ffnNormData, int ffnNormCount, float eps, @@ -1232,8 +1254,9 @@ private static extern int TSGgml_FusedOutProjFFNQuantF32( IntPtr dnData, int dnType, long dnNe0, long dnNe1, long dnRawBytes, int halfDim); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedVisionAttentionF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedVisionAttentionF32( GgmlTensorView2D hidden, IntPtr lnW, IntPtr lnB, int lnDim, float eps, IntPtr qkvW, int qkvNe0, int qkvNe1, long qkvBytes, @@ -1244,8 +1267,9 @@ private static extern int TSGgml_FusedVisionAttentionF32( int numPatches, int numHeads, int headDim, int halfDim, float attnScale); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35VisionEncoderF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35VisionEncoderF32( GgmlTensorView2D hidden, int blockCount, float eps, float attnScale, int numPatches, int numHeads, int headDim, int halfDim, @@ -1262,8 +1286,9 @@ private static extern int TSGgml_Qwen35VisionEncoderF32( int upNe0, int upNe1, long upBytes, int upBDim, int downNe0, int downNe1, long downBytes, int downBDim); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedGemma4VisionBlockF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedGemma4VisionBlockF32( GgmlTensorView2D hidden, float eps, IntPtr ln1W, IntPtr qW, int qNe0, int qNe1, long qBytes, @@ -1281,8 +1306,9 @@ private static extern int TSGgml_FusedGemma4VisionBlockF32( IntPtr clamps, int numPatches, int numHeads, int headDim); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetRowsQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetRowsQuantF32( GgmlTensorView2D result, IntPtr srcData, int srcGgmlType, @@ -1291,8 +1317,9 @@ private static extern int TSGgml_GetRowsQuantF32( long srcRawBytes, GgmlContiguousTensor indices); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_MoEExpertsForwardF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_MoEExpertsForwardF32( GgmlTensorView2D result, GgmlTensorView2D input, int numExperts, @@ -1308,8 +1335,9 @@ private static extern int TSGgml_MoEExpertsForwardF32( long downRawBytesEach, float[] routeWeights); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_MoEExpertsSwiGLUForwardF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_MoEExpertsSwiGLUForwardF32( GgmlTensorView2D result, GgmlTensorView2D input, int numExperts, @@ -1330,8 +1358,9 @@ private static extern int TSGgml_MoEExpertsSwiGLUForwardF32( long downRawBytesEach, float[] routeWeights); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_MoEExpertsSwiGLUResidualF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_MoEExpertsSwiGLUResidualF32( GgmlTensorView2D residual, GgmlTensorView2D input, int numExperts, @@ -1369,8 +1398,9 @@ private static extern int TSGgml_MoEExpertsSwiGLUResidualF32( long sharedDownRawBytes, float sharedScalar); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_AddmmQuantBatchF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_AddmmQuantBatchF32( GgmlTensorView2D result, GgmlTensorView2D m1, IntPtr m2Data, @@ -1381,8 +1411,9 @@ private static extern int TSGgml_AddmmQuantBatchF32( long[] weightOffsets, long[] weightNe1Arr); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_AddmmBatchF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_AddmmBatchF32( GgmlTensorView3D result, GgmlTensorView3D src, GgmlTensorView3D m1, @@ -1390,8 +1421,9 @@ private static extern int TSGgml_AddmmBatchF32( float beta, float alpha); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_MulMatIdF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_MulMatIdF32( GgmlTensorView3D result, GgmlTensorView3D expertWeights, GgmlTensorView3D input, @@ -1399,8 +1431,9 @@ private static extern int TSGgml_MulMatIdF32( int idsRows, int idsCols); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_AddIdF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_AddIdF32( GgmlTensorView3D result, GgmlTensorView3D src, GgmlTensorView2D bias, @@ -1408,28 +1441,32 @@ private static extern int TSGgml_AddIdF32( int idsRows, int idsCols); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_ReduceLastDimF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_ReduceLastDimF32( int op, GgmlTensorView4D result, GgmlTensorView4D src); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_IndexReductionF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_IndexReductionF32( int op, GgmlTensorView4D result, GgmlTensorView4D src); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_SoftmaxF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_SoftmaxF32( GgmlTensorView4D result, GgmlTensorView4D src); // In-place softmax with causal+SWA mask and optional attention sinks. // Replaces the GptOss CPU softmax-with-sinks loop. See native side: // attention_softmax_with_sinks_f32_impl in ggml_ops_norm_attn.cpp. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_AttentionSoftmaxWithSinksF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_AttentionSoftmaxWithSinksF32( GgmlTensorView3D scores, IntPtr sinksData, // float* [num_heads], or IntPtr.Zero for no sinks int numHeads, @@ -1443,8 +1480,9 @@ private static extern int TSGgml_AttentionSoftmaxWithSinksF32( // Collapses an entire layer's MoE forward (gate + up + SwiGLU + down + // expert weighting + aggregation) into one GGML graph dispatch. // See native side: TSGgml_MoEFFNPrefillSwiGLUQuantF32 in ggml_ops_moe.cpp. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_MoEFFNPrefillSwiGLUQuantF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_MoEFFNPrefillSwiGLUQuantF32( IntPtr hiddenIn, IntPtr hiddenOut, int seqLen, @@ -1471,8 +1509,9 @@ private static extern int TSGgml_MoEFFNPrefillSwiGLUQuantF32( // TSGgml_MoEFFNPrefillSwiGLUQuantF32 ABI but adds the residual buffer, // the post_ffw_norm_2 weight, and an RMSNorm epsilon. // See native side: TSGgml_Gemma4MoEGEGLUResidualF32 in ggml_ops_moe.cpp. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4MoEGEGLUResidualF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4MoEGEGLUResidualF32( IntPtr hiddenIn, IntPtr residualInOut, // float* [seqLen, hiddenDim] - dense FFN result; kernel adds normed MoE output to it in place IntPtr postNormW, // float* [hiddenDim] - post_ffw_norm_2.weight @@ -1495,8 +1534,9 @@ private static extern int TSGgml_Gemma4MoEGEGLUResidualF32( float oaiLimit, int runOnCpu); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_ScaledDotProductAttentionF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_ScaledDotProductAttentionF32( GgmlTensorView4D result, GgmlTensorView4D query, GgmlTensorView4D key, @@ -1505,23 +1545,26 @@ private static extern int TSGgml_ScaledDotProductAttentionF32( int hasMask, float scale); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_SoftmaxGradF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_SoftmaxGradF32( GgmlTensorView4D result, GgmlTensorView4D adj, GgmlTensorView4D val, int addGrad); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_CrossEntropyLossF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_CrossEntropyLossF32( out float lossValue, GgmlTensorView4D probs, GgmlContiguousTensor targetIndices, float smooth, float labelSmooth); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_CrossEntropyLossBackwardF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_CrossEntropyLossBackwardF32( GgmlTensorView4D grad, GgmlTensorView4D probs, GgmlContiguousTensor targetIndices, @@ -1530,8 +1573,9 @@ private static extern int TSGgml_CrossEntropyLossBackwardF32( float labelSmooth, int addGrad); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_AdamF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_AdamF32( GgmlContiguousTensor weight, GgmlContiguousTensor gradient, GgmlContiguousTensor v, @@ -1545,8 +1589,9 @@ private static extern int TSGgml_AdamF32( int iter, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_TransformerLayerDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_TransformerLayerDecode( IntPtr hiddenData, int hiddenSize, IntPtr attnNormData, IntPtr qkvData, int qkvType, long qkvNe0, long qkvNe1, long qkvBytes, @@ -1563,8 +1608,9 @@ private static extern int TSGgml_TransformerLayerDecode( int intermediateSize, int ropeMode, int kvCacheType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4LayerPrefill( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4LayerPrefill( IntPtr hiddenData, int hiddenSize, int seqLen, IntPtr attnNormW, IntPtr qkvW, int qkvType, long qkvNe0, long qkvNe1, long qkvBytes, @@ -1649,24 +1695,27 @@ public static void Gemma4LayerPrefill( kvCacheType), "gemma4_layer_prefill"); } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedPrefillAttentionF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedPrefillAttentionF32( IntPtr qData, IntPtr kData, IntPtr vData, IntPtr outData, int numHeads, int numKvHeads, int headDim, int seqLen, int kvLen, int maskStartPos, int slidingWindow, float scale, int inputFormat); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedPrefillAttentionF16KV( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedPrefillAttentionF16KV( IntPtr qData, IntPtr kData, IntPtr vData, IntPtr outData, int numHeads, int numKvHeads, int headDim, int seqLen, int kvLen, int kvCacheLen, int maskStartPos, int slidingWindow, float scale); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FlashAttnDecodeF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FlashAttnDecodeF32( IntPtr qData, IntPtr kData, IntPtr vData, IntPtr kCacheData, IntPtr vCacheData, IntPtr outData, @@ -1674,8 +1723,9 @@ private static extern int TSGgml_FlashAttnDecodeF32( int maxSeqLen, int position, float scale, int kvCacheType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_PagedAttentionForward( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_PagedAttentionForward( IntPtr qData, IntPtr pagedKData, IntPtr pagedVData, @@ -1694,8 +1744,9 @@ private static extern int TSGgml_PagedAttentionForward( int slidingWindow, float scale); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_PagedAttentionForwardWithSinks( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_PagedAttentionForwardWithSinks( IntPtr qData, IntPtr pagedKData, IntPtr pagedVData, @@ -1720,8 +1771,9 @@ private static extern int TSGgml_PagedAttentionForwardWithSinks( // instead of round-tripping through host arrays + ggml_backend_synchronize. // Eliminates the per-layer queue drain that GetElementsAsFloat would // otherwise force. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_PagedAttentionForwardDevice( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_PagedAttentionForwardDevice( IntPtr qData, IntPtr pagedKData, IntPtr pagedVData, @@ -1740,8 +1792,9 @@ private static extern int TSGgml_PagedAttentionForwardDevice( int slidingWindow, float scale); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_PagedAttentionForwardDeviceWithSinks( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_PagedAttentionForwardDeviceWithSinks( IntPtr qData, IntPtr pagedKData, IntPtr pagedVData, @@ -1761,8 +1814,9 @@ private static extern int TSGgml_PagedAttentionForwardDeviceWithSinks( float scale, IntPtr sinksData); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35AttentionLayerDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35AttentionLayerDecode( IntPtr residualData, int hiddenSize, IntPtr attnNormData, IntPtr qkvData, int qkvType, long qkvNe0, long qkvNe1, long qkvBytes, @@ -1774,8 +1828,9 @@ private static extern int TSGgml_Qwen35AttentionLayerDecode( float eps, float ropeBase, float ropeFreqScale, int ropeMode, int kvCacheType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GptOssAttentionLayerPrefill( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GptOssAttentionLayerPrefill( IntPtr hiddenData, int hiddenSize, int seqLen, IntPtr attnNormW, IntPtr qkvW, int qkvType, long qkvNe0, long qkvNe1, long qkvBytes, @@ -1797,8 +1852,9 @@ private static extern int TSGgml_GptOssAttentionLayerPrefill( int kvCacheType, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35AttentionLayerPrefill( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35AttentionLayerPrefill( IntPtr hiddenData, int hiddenSize, int seqLen, IntPtr attnNormW, IntPtr qkvW, int qkvType, long qkvNe0, long qkvNe1, long qkvBytes, @@ -1813,8 +1869,9 @@ private static extern int TSGgml_Qwen35AttentionLayerPrefill( float eps, int tpDegree, out IntPtr tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Qwen35ReleaseAttentionTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Qwen35ReleaseAttentionTpGraphs(); public static void Qwen35ReleaseAttentionTpGraphs() { @@ -1898,8 +1955,9 @@ public static void GptOssAttentionLayerPrefill( eps), "gpt_oss_attention_layer_prefill"); } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_TransformerModelDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_TransformerModelDecode( IntPtr hiddenData, int hiddenSize, int numLayers, IntPtr[] attnNormArr, IntPtr[] qkvArr, IntPtr[] qNormArr, IntPtr[] kNormArr, IntPtr[] oArr, IntPtr[] ffnNormArr, IntPtr[] guArr, IntPtr[] downArr, @@ -1921,8 +1979,9 @@ private static extern int TSGgml_TransformerModelDecode( int intermediateSize, int ropeMode, int kvCacheType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4ModelDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4ModelDecode( IntPtr hiddenData, int hiddenSize, int numLayers, IntPtr[] attnNormArr, IntPtr[] qkvArr, IntPtr[] qNormArr, IntPtr[] kNormArr, IntPtr[] oArr, IntPtr[] postAttnNormArr, @@ -1958,8 +2017,9 @@ private static extern int TSGgml_Gemma4ModelDecode( int tpDegree, out IntPtr tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DFlashInject( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DFlashInject( float[] featRows, int featureSize, int nRows, long[] ringRowsIdx, int[] positions, int numLayers, int hiddenSize, int headDim, int numKvHeads, int ringRows, @@ -1972,8 +2032,9 @@ private static extern int TSGgml_DFlashInject( IntPtr[] ringKArr, IntPtr[] ringVArr, int ringDtype); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DFlashDraftBlock( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DFlashDraftBlock( int[] blockIds, int blockLen, int[] positions, int numLayers, int hiddenSize, int headDim, int numHeads, int numKvHeads, int ringRows, float eps, float ropeBase, float ropeFreqScale, float kqScale, @@ -1994,8 +2055,9 @@ private static extern int TSGgml_DFlashDraftBlock( IntPtr lmHeadData, int lmHeadType, long lmHeadNe0, long lmHeadNe1, long lmHeadBytes, int vocabSize, int[] idsOut, float[] confOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_DFlashResetCaches(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_DFlashResetCaches(); /// DFlash PASS A+B in one graph. False = declined, caller falls back. public static bool DFlashInject( @@ -2058,8 +2120,9 @@ public static bool DFlashDraftBlock( /// Drop the persistent DFlash graphs (ring reallocation / KV reset). public static void DFlashResetCaches() => TSGgml_DFlashResetCaches(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_MuseGlimmerModelForward( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_MuseGlimmerModelForward( IntPtr hiddenData, int hiddenSize, int nTokens, int numLayers, IntPtr[] attnNormArr, IntPtr[] qArr, IntPtr[] kArr, IntPtr[] vArr, IntPtr[] gateArr, @@ -2090,11 +2153,13 @@ private static extern int TSGgml_MuseGlimmerModelForward( int[] tokenIds, int allLogitsRows, int tpDegree, [In, Out] IntPtr[] tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_MuseGlimmerResetDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_MuseGlimmerResetDecodeCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_MuseGlimmerReleaseTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_MuseGlimmerReleaseTpGraphs(); /// /// Whole-model Muse-Glimmer forward in a single GGML graph. nTokens == 1 uses @@ -2182,8 +2247,9 @@ public static void MuseGlimmerReleaseTpGraphs() catch (EntryPointNotFoundException) { } } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4ModelDecodeBatched( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4ModelDecodeBatched( IntPtr hiddenData, int hiddenSize, int numLayers, int nSeqs, IntPtr[] attnNormArr, IntPtr[] qkvArr, IntPtr[] qNormArr, IntPtr[] kNormArr, IntPtr[] oArr, IntPtr[] postAttnNormArr, @@ -2206,8 +2272,9 @@ private static extern int TSGgml_Gemma4ModelDecodeBatched( IntPtr lmHeadData, int lmHeadType, long lmHeadNe0, long lmHeadNe1, long lmHeadBytes, IntPtr finalNormData, float logitSoftcap); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4ModelVerify( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4ModelVerify( IntPtr hiddenData, int hiddenSize, int numLayers, int numTokens, IntPtr[] attnNormArr, IntPtr[] qkvArr, IntPtr[] qNormArr, IntPtr[] kNormArr, IntPtr[] oArr, IntPtr[] postAttnNormArr, @@ -2240,8 +2307,9 @@ private static extern int TSGgml_Gemma4ModelVerify( IntPtr pleProjNormData, int tpDegree, out IntPtr tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4DraftStep( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4DraftStep( int token, IntPtr hPrev, int fixedPos, int backbone, int draftHidden, int numDLayers, int numHeads, int vocab, float eps, int kvCacheType, @@ -2309,24 +2377,28 @@ public static unsafe bool Gemma4DraftStep( return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4MoELayerDecode(in Gemma4MoELayerDecodeArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4MoELayerDecode(in Gemma4MoELayerDecodeArgs desc); public static void Gemma4MoELayerDecode(in Gemma4MoELayerDecodeArgs desc) { CheckResult(TSGgml_Gemma4MoELayerDecode(in desc), nameof(TSGgml_Gemma4MoELayerDecode)); } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DiffusionDecodeLayer(in DiffusionDecodeLayerArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DiffusionDecodeLayer(in DiffusionDecodeLayerArgs desc); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_QwenImageModMlp(in QwenImageModMlpArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_QwenImageModMlp(in QwenImageModMlpArgs desc); public static bool TryQwenImageModMlp(in QwenImageModMlpArgs desc) => TSGgml_QwenImageModMlp(in desc) != 0; - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_QwenImageJointAttn(in QwenImageJointAttnArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_QwenImageJointAttn(in QwenImageJointAttnArgs desc); public static bool TryQwenImageJointAttn(in QwenImageJointAttnArgs desc) { @@ -2336,8 +2408,9 @@ public static bool TryQwenImageJointAttn(in QwenImageJointAttnArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_QwenImageBlock(in QwenImageBlockArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_QwenImageBlock(in QwenImageBlockArgs desc); public static bool TryQwenImageBlock(in QwenImageBlockArgs desc) { @@ -2347,8 +2420,9 @@ public static bool TryQwenImageBlock(in QwenImageBlockArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_QwenImageBlockCfg(in QwenImageBlockArgs condDesc, in QwenImageBlockArgs negDesc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_QwenImageBlockCfg(in QwenImageBlockArgs condDesc, in QwenImageBlockArgs negDesc); // CFG-batched block: both true-CFG branches in one dispatch sharing the weights. public static bool TryQwenImageBlockCfg(in QwenImageBlockArgs condDesc, in QwenImageBlockArgs negDesc) @@ -2359,8 +2433,9 @@ public static bool TryQwenImageBlockCfg(in QwenImageBlockArgs condDesc, in QwenI return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_QwenImageForward(in QwenImageForwardArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_QwenImageForward(in QwenImageForwardArgs desc); // Whole 60-block DiT forward in one resident-weight graph (in-graph modulation). public static bool TryQwenImageForward(in QwenImageForwardArgs desc) @@ -2371,8 +2446,9 @@ public static bool TryQwenImageForward(in QwenImageForwardArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_WanT5Encode(in WanT5EncodeArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_WanT5Encode(in WanT5EncodeArgs desc); // Whole UMT5-XXL encoder forward in one resident-weight graph. public static bool TryWanT5Encode(in WanT5EncodeArgs desc) @@ -2383,8 +2459,9 @@ public static bool TryWanT5Encode(in WanT5EncodeArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_WanDitForward(in WanDitForwardArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_WanDitForward(in WanDitForwardArgs desc); // Whole Wan DiT forward (one denoising-step velocity prediction) in one // resident-weight graph; persistent + CUDA-graph-captured per shape on CUDA. @@ -2396,8 +2473,9 @@ public static bool TryWanDitForward(in WanDitForwardArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_WanVaeDecode(in WanVaeDecodeArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_WanVaeDecode(in WanVaeDecodeArgs desc); // Whole Wan 2.1 video VAE decode (chunked causal 3D decoder) in one graph. public static bool TryWanVaeDecode(in WanVaeDecodeArgs desc) @@ -2408,8 +2486,9 @@ public static bool TryWanVaeDecode(in WanVaeDecodeArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_WanVaeEncode(in WanVaeEncodeArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_WanVaeEncode(in WanVaeEncodeArgs desc); // Whole Wan video VAE encode (chunked causal 3D encoder) in one graph. public static bool TryWanVaeEncode(in WanVaeEncodeArgs desc) @@ -2420,8 +2499,9 @@ public static bool TryWanVaeEncode(in WanVaeEncodeArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_QwenImageSetOffload(int on); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_QwenImageSetOffload(int on); // CPU-offload mode for the Qwen-Image DiT kernels: disables the persistent / // CUDA-graph-captured entries (whose one-time resident weight upload is their @@ -2429,8 +2509,9 @@ public static bool TryWanVaeEncode(in WanVaeEncodeArgs desc) // call. Set per request by the pipeline with the device-copy residency budget. public static void QwenImageSetOffload(bool on) => TSGgml_QwenImageSetOffload(on ? 1 : 0); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Conv2d(in Conv2dArgs desc); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Conv2d(in Conv2dArgs desc); public static bool TryConv2d(in Conv2dArgs desc) { @@ -2448,8 +2529,9 @@ public static bool TryDiffusionDecodeLayer(in DiffusionDecodeLayerArgs desc) return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DiffusionLmHead( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DiffusionLmHead( IntPtr hidden, int hiddenSize, int canvasLen, IntPtr outputNormW, IntPtr lmHeadW, int lmHeadType, long lmHeadNe0, long lmHeadNe1, long lmHeadBytes, @@ -2467,8 +2549,9 @@ public static bool TryDiffusionLmHead( return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DiffusionLmHeadSample( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DiffusionLmHeadSample( IntPtr hidden, int hiddenSize, int canvasLen, IntPtr outputNormW, IntPtr lmHeadW, int lmHeadType, long lmHeadNe0, long lmHeadNe1, long lmHeadBytes, @@ -2495,8 +2578,9 @@ public static bool TryDiffusionLmHeadSample( return r != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DiffusionModelDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DiffusionModelDecode( [In] DiffusionDecodeLayerArgs[] layers, int numLayers, IntPtr hidden, int hiddenSize, int canvasLen, int promptLen, IntPtr outputNormW, @@ -2520,8 +2604,9 @@ public static bool TryDiffusionModelDecode( // Model-wide MoE decode: the whole transformer as one graph/token. // GPT-OSS whole-model decode: all layers + MoE + folded final norm/LM head // in ONE graph dispatch per token (see ggml_ops_gptoss_decode.cpp). - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GptOssModelDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GptOssModelDecode( [In] GptOssLayerDecodeArgs[] layers, int numLayers, IntPtr hidden, int hiddenSize, int position, IntPtr logits, int vocabSize, @@ -2542,8 +2627,9 @@ public static bool TryGptOssModelDecode( // GPT-OSS whole-model prefill: N tokens through every layer + MoE + // folded final norm/LM head in ONE graph (see ggml_ops_gptoss_prefill.cpp). - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GptOssModelPrefill( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GptOssModelPrefill( [In] GptOssLayerDecodeArgs[] layers, int numLayers, IntPtr hidden, int hiddenSize, int numTokens, int startPos, IntPtr logits, int vocabSize, @@ -2564,8 +2650,9 @@ public static bool TryGptOssModelPrefill( => TSGgml_GptOssModelPrefill(layers, numLayers, hidden, hiddenSize, numTokens, startPos, logits, vocabSize, lmHead, lmHeadType, lmHeadNe0, lmHeadNe1, lmHeadBytes, finalNorm) != 0; - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_GptOssResetDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_GptOssResetDecodeCache(); /// /// Drops every cached GPT-OSS whole-model decode graph. Call before a @@ -2574,8 +2661,9 @@ public static bool TryGptOssModelPrefill( /// public static void GptOssResetDecodeCache() => TSGgml_GptOssResetDecodeCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GptOssSyncKvCacheToHost( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GptOssSyncKvCacheToHost( IntPtr kCache, IntPtr vCache, int cacheSize, int rows); /// @@ -2589,8 +2677,9 @@ public static void GptOssSyncKvCacheToHost(IntPtr kCache, IntPtr vCache, int cac // `layers` is one Gemma4MoELayerDecodeArgs per layer (blittable, marshalled // as a contiguous TSGgmlGemma4MoELayerDesc array). hidden/position come from // the explicit params; the per-element Hidden/Position fields are ignored. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4MoEModelDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4MoEModelDecode( [In] Gemma4MoELayerDecodeArgs[] layers, int numLayers, IntPtr hidden, int hiddenSize, int position, IntPtr logits, int vocabSize, @@ -2623,8 +2712,9 @@ public static void Gemma4MoEModelDecode(Gemma4MoELayerDecodeArgs[] layers, int n // TRUE token-batched MoE decode: N concurrent sequences, one token each, in // one captured graph. Reuses the per-layer descriptor array for weights; // KV caches are per-(layer,seq) [layer*nSeqs+seq]; positions per seq. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4MoEModelDecodeBatched( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4MoEModelDecodeBatched( [In] Gemma4MoELayerDecodeArgs[] layers, int numLayers, int nSeqs, IntPtr hidden, IntPtr[] kCacheArr, IntPtr[] vCacheArr, @@ -2644,23 +2734,26 @@ public static bool Gemma4MoEModelDecodeBatched(Gemma4MoELayerDecodeArgs[] layers return rc != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Gemma4ResetMoEBatchedDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Gemma4ResetMoEBatchedDecodeCache(); public static void Gemma4ResetMoEBatchedDecodeCache() => TSGgml_Gemma4ResetMoEBatchedDecodeCache(); // Model-wide MoE multi-token verify: the whole MoE transformer over N tokens // as one graph. Reuses the same descriptor array as the decode; start_pos + // num_tokens are explicit. Returns 0 (false) when the kernel cannot handle // the shape so the caller falls back to the per-op verify. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Gemma4MoEModelVerify( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Gemma4MoEModelVerify( [In] Gemma4MoELayerDecodeArgs[] layers, int numLayers, IntPtr hidden, int hiddenSize, int startPos, int numTokens, byte[] mmIsExcept, int tpDegree, out IntPtr tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Gemma4MoEReleaseVerifyTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Gemma4MoEReleaseVerifyTpGraphs(); public static void Gemma4MoEReleaseVerifyTpGraphs() { @@ -2681,8 +2774,9 @@ public static bool Gemma4MoEModelVerify(Gemma4MoELayerDecodeArgs[] layers, int n // Qwen3.5/3.6 full-model decode: the whole hybrid transformer (full-attention // + GatedDeltaNet recurrent layers + per-layer FFN) as one graph/token. // Returns 0 when it cannot handle the shape so the caller falls back to per-op. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35ModelDecode( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35ModelDecode( [In] Qwen35LayerDecodeArgs[] layers, int numLayers, [MarshalAs(UnmanagedType.Bool)] bool reseedState, IntPtr hidden, int hiddenSize, int position, @@ -2697,8 +2791,9 @@ private static extern int TSGgml_Qwen35ModelDecode( IntPtr finalNorm, int tpDegree, [In, Out] IntPtr[] tpPlanOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35ModelDecodeToken( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35ModelDecodeToken( [In] Qwen35LayerDecodeArgs[] layers, int numLayers, [MarshalAs(UnmanagedType.Bool)] bool reseedState, int tokenId, @@ -2715,16 +2810,19 @@ private static extern int TSGgml_Qwen35ModelDecodeToken( IntPtr lmHead, int lmHeadType, long lmHeadNe0, long lmHeadNe1, long lmHeadBytes, IntPtr finalNorm); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Qwen35ResetDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Qwen35ResetDecodeCache(); public static void Qwen35ResetDecodeCache() => TSGgml_Qwen35ResetDecodeCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Gemma4ResetDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Gemma4ResetDecodeCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Gemma4ReleaseVerifyTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Gemma4ReleaseVerifyTpGraphs(); public static void Gemma4ReleaseVerifyTpGraphs() { @@ -2734,18 +2832,21 @@ public static void Gemma4ReleaseVerifyTpGraphs() public static void Gemma4ResetDecodeCache() => TSGgml_Gemma4ResetDecodeCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Gemma4ResetBatchedDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Gemma4ResetBatchedDecodeCache(); public static void Gemma4ResetBatchedDecodeCache() => TSGgml_Gemma4ResetBatchedDecodeCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Gemma4MoEResetDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Gemma4MoEResetDecodeCache(); public static void Gemma4MoEResetDecodeCache() => TSGgml_Gemma4MoEResetDecodeCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Qwen35ResetVerifyCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Qwen35ResetVerifyCache(); public static void Qwen35ResetVerifyCache() => TSGgml_Qwen35ResetVerifyCache(); @@ -2755,8 +2856,9 @@ public static void Gemma4ReleaseVerifyTpGraphs() // when it cannot handle the shape so the caller falls back to op-by-op. // padKv = fixed per-seq gather length (round_up(maxSeqLen, stride)); gatherIdx is // [nSeqs*padKv] (real slots then pad), seqLens [nSeqs] drives the per-seq attn mask. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35ModelDecodeBatched( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35ModelDecodeBatched( [In] Qwen35LayerDecodeArgs[] layers, int numLayers, IntPtr hidden, int hiddenSize, int nTokens, int nSeqs, IntPtr positions, IntPtr slotMapping, @@ -2768,8 +2870,9 @@ private static extern int TSGgml_Qwen35ModelDecodeBatched( int numExperts, int numExpertsUsed, int expertFf, int sharedFf, int normTopk, float expertWeightsScale); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Qwen35ResetBatchedDecodeCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Qwen35ResetBatchedDecodeCache(); public static void Qwen35ResetBatchedDecodeCache() => TSGgml_Qwen35ResetBatchedDecodeCache(); @@ -2865,8 +2968,9 @@ public static bool Qwen35ModelDecodeToken( // [hidden, N] (normedOut, for the MTP draft head). GDN state advances from // each layer's ConvStateIn/DeltaStateIn to ConvStateOut/DeltaStateOut. // Returns 0 when it cannot handle the shape so the caller falls back to per-op. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35ModelVerify( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35ModelVerify( [In] Qwen35LayerDecodeArgs[] layers, int numLayers, IntPtr hidden, int hiddenSize, int startPos, int numTokens, int numHeads, int numKvHeads, int headDim, int cacheSize, @@ -2910,8 +3014,9 @@ public static bool Qwen35ModelVerify( tpDegree, tpPlanOut) != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Qwen35ReleaseVerifyTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Qwen35ReleaseVerifyTpGraphs(); public static void Qwen35ReleaseVerifyTpGraphs() { @@ -2919,8 +3024,9 @@ public static void Qwen35ReleaseVerifyTpGraphs() catch (EntryPointNotFoundException) { } } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35RecurrentLayerPrefill( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35RecurrentLayerPrefill( IntPtr hiddenData, int hiddenSize, int n, IntPtr attnNormW, IntPtr gdnQkvW, int gdnQkvType, long gdnQkvNe0, long gdnQkvNe1, long gdnQkvBytes, @@ -2960,8 +3066,9 @@ public static bool Qwen35RecurrentLayerPrefill( convKernel, headKDim, headVDim, numKHeads, numVHeads, eps) != 0; } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_Qwen35GdnLayerTP( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_Qwen35GdnLayerTP( IntPtr hiddenData, int hiddenSize, int n, IntPtr attnNormW, IntPtr inprojW, int inprojType, long inprojNe0, long inprojNe1, long inprojBytes, @@ -2999,14 +3106,16 @@ public static void Qwen35GdnLayerTP( convKernel, eps), "qwen35_gdn_layer_tp"); } - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Qwen35GdnDropTpGraphs(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Qwen35GdnDropTpGraphs(); /// Free every cached per-rank TP GatedDeltaNet graph. public static void Qwen35GdnDropTpGraphs() => TSGgml_Qwen35GdnDropTpGraphs(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GatedDeltaNetChunkedF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GatedDeltaNetChunkedF32( GgmlTensorView3D q, GgmlTensorView3D k, GgmlTensorView3D v, @@ -3023,8 +3132,9 @@ private static extern int TSGgml_GatedDeltaNetChunkedF32( // Mirrors NemoMamba2BatchedSeqDesc in ggml_ops_mamba2.cpp; same 32-byte // POD layout on 64-bit (two ints, two padding ints, two pointers). - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_NemotronMamba2BatchedStepF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_NemotronMamba2BatchedStepF32( int numSeqs, [In, Out] NemoMamba2BatchedSeqDesc[] seqs, int numTokens, @@ -3045,8 +3155,9 @@ private static extern int TSGgml_NemotronMamba2BatchedStepF32( float eps, IntPtr outBatched); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GatedDeltaNetBatchedStepF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GatedDeltaNetBatchedStepF32( int numSeqs, [In, Out] GdnBatchedSeqDesc[] seqs, int numTokens, @@ -3069,8 +3180,9 @@ private static extern int TSGgml_GatedDeltaNetBatchedStepF32( float eps, IntPtr gatedOut); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_NemotronMamba2PrefillF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_NemotronMamba2PrefillF32( GgmlTensorView2D projected, GgmlTensorView2D hiddenOut, IntPtr convStateData, @@ -3091,8 +3203,9 @@ private static extern int TSGgml_NemotronMamba2PrefillF32( int dConv, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_NemotronMamba2DecodeF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_NemotronMamba2DecodeF32( ulong stateKey, GgmlTensorView2D projected, GgmlTensorView2D hiddenOut, @@ -3116,157 +3229,195 @@ private static extern int TSGgml_NemotronMamba2DecodeF32( int dConv, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_NemotronMamba2DecodeClear(ulong modelKey); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_NemotronMamba2DecodeClear(ulong modelKey); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern IntPtr TSGgml_AlignedAlloc(UIntPtr size); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr TSGgml_AlignedAlloc(UIntPtr size); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_AlignedFree(IntPtr ptr); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_AlignedFree(IntPtr ptr); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_ClearHostBufferCache(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_ClearHostBufferCache(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_Shutdown(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_Shutdown(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_ReleaseReuseComputeBuffers(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_ReleaseReuseComputeBuffers(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_InvalidateHostBuffer(IntPtr ptr); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_InvalidateHostBuffer(IntPtr ptr); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_SyncHostBuffer(IntPtr ptr, long byteCount); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_SyncHostBuffer(IntPtr ptr, long byteCount); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern long TSGgml_DeviceCopyCacheResidentBytes(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial long TSGgml_DeviceCopyCacheResidentBytes(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetBackendMemory(out long freeBytes, out long totalBytes); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetBackendMemory(out long freeBytes, out long totalBytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_IsActiveDeviceIntegrated(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_IsActiveDeviceIntegrated(); // Async dispatch (deferred ggml_backend_synchronize). When enabled, per-op // kernels return without waiting on the Metal command buffer; subsequent ops // chain through the Metal command queue, and host-side reads must call // TSGgml_HostReadBarrier first to drain pending GPU work. See // GgmlStorage.EnsureHostReadable for the C# entry point that triggers this. - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_SetAsyncCompute(int enabled); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_SetAsyncCompute(int enabled); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_SetHostMoeThreads(int threads); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_SetHostMoeThreads(int threads); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetAsyncCompute(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetAsyncCompute(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_HostReadBarrier(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_HostReadBarrier(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_PreloadQuantizedWeight(IntPtr cacheKey, IntPtr hostData, int ggmlType, long ne0, long ne1, long rawBytes); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_PreloadQuantizedWeight(IntPtr cacheKey, IntPtr hostData, int ggmlType, long ne0, long ne1, long rawBytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_RegisterOffloadable(IntPtr key); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_RegisterOffloadable(IntPtr key); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_SetOffloadableBudget(long bytes); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_SetOffloadableBudget(long bytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_ClearOffloadableState(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_ClearOffloadableState(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_SetDeviceCopyBudget(long bytes); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_SetDeviceCopyBudget(long bytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DeviceMemoryInfo(out long freeBytes, out long totalBytes); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DeviceMemoryInfo(out long freeBytes, out long totalBytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_RegisterPinnedHostBuffer(IntPtr ptr, long bytes); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_RegisterPinnedHostBuffer(IntPtr ptr, long bytes); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void TSGgml_UnregisterPinnedHostBuffer(IntPtr ptr); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void TSGgml_UnregisterPinnedHostBuffer(IntPtr ptr); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern UIntPtr TSGgml_RowSize(int ggmlType, long ne); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial UIntPtr TSGgml_RowSize(int ggmlType, long ne); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_DequantizeToF32(int ggmlType, IntPtr src, long numElements, IntPtr dst); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_DequantizeToF32(int ggmlType, IntPtr src, long numElements, IntPtr dst); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_ApplyLoraDelta(IntPtr w, int ggmlType, long ne0, long ne1, + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_ApplyLoraDelta(IntPtr w, int ggmlType, long ne0, long ne1, IntPtr up, IntPtr down, int rank, float scale, int nThreads); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern void ggml_quantize_init(int type); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void ggml_quantize_init(int type); - [DllImport(DllName, CallingConvention = CallingConventionType)] + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] [return: MarshalAs(UnmanagedType.I1)] - private static extern bool ggml_quantize_requires_imatrix(int type); + private static partial bool ggml_quantize_requires_imatrix(int type); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern UIntPtr ggml_quantize_chunk(int type, IntPtr src, IntPtr dst, + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial UIntPtr ggml_quantize_chunk(int type, IntPtr src, IntPtr dst, long start, long nrows, long nPerRow, IntPtr imatrix); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_QwenVaeRun(in QwenVaeArgs args); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_QwenVaeRun(in QwenVaeArgs args); /// Run a whole VAE encode/decode op-list as ONE device graph (see QwenVaeArgs). /// Returns false when the backend can't run it (caller falls back to the per-conv path). internal static bool TryQwenVaeRun(in QwenVaeArgs args) => TSGgml_QwenVaeRun(in args) != 0; - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_QwenTeTrunk(in QwenTeTrunkArgs args); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_QwenTeTrunk(in QwenTeTrunkArgs args); /// Run a whole conditioning-encoder transformer trunk as ONE device graph /// (see QwenTeTrunkArgs). Returns false when the backend can't run it (caller falls /// back to the per-op path). internal static bool TryQwenTeTrunk(in QwenTeTrunkArgs args) => TSGgml_QwenTeTrunk(in args) != 0; - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_CopyF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_CopyF32( GgmlTensorView4D result, GgmlTensorView4D src); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_UnaryF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_UnaryF32( int op, GgmlTensorView4D result, GgmlTensorView4D src); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_BinaryTensorF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_BinaryTensorF32( int op, GgmlTensorView4D result, GgmlTensorView4D lhs, GgmlTensorView4D rhs); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedActMulF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedActMulF32( int op, GgmlTensorView4D result, GgmlTensorView4D a, GgmlTensorView4D b); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_FusedActMulSplitF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_FusedActMulSplitF32( int op, GgmlTensorView2D result, GgmlTensorView2D gateUp, int halfDim); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_BinaryScalarF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_BinaryScalarF32( int op, GgmlTensorView4D result, GgmlTensorView4D src, float scalar); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_ActivationGradF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_ActivationGradF32( int op, GgmlTensorView4D result, GgmlTensorView4D src, @@ -3274,8 +3425,9 @@ private static extern int TSGgml_ActivationGradF32( GgmlTensorView4D accumulation, int hasAccumulation); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_NormF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_NormF32( int op, GgmlTensorView4D result, GgmlTensorView4D src, @@ -3284,8 +3436,9 @@ private static extern int TSGgml_NormF32( int hasBeta, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_NormGradF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_NormGradF32( int op, GgmlTensorView4D result, GgmlTensorView4D gradGamma, @@ -3296,21 +3449,24 @@ private static extern int TSGgml_NormGradF32( int hasGradBeta, float eps); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_IndexSelectF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_IndexSelectF32( GgmlTensorView2D result, GgmlTensorView2D src, GgmlContiguousTensor indices, int addToResult); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_IndexSelectGradF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_IndexSelectGradF32( GgmlTensorView2D grad, GgmlTensorView2D adj, GgmlContiguousTensor indices); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_RoPEF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_RoPEF32( GgmlTensorView4D result, GgmlTensorView4D src, int seqLen, @@ -3318,8 +3474,9 @@ private static extern int TSGgml_RoPEF32( int addToResult, int invertPositions); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_RoPEExF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_RoPEExF32( GgmlTensorView4D result, GgmlTensorView4D src, GgmlContiguousTensor positions, @@ -3335,8 +3492,9 @@ private static extern int TSGgml_RoPEExF32( int addToResult, int invertPositions); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_RoPEMRoPEF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_RoPEMRoPEF32( GgmlTensorView4D result, GgmlTensorView4D src, GgmlContiguousTensor positions, @@ -3351,8 +3509,9 @@ private static extern int TSGgml_RoPEMRoPEF32( float betaFast, float betaSlow); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_RoPEExFreqFactorsF32( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_RoPEExFreqFactorsF32( GgmlTensorView4D result, GgmlTensorView4D src, GgmlContiguousTensor positions, diff --git a/TensorSharp.Backends.GGML/GgmlTensorParallel.cs b/TensorSharp.Backends.GGML/GgmlTensorParallel.cs index c9304db1..3b1a0598 100644 --- a/TensorSharp.Backends.GGML/GgmlTensorParallel.cs +++ b/TensorSharp.Backends.GGML/GgmlTensorParallel.cs @@ -12,47 +12,60 @@ // rank selection, the cross-GPU AllReduce, and the fused multi-rank matmul that // backs column-/row-parallel linear layers. using System; +using System.Runtime.CompilerServices; using System.Runtime.InteropServices; namespace TensorSharp.GGML { internal static partial class GgmlNative { - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetGpuDeviceCount(int backendType); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetGpuDeviceCount(int backendType); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetGpuDeviceDescription(int backendType, int deviceIndex, byte[] description, int descriptionSize); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetGpuDeviceDescription(int backendType, int deviceIndex, byte[] description, int descriptionSize); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_TensorParallelInit(int backendType, int[] deviceIndices, int count, int concurrentRanks); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_TensorParallelInit(int backendType, int[] deviceIndices, int count, int concurrentRanks); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_SetActiveDevice(int rank); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_SetActiveDevice(int rank); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetActiveDevice(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetActiveDevice(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_GetTensorParallelDegree(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_GetTensorParallelDegree(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_TensorParallelHasDeviceAllReduce(); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_TensorParallelHasDeviceAllReduce(); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern unsafe int TSGgml_TensorParallelAllReduceHost(float** buffers, int rankCount, long count); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static unsafe partial int TSGgml_TensorParallelAllReduceHost(float** buffers, int rankCount, long count); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern unsafe int TSGgml_TensorParallelAllReduceDevice(float** buffers, int rankCount, long count); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static unsafe partial int TSGgml_TensorParallelAllReduceDevice(float** buffers, int rankCount, long count); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_TensorParallelFusedAvailable(int rankCount); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_TensorParallelFusedAvailable(int rankCount); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern int TSGgml_TensorParallelExecutePlans(IntPtr[] plans, int rankCount); + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int TSGgml_TensorParallelExecutePlans(IntPtr[] plans, int rankCount); - [DllImport(DllName, CallingConvention = CallingConventionType)] - private static extern unsafe int TSGgml_TensorParallelMatmul( + [LibraryImport(DllName)] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static unsafe partial int TSGgml_TensorParallelMatmul( GgmlTensorView2D* results, GgmlTensorView2D* inputs, IntPtr* weightData, From 754ec24f2d9eaff6b6e823731267c17e0ad4af2f Mon Sep 17 00:00:00 2001 From: craig-b <3112503+craig-b@users.noreply.github.com> Date: Sat, 15 Aug 2026 14:31:40 +0100 Subject: [PATCH 3/5] Migrate MLX backend P/Invokes to LibraryImport The MlxOptional* structs cross by value with bool fields, so the assembly disables runtime marshalling; their managed layout then matches the mlx-c definitions exactly (1-byte bool). --- TensorSharp.Backends.MLX/MlxNative.cs | 523 +++++++++++-------- TensorSharp.Backends.MLX/MoeExpertOffload.cs | 6 +- 2 files changed, 319 insertions(+), 210 deletions(-) diff --git a/TensorSharp.Backends.MLX/MlxNative.cs b/TensorSharp.Backends.MLX/MlxNative.cs index 9643276d..c2d98bad 100644 --- a/TensorSharp.Backends.MLX/MlxNative.cs +++ b/TensorSharp.Backends.MLX/MlxNative.cs @@ -2,11 +2,17 @@ using System.Collections.Generic; using System.IO; using System.Reflection; +using System.Runtime.CompilerServices; using System.Runtime.InteropServices; +// Structs with bool fields (MlxOptional*) cross by value; with runtime +// marshalling disabled they pass with their managed layout, which matches +// the mlx-c definitions (1-byte bool). +[assembly: DisableRuntimeMarshalling] + namespace TensorSharp.MLX { - internal static class MlxNative + internal static partial class MlxNative { private const string LibraryName = "mlxc"; private const int MlxGpu = 1; @@ -9282,275 +9288,365 @@ internal enum MlxBinaryOp Maximum, } - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_metal_is_available")] - private static extern int mlx_metal_is_available([MarshalAs(UnmanagedType.I1)] out bool res); + [LibraryImport(LibraryName, EntryPoint = "mlx_metal_is_available")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_metal_is_available([MarshalAs(UnmanagedType.I1)] out bool res); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_set_error_handler")] - private static extern void mlx_set_error_handler(MlxErrorHandler handler, IntPtr data, IntPtr destructor); + [LibraryImport(LibraryName, EntryPoint = "mlx_set_error_handler")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial void mlx_set_error_handler(MlxErrorHandler handler, IntPtr data, IntPtr destructor); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_device_new_type")] - private static extern MlxDevice mlx_device_new_type(int type, int index); + [LibraryImport(LibraryName, EntryPoint = "mlx_device_new_type")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxDevice mlx_device_new_type(int type, int index); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_device_free")] - private static extern int mlx_device_free(MlxDevice dev); + [LibraryImport(LibraryName, EntryPoint = "mlx_device_free")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_device_free(MlxDevice dev); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_device_is_available")] - private static extern int mlx_device_is_available([MarshalAs(UnmanagedType.I1)] out bool avail, MlxDevice dev); + [LibraryImport(LibraryName, EntryPoint = "mlx_device_is_available")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_device_is_available([MarshalAs(UnmanagedType.I1)] out bool avail, MlxDevice dev); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_set_default_device")] - private static extern int mlx_set_default_device(MlxDevice dev); + [LibraryImport(LibraryName, EntryPoint = "mlx_set_default_device")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_set_default_device(MlxDevice dev); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_clear_cache")] - private static extern int mlx_clear_cache(); + [LibraryImport(LibraryName, EntryPoint = "mlx_clear_cache")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_clear_cache(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_get_active_memory")] - private static extern int mlx_get_active_memory(ref nuint res); + [LibraryImport(LibraryName, EntryPoint = "mlx_get_active_memory")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_get_active_memory(ref nuint res); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_get_cache_memory")] - private static extern int mlx_get_cache_memory(ref nuint res); + [LibraryImport(LibraryName, EntryPoint = "mlx_get_cache_memory")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_get_cache_memory(ref nuint res); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_get_peak_memory")] - private static extern int mlx_get_peak_memory(ref nuint res); + [LibraryImport(LibraryName, EntryPoint = "mlx_get_peak_memory")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_get_peak_memory(ref nuint res); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_reset_peak_memory")] - private static extern int mlx_reset_peak_memory(); + [LibraryImport(LibraryName, EntryPoint = "mlx_reset_peak_memory")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_reset_peak_memory(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_set_cache_limit")] - private static extern int mlx_set_cache_limit(ref nuint previous, nuint limit); + [LibraryImport(LibraryName, EntryPoint = "mlx_set_cache_limit")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_set_cache_limit(ref nuint previous, nuint limit); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_set_wired_limit")] - private static extern int mlx_set_wired_limit(ref nuint previous, nuint limit); + [LibraryImport(LibraryName, EntryPoint = "mlx_set_wired_limit")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_set_wired_limit(ref nuint previous, nuint limit); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_set_memory_limit")] - private static extern int mlx_set_memory_limit(ref nuint previous, nuint limit); + [LibraryImport(LibraryName, EntryPoint = "mlx_set_memory_limit")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_set_memory_limit(ref nuint previous, nuint limit); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_get_default_stream")] - private static extern int mlx_get_default_stream(out MlxStream stream, MlxDevice dev); + [LibraryImport(LibraryName, EntryPoint = "mlx_get_default_stream")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_get_default_stream(out MlxStream stream, MlxDevice dev); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_stream_free")] - private static extern int mlx_stream_free(MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_stream_free")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_stream_free(MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_new_data")] - private static extern MlxArray mlx_array_new_data(IntPtr data, int[] shape, int dim, int dtype); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_new_data")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxArray mlx_array_new_data(IntPtr data, int[] shape, int dim, int dtype); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_new_data_managed")] - private static extern MlxArray mlx_array_new_data_managed(IntPtr data, int[] shape, int dim, int dtype, IntPtr dtor); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_new_data_managed")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxArray mlx_array_new_data_managed(IntPtr data, int[] shape, int dim, int dtype, IntPtr dtor); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_new_float32")] - private static extern MlxArray mlx_array_new_float32(float value); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_new_float32")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxArray mlx_array_new_float32(float value); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_new_int")] - private static extern MlxArray mlx_array_new_int(int value); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_new_int")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxArray mlx_array_new_int(int value); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_free")] - private static extern int mlx_array_free(MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_free")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_array_free(MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_size")] - private static extern nuint mlx_array_size_native(MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_size")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial nuint mlx_array_size_native(MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "_mlx_array_is_row_contiguous")] - private static extern int mlx_array_is_row_contiguous_native([MarshalAs(UnmanagedType.I1)] out bool result, MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "_mlx_array_is_row_contiguous")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_array_is_row_contiguous_native([MarshalAs(UnmanagedType.I1)] out bool result, MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_astype")] - private static extern int mlx_astype(out MlxArray result, MlxArray array, int dtype, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_astype")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_astype(out MlxArray result, MlxArray array, int dtype, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_data_float32")] - private static extern IntPtr mlx_array_data_float32(MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_data_float32")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr mlx_array_data_float32(MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_data_float64")] - private static extern IntPtr mlx_array_data_float64(MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_data_float64")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr mlx_array_data_float64(MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_data_float16")] - private static extern IntPtr mlx_array_data_float16(MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_data_float16")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr mlx_array_data_float16(MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_data_int32")] - private static extern IntPtr mlx_array_data_int32(MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_data_int32")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr mlx_array_data_int32(MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_data_uint8")] - private static extern IntPtr mlx_array_data_uint8(MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_data_uint8")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial IntPtr mlx_array_data_uint8(MlxArray array); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_array_new")] - private static extern MlxVectorArray mlx_vector_array_new(); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_array_new")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxVectorArray mlx_vector_array_new(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_array_new_data")] - private static extern MlxVectorArray mlx_vector_array_new_data(IntPtr values, nuint size); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_array_new_data")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxVectorArray mlx_vector_array_new_data(IntPtr values, nuint size); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_array_free")] - private static extern int mlx_vector_array_free(MlxVectorArray vector); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_array_free")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_vector_array_free(MlxVectorArray vector); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_array_append_value")] - private static extern int mlx_vector_array_append_value(MlxVectorArray vector, MlxArray value); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_array_append_value")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_vector_array_append_value(MlxVectorArray vector, MlxArray value); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_array_size")] - private static extern nuint mlx_vector_array_size(MlxVectorArray vector); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_array_size")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial nuint mlx_vector_array_size(MlxVectorArray vector); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_array_get")] - private static extern int mlx_vector_array_get(out MlxArray result, MlxVectorArray vector, nuint index); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_array_get")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_vector_array_get(out MlxArray result, MlxVectorArray vector, nuint index); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_string_new")] - private static extern MlxVectorString mlx_vector_string_new(); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_string_new")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxVectorString mlx_vector_string_new(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_string_free")] - private static extern int mlx_vector_string_free(MlxVectorString vector); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_string_free")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_vector_string_free(MlxVectorString vector); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_string_append_value")] - private static extern int mlx_vector_string_append_value(MlxVectorString vector, IntPtr value); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_string_append_value")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_vector_string_append_value(MlxVectorString vector, IntPtr value); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_eval")] - private static extern int mlx_eval(MlxVectorArray outputs); + [LibraryImport(LibraryName, EntryPoint = "mlx_eval")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_eval(MlxVectorArray outputs); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_async_eval")] - private static extern int mlx_async_eval(MlxVectorArray outputs); + [LibraryImport(LibraryName, EntryPoint = "mlx_async_eval")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_async_eval(MlxVectorArray outputs); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_as_strided")] - private static extern int mlx_as_strided(out MlxArray result, MlxArray array, int[] shape, nuint shapeCount, long[] strides, nuint stridesCount, nuint offset, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_as_strided")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_as_strided(out MlxArray result, MlxArray array, int[] shape, nuint shapeCount, long[] strides, nuint stridesCount, nuint offset, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_reshape")] - private static extern int mlx_reshape(out MlxArray result, MlxArray array, int[] shape, nuint shapeCount, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_reshape")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_reshape(out MlxArray result, MlxArray array, int[] shape, nuint shapeCount, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_contiguous")] - private static extern int mlx_contiguous(out MlxArray result, MlxArray array, [MarshalAs(UnmanagedType.I1)] bool allowColMajor, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_contiguous")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_contiguous(out MlxArray result, MlxArray array, [MarshalAs(UnmanagedType.I1)] bool allowColMajor, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_concatenate_axis")] - private static extern int mlx_concatenate_axis(out MlxArray result, MlxVectorArray arrays, int axis, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_concatenate_axis")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_concatenate_axis(out MlxArray result, MlxVectorArray arrays, int axis, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_full")] - private static extern int mlx_full(out MlxArray result, int[] shape, nuint shapeCount, MlxArray values, int dtype, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_full")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_full(out MlxArray result, int[] shape, nuint shapeCount, MlxArray values, int dtype, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_arange")] - private static extern int mlx_arange(out MlxArray result, double start, double stop, double step, int dtype, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_arange")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_arange(out MlxArray result, double start, double stop, double step, int dtype, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_transpose_axes")] - private static extern int mlx_transpose_axes(out MlxArray result, MlxArray array, int[] axes, nuint axesCount, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_transpose_axes")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_transpose_axes(out MlxArray result, MlxArray array, int[] axes, nuint axesCount, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_abs")] - private static extern int mlx_abs(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_abs")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_abs(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_negative")] - private static extern int mlx_negative(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_negative")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_negative(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_sqrt")] - private static extern int mlx_sqrt(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_sqrt")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_sqrt(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_rsqrt")] - private static extern int mlx_rsqrt(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_rsqrt")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_rsqrt(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_exp")] - private static extern int mlx_exp(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_exp")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_exp(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_log")] - private static extern int mlx_log(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_log")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_log(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_log1p")] - private static extern int mlx_log1p(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_log1p")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_log1p(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_floor")] - private static extern int mlx_floor(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_floor")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_floor(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_ceil")] - private static extern int mlx_ceil(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_ceil")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_ceil(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_sin")] - private static extern int mlx_sin(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_sin")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_sin(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_cos")] - private static extern int mlx_cos(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_cos")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_cos(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_tanh")] - private static extern int mlx_tanh(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_tanh")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_tanh(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_sigmoid")] - private static extern int mlx_sigmoid(out MlxArray result, MlxArray input, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_sigmoid")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_sigmoid(out MlxArray result, MlxArray input, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_add")] - private static extern int mlx_add(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_add")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_add(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_subtract")] - private static extern int mlx_subtract(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_subtract")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_subtract(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_multiply")] - private static extern int mlx_multiply(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_multiply")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_multiply(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_divide")] - private static extern int mlx_divide(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_divide")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_divide(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_maximum")] - private static extern int mlx_maximum(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_maximum")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_maximum(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_remainder")] - private static extern int mlx_remainder(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_remainder")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_remainder(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_greater")] - private static extern int mlx_greater(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_greater")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_greater(out MlxArray result, MlxArray lhs, MlxArray rhs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_where")] - private static extern int mlx_where(out MlxArray result, MlxArray condition, MlxArray whenTrue, MlxArray whenFalse, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_where")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_where(out MlxArray result, MlxArray condition, MlxArray whenTrue, MlxArray whenFalse, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_addmm")] - private static extern int mlx_addmm(out MlxArray result, MlxArray src, MlxArray m1, MlxArray m2, float alpha, float beta, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_addmm")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_addmm(out MlxArray result, MlxArray src, MlxArray m1, MlxArray m2, float alpha, float beta, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_softmax_axis")] - private static extern int mlx_softmax_axis(out MlxArray result, MlxArray input, int axis, [MarshalAs(UnmanagedType.I1)] bool precise, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_softmax_axis")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_softmax_axis(out MlxArray result, MlxArray input, int axis, [MarshalAs(UnmanagedType.I1)] bool precise, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_repeat_axis")] - private static extern int mlx_repeat_axis(out MlxArray result, MlxArray input, int repeats, int axis, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_repeat_axis")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_repeat_axis(out MlxArray result, MlxArray input, int repeats, int axis, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_layer_norm")] - private static extern int mlx_fast_layer_norm(out MlxArray result, MlxArray input, MlxArray weight, MlxArray bias, float eps, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_layer_norm")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_layer_norm(out MlxArray result, MlxArray input, MlxArray weight, MlxArray bias, float eps, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_rms_norm")] - private static extern int mlx_fast_rms_norm(out MlxArray result, MlxArray input, MlxArray weight, float eps, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_rms_norm")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_rms_norm(out MlxArray result, MlxArray input, MlxArray weight, float eps, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_scaled_dot_product_attention")] - private static extern int mlx_fast_scaled_dot_product_attention(out MlxArray result, MlxArray query, MlxArray key, MlxArray value, float scale, IntPtr maskMode, MlxArray mask, MlxArray sinks, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_scaled_dot_product_attention")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_scaled_dot_product_attention(out MlxArray result, MlxArray query, MlxArray key, MlxArray value, float scale, IntPtr maskMode, MlxArray mask, MlxArray sinks, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_rope_dynamic")] - private static extern int mlx_fast_rope_dynamic(out MlxArray result, MlxArray input, int dims, [MarshalAs(UnmanagedType.I1)] bool traditional, MlxOptionalFloat baseValue, float scale, MlxArray offsets, MlxArray freqs, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_rope_dynamic")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_rope_dynamic(out MlxArray result, MlxArray input, int dims, [MarshalAs(UnmanagedType.I1)] bool traditional, MlxOptionalFloat baseValue, float scale, MlxArray offsets, MlxArray freqs, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_take_axis")] - private static extern int mlx_take_axis(out MlxArray result, MlxArray input, MlxArray indices, int axis, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_take_axis")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_take_axis(out MlxArray result, MlxArray input, MlxArray indices, int axis, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_take_along_axis")] - private static extern int mlx_take_along_axis(out MlxArray result, MlxArray input, MlxArray indices, int axis, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_take_along_axis")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_take_along_axis(out MlxArray result, MlxArray input, MlxArray indices, int axis, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_argmax_axis")] - private static extern int mlx_argmax_axis(out MlxArray result, MlxArray input, int axis, [MarshalAs(UnmanagedType.I1)] bool keepdims, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_argmax_axis")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_argmax_axis(out MlxArray result, MlxArray input, int axis, [MarshalAs(UnmanagedType.I1)] bool keepdims, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_argpartition_axis")] - private static extern int mlx_argpartition_axis(out MlxArray result, MlxArray input, int kth, int axis, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_argpartition_axis")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_argpartition_axis(out MlxArray result, MlxArray input, int kth, int axis, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_slice_update")] - private static extern int mlx_slice_update(out MlxArray result, MlxArray input, MlxArray update, int[] starts, nuint startCount, int[] stops, nuint stopCount, int[] strides, nuint strideCount, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_slice_update")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_slice_update(out MlxArray result, MlxArray input, MlxArray update, int[] starts, nuint startCount, int[] stops, nuint stopCount, int[] strides, nuint strideCount, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_slice")] - private static extern int mlx_slice(out MlxArray result, MlxArray input, int[] starts, nuint startCount, int[] stops, nuint stopCount, int[] strides, nuint strideCount, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_slice")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_slice(out MlxArray result, MlxArray input, int[] starts, nuint startCount, int[] stops, nuint stopCount, int[] strides, nuint strideCount, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_quantized_matmul")] - private static extern int mlx_quantized_matmul(out MlxArray result, MlxArray input, MlxArray weight, MlxArray scales, MlxArray biases, [MarshalAs(UnmanagedType.I1)] bool transpose, MlxOptionalInt groupSize, MlxOptionalInt bits, IntPtr mode, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_quantized_matmul")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_quantized_matmul(out MlxArray result, MlxArray input, MlxArray weight, MlxArray scales, MlxArray biases, [MarshalAs(UnmanagedType.I1)] bool transpose, MlxOptionalInt groupSize, MlxOptionalInt bits, IntPtr mode, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_dequantize")] - private static extern int mlx_dequantize(out MlxArray result, MlxArray weight, MlxArray scales, MlxArray biases, MlxOptionalInt groupSize, MlxOptionalInt bits, IntPtr mode, MlxArray globalScale, MlxOptionalDType dtype, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_dequantize")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_dequantize(out MlxArray result, MlxArray weight, MlxArray scales, MlxArray biases, MlxOptionalInt groupSize, MlxOptionalInt bits, IntPtr mode, MlxArray globalScale, MlxOptionalDType dtype, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_config_new")] - private static extern MlxFastMetalKernelConfig mlx_fast_metal_kernel_config_new(); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_config_new")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxFastMetalKernelConfig mlx_fast_metal_kernel_config_new(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_config_free")] - private static extern int mlx_fast_metal_kernel_config_free(MlxFastMetalKernelConfig config); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_config_free")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_metal_kernel_config_free(MlxFastMetalKernelConfig config); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_config_add_output_arg")] - private static extern int mlx_fast_metal_kernel_config_add_output_arg(MlxFastMetalKernelConfig config, int[] shape, nuint size, int dtype); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_config_add_output_arg")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_metal_kernel_config_add_output_arg(MlxFastMetalKernelConfig config, int[] shape, nuint size, int dtype); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_config_set_grid")] - private static extern int mlx_fast_metal_kernel_config_set_grid(MlxFastMetalKernelConfig config, int grid1, int grid2, int grid3); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_config_set_grid")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_metal_kernel_config_set_grid(MlxFastMetalKernelConfig config, int grid1, int grid2, int grid3); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_config_set_thread_group")] - private static extern int mlx_fast_metal_kernel_config_set_thread_group(MlxFastMetalKernelConfig config, int thread1, int thread2, int thread3); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_config_set_thread_group")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_metal_kernel_config_set_thread_group(MlxFastMetalKernelConfig config, int thread1, int thread2, int thread3); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_config_add_template_arg_int")] - private static extern int mlx_fast_metal_kernel_config_add_template_arg_int(MlxFastMetalKernelConfig config, IntPtr name, int value); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_config_add_template_arg_int")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_metal_kernel_config_add_template_arg_int(MlxFastMetalKernelConfig config, IntPtr name, int value); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_new")] - private static extern MlxFastMetalKernel mlx_fast_metal_kernel_new( + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_new")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxFastMetalKernel mlx_fast_metal_kernel_new( IntPtr name, MlxVectorString inputNames, MlxVectorString outputNames, @@ -9559,47 +9655,60 @@ private static extern MlxFastMetalKernel mlx_fast_metal_kernel_new( [MarshalAs(UnmanagedType.I1)] bool ensureRowContiguous, [MarshalAs(UnmanagedType.I1)] bool atomicOutputs); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_fast_metal_kernel_apply")] - private static extern int mlx_fast_metal_kernel_apply(ref MlxVectorArray outputs, MlxFastMetalKernel kernel, MlxVectorArray inputs, MlxFastMetalKernelConfig config, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_fast_metal_kernel_apply")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_fast_metal_kernel_apply(ref MlxVectorArray outputs, MlxFastMetalKernel kernel, MlxVectorArray inputs, MlxFastMetalKernelConfig config, MlxStream stream); // ----- mlx_compile / mlx_closure_* (graph compilation) ----- - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_compile")] - private static extern int mlx_compile(out MlxClosure result, MlxClosure fun, [MarshalAs(UnmanagedType.I1)] bool shapeless); + [LibraryImport(LibraryName, EntryPoint = "mlx_compile")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_compile(out MlxClosure result, MlxClosure fun, [MarshalAs(UnmanagedType.I1)] bool shapeless); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_enable_compile")] - private static extern int mlx_enable_compile(); + [LibraryImport(LibraryName, EntryPoint = "mlx_enable_compile")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_enable_compile(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_disable_compile")] - private static extern int mlx_disable_compile(); + [LibraryImport(LibraryName, EntryPoint = "mlx_disable_compile")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_disable_compile(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_set_compile_mode")] - private static extern int mlx_set_compile_mode(int mode); + [LibraryImport(LibraryName, EntryPoint = "mlx_set_compile_mode")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_set_compile_mode(int mode); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_closure_new")] - private static extern MlxClosure mlx_closure_new(); + [LibraryImport(LibraryName, EntryPoint = "mlx_closure_new")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxClosure mlx_closure_new(); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_closure_free")] - private static extern int mlx_closure_free(MlxClosure closure); + [LibraryImport(LibraryName, EntryPoint = "mlx_closure_free")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_closure_free(MlxClosure closure); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_closure_new_func_payload")] - private static extern MlxClosure mlx_closure_new_func_payload(IntPtr fun, IntPtr payload, IntPtr destructor); + [LibraryImport(LibraryName, EntryPoint = "mlx_closure_new_func_payload")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial MlxClosure mlx_closure_new_func_payload(IntPtr fun, IntPtr payload, IntPtr destructor); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_closure_apply")] - private static extern int mlx_closure_apply(ref MlxVectorArray result, MlxClosure closure, MlxVectorArray input); + [LibraryImport(LibraryName, EntryPoint = "mlx_closure_apply")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_closure_apply(ref MlxVectorArray result, MlxClosure closure, MlxVectorArray input); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_vector_array_set_data")] - private static extern int mlx_vector_array_set_data(ref MlxVectorArray vec, IntPtr data, nuint size); + [LibraryImport(LibraryName, EntryPoint = "mlx_vector_array_set_data")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_vector_array_set_data(ref MlxVectorArray vec, IntPtr data, nuint size); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_array_dtype")] - private static extern int mlx_array_dtype(out int dtype, MlxArray array); + [LibraryImport(LibraryName, EntryPoint = "mlx_array_dtype")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_array_dtype(out int dtype, MlxArray array); // ----- mlx_gather_mm / mlx_gather_qmm (MoE-friendly batched matmul) ----- - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_gather_mm")] - private static extern int mlx_gather_mm(out MlxArray result, MlxArray a, MlxArray b, MlxArray lhsIndices, MlxArray rhsIndices, [MarshalAs(UnmanagedType.I1)] bool sortedIndices, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_gather_mm")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_gather_mm(out MlxArray result, MlxArray a, MlxArray b, MlxArray lhsIndices, MlxArray rhsIndices, [MarshalAs(UnmanagedType.I1)] bool sortedIndices, MlxStream stream); - [DllImport(LibraryName, CallingConvention = CallingConvention.Cdecl, EntryPoint = "mlx_gather_qmm")] - private static extern int mlx_gather_qmm(out MlxArray result, MlxArray x, MlxArray w, MlxArray scales, MlxArray biases, MlxArray lhsIndices, MlxArray rhsIndices, [MarshalAs(UnmanagedType.I1)] bool transpose, MlxOptionalInt groupSize, MlxOptionalInt bits, IntPtr mode, [MarshalAs(UnmanagedType.I1)] bool sortedIndices, MlxStream stream); + [LibraryImport(LibraryName, EntryPoint = "mlx_gather_qmm")] + [UnmanagedCallConv(CallConvs = new[] { typeof(CallConvCdecl) })] + private static partial int mlx_gather_qmm(out MlxArray result, MlxArray x, MlxArray w, MlxArray scales, MlxArray biases, MlxArray lhsIndices, MlxArray rhsIndices, [MarshalAs(UnmanagedType.I1)] bool transpose, MlxOptionalInt groupSize, MlxOptionalInt bits, IntPtr mode, [MarshalAs(UnmanagedType.I1)] bool sortedIndices, MlxStream stream); } } diff --git a/TensorSharp.Backends.MLX/MoeExpertOffload.cs b/TensorSharp.Backends.MLX/MoeExpertOffload.cs index e06e93f4..af0fe001 100644 --- a/TensorSharp.Backends.MLX/MoeExpertOffload.cs +++ b/TensorSharp.Backends.MLX/MoeExpertOffload.cs @@ -31,7 +31,7 @@ namespace TensorSharp.MLX /// are small in aggregate, hot on every forward, and remain permanently /// device-resident under the existing preload path. /// - public static class MoeExpertOffload + public static partial class MoeExpertOffload { private const string EnvVarMb = "TS_MLX_EXPERT_OFFLOAD_MB"; private static readonly long _maxCacheBytes = ParseLimit(); @@ -138,8 +138,8 @@ public static unsafe void AdvisePagesNotNeeded(IntPtr data, long byteCount) private const int MadvDontNeed = 4; - [DllImport("libc", SetLastError = true, EntryPoint = "madvise")] - private static extern unsafe int madvise(void* addr, nuint len, int advice); + [LibraryImport("libc", EntryPoint = "madvise", SetLastError = true)] + private static unsafe partial int madvise(void* addr, nuint len, int advice); private static long ParseLimit() { From 0c2ff16691ab8f5d44eea5e10caee4f3ea6c0aea Mon Sep 17 00:00:00 2001 From: craig-b <3112503+craig-b@users.noreply.github.com> Date: Sat, 15 Aug 2026 14:31:40 +0100 Subject: [PATCH 4/5] Migrate CUDA interop P/Invokes to LibraryImport cuModuleGetFunction's kernel-name parameter now marshals explicitly as UTF-8. --- .../Interop/CublasApi.cs | 30 +-- .../Interop/CudaDriverApi.cs | 194 +++++++++--------- 2 files changed, 112 insertions(+), 112 deletions(-) diff --git a/TensorSharp.Backends.Cuda/Interop/CublasApi.cs b/TensorSharp.Backends.Cuda/Interop/CublasApi.cs index 9b5f95b1..a9ff605a 100644 --- a/TensorSharp.Backends.Cuda/Interop/CublasApi.cs +++ b/TensorSharp.Backends.Cuda/Interop/CublasApi.cs @@ -3,24 +3,24 @@ namespace TensorSharp.Cuda.Interop { - internal static class CublasApi + internal static partial class CublasApi { private const string LibName = "cublas"; - [DllImport(LibName, EntryPoint = "cublasCreate_v2")] - public static extern int cublasCreate(out IntPtr handle); + [LibraryImport(LibName, EntryPoint = "cublasCreate_v2")] + public static partial int cublasCreate(out IntPtr handle); - [DllImport(LibName, EntryPoint = "cublasDestroy_v2")] - public static extern int cublasDestroy(IntPtr handle); + [LibraryImport(LibName, EntryPoint = "cublasDestroy_v2")] + public static partial int cublasDestroy(IntPtr handle); - [DllImport(LibName, EntryPoint = "cublasSetStream_v2")] - public static extern int cublasSetStream(IntPtr handle, IntPtr stream); + [LibraryImport(LibName, EntryPoint = "cublasSetStream_v2")] + public static partial int cublasSetStream(IntPtr handle, IntPtr stream); - [DllImport(LibName)] - public static extern int cublasSetMathMode(IntPtr handle, int mode); + [LibraryImport(LibName)] + public static partial int cublasSetMathMode(IntPtr handle, int mode); - [DllImport(LibName, EntryPoint = "cublasSgemm_v2")] - public static extern int cublasSgemm( + [LibraryImport(LibName, EntryPoint = "cublasSgemm_v2")] + public static partial int cublasSgemm( IntPtr handle, int transa, int transb, @@ -36,8 +36,8 @@ public static extern int cublasSgemm( IntPtr c, int ldc); - [DllImport(LibName)] - public static extern int cublasSgemmStridedBatched( + [LibraryImport(LibName)] + public static partial int cublasSgemmStridedBatched( IntPtr handle, int transa, int transb, @@ -57,8 +57,8 @@ public static extern int cublasSgemmStridedBatched( long strideC, int batchCount); - [DllImport(LibName)] - public static extern int cublasGemmEx( + [LibraryImport(LibName)] + public static partial int cublasGemmEx( IntPtr handle, int transa, int transb, diff --git a/TensorSharp.Backends.Cuda/Interop/CudaDriverApi.cs b/TensorSharp.Backends.Cuda/Interop/CudaDriverApi.cs index 76fbb09a..51423496 100644 --- a/TensorSharp.Backends.Cuda/Interop/CudaDriverApi.cs +++ b/TensorSharp.Backends.Cuda/Interop/CudaDriverApi.cs @@ -3,107 +3,107 @@ namespace TensorSharp.Cuda.Interop { - internal static class CudaDriverApi + internal static partial class CudaDriverApi { private const string LibName = "cuda"; - [DllImport(LibName)] - public static extern int cuInit(uint flags); + [LibraryImport(LibName)] + public static partial int cuInit(uint flags); - [DllImport(LibName)] - public static extern int cuDeviceGet(out int device, int ordinal); + [LibraryImport(LibName)] + public static partial int cuDeviceGet(out int device, int ordinal); - [DllImport(LibName)] - public static extern int cuDeviceGetCount(out int count); + [LibraryImport(LibName)] + public static partial int cuDeviceGetCount(out int count); - [DllImport(LibName)] - public static extern int cuDeviceGetName(byte[] name, int len, int device); + [LibraryImport(LibName)] + public static partial int cuDeviceGetName(byte[] name, int len, int device); - [DllImport(LibName, EntryPoint = "cuDeviceTotalMem_v2")] - public static extern int cuDeviceTotalMem(out UIntPtr bytes, int device); + [LibraryImport(LibName, EntryPoint = "cuDeviceTotalMem_v2")] + public static partial int cuDeviceTotalMem(out UIntPtr bytes, int device); - [DllImport(LibName, EntryPoint = "cuMemGetInfo_v2")] - public static extern int cuMemGetInfo(out UIntPtr free, out UIntPtr total); + [LibraryImport(LibName, EntryPoint = "cuMemGetInfo_v2")] + public static partial int cuMemGetInfo(out UIntPtr free, out UIntPtr total); - [DllImport(LibName)] - public static extern int cuDeviceGetAttribute(out int value, int attribute, int device); + [LibraryImport(LibName)] + public static partial int cuDeviceGetAttribute(out int value, int attribute, int device); - [DllImport(LibName, EntryPoint = "cuCtxCreate_v2")] - public static extern int cuCtxCreate(out IntPtr ctx, uint flags, int device); + [LibraryImport(LibName, EntryPoint = "cuCtxCreate_v2")] + public static partial int cuCtxCreate(out IntPtr ctx, uint flags, int device); - [DllImport(LibName, EntryPoint = "cuCtxDestroy_v2")] - public static extern int cuCtxDestroy(IntPtr ctx); + [LibraryImport(LibName, EntryPoint = "cuCtxDestroy_v2")] + public static partial int cuCtxDestroy(IntPtr ctx); - [DllImport(LibName)] - public static extern int cuCtxSetCurrent(IntPtr ctx); + [LibraryImport(LibName)] + public static partial int cuCtxSetCurrent(IntPtr ctx); - [DllImport(LibName)] - public static extern int cuCtxGetCurrent(out IntPtr ctx); + [LibraryImport(LibName)] + public static partial int cuCtxGetCurrent(out IntPtr ctx); - [DllImport(LibName)] - public static extern int cuDevicePrimaryCtxRetain(out IntPtr ctx, int device); + [LibraryImport(LibName)] + public static partial int cuDevicePrimaryCtxRetain(out IntPtr ctx, int device); - [DllImport(LibName)] - public static extern int cuDevicePrimaryCtxRelease(int device); + [LibraryImport(LibName)] + public static partial int cuDevicePrimaryCtxRelease(int device); - [DllImport(LibName, EntryPoint = "cuMemAlloc_v2")] - public static extern int cuMemAlloc(out IntPtr devicePtr, UIntPtr byteSize); + [LibraryImport(LibName, EntryPoint = "cuMemAlloc_v2")] + public static partial int cuMemAlloc(out IntPtr devicePtr, UIntPtr byteSize); - [DllImport(LibName, EntryPoint = "cuMemFree_v2")] - public static extern int cuMemFree(IntPtr devicePtr); + [LibraryImport(LibName, EntryPoint = "cuMemFree_v2")] + public static partial int cuMemFree(IntPtr devicePtr); /// Page-locked (pinned) host allocation. A captured HtoD memcpy /// node whose source is pinned host memory re-reads the buffer on every /// graph launch, which is how decode-graph replays receive per-token /// parameters (see CudaDecodeDynParams). - [DllImport(LibName, EntryPoint = "cuMemHostAlloc")] - public static extern int cuMemHostAlloc(out IntPtr hostPtr, UIntPtr byteSize, uint flags); + [LibraryImport(LibName, EntryPoint = "cuMemHostAlloc")] + public static partial int cuMemHostAlloc(out IntPtr hostPtr, UIntPtr byteSize, uint flags); - [DllImport(LibName, EntryPoint = "cuMemFreeHost")] - public static extern int cuMemFreeHost(IntPtr hostPtr); + [LibraryImport(LibName, EntryPoint = "cuMemFreeHost")] + public static partial int cuMemFreeHost(IntPtr hostPtr); - [DllImport(LibName, EntryPoint = "cuMemcpyHtoD_v2")] - public static extern int cuMemcpyHtoD(IntPtr dstDevice, IntPtr srcHost, UIntPtr byteCount); + [LibraryImport(LibName, EntryPoint = "cuMemcpyHtoD_v2")] + public static partial int cuMemcpyHtoD(IntPtr dstDevice, IntPtr srcHost, UIntPtr byteCount); - [DllImport(LibName, EntryPoint = "cuMemcpyHtoDAsync_v2")] - public static extern int cuMemcpyHtoDAsync(IntPtr dstDevice, IntPtr srcHost, UIntPtr byteCount, IntPtr stream); + [LibraryImport(LibName, EntryPoint = "cuMemcpyHtoDAsync_v2")] + public static partial int cuMemcpyHtoDAsync(IntPtr dstDevice, IntPtr srcHost, UIntPtr byteCount, IntPtr stream); - [DllImport(LibName, EntryPoint = "cuMemcpyDtoH_v2")] - public static extern int cuMemcpyDtoH(IntPtr dstHost, IntPtr srcDevice, UIntPtr byteCount); + [LibraryImport(LibName, EntryPoint = "cuMemcpyDtoH_v2")] + public static partial int cuMemcpyDtoH(IntPtr dstHost, IntPtr srcDevice, UIntPtr byteCount); - [DllImport(LibName, EntryPoint = "cuMemcpyDtoHAsync_v2")] - public static extern int cuMemcpyDtoHAsync(IntPtr dstHost, IntPtr srcDevice, UIntPtr byteCount, IntPtr stream); + [LibraryImport(LibName, EntryPoint = "cuMemcpyDtoHAsync_v2")] + public static partial int cuMemcpyDtoHAsync(IntPtr dstHost, IntPtr srcDevice, UIntPtr byteCount, IntPtr stream); - [DllImport(LibName, EntryPoint = "cuMemcpyDtoD_v2")] - public static extern int cuMemcpyDtoD(IntPtr dstDevice, IntPtr srcDevice, UIntPtr byteCount); + [LibraryImport(LibName, EntryPoint = "cuMemcpyDtoD_v2")] + public static partial int cuMemcpyDtoD(IntPtr dstDevice, IntPtr srcDevice, UIntPtr byteCount); - [DllImport(LibName, EntryPoint = "cuMemcpyDtoDAsync_v2")] - public static extern int cuMemcpyDtoDAsync(IntPtr dstDevice, IntPtr srcDevice, UIntPtr byteCount, IntPtr stream); + [LibraryImport(LibName, EntryPoint = "cuMemcpyDtoDAsync_v2")] + public static partial int cuMemcpyDtoDAsync(IntPtr dstDevice, IntPtr srcDevice, UIntPtr byteCount, IntPtr stream); - [DllImport(LibName, EntryPoint = "cuMemsetD8_v2")] - public static extern int cuMemsetD8(IntPtr dstDevice, byte value, UIntPtr count); + [LibraryImport(LibName, EntryPoint = "cuMemsetD8_v2")] + public static partial int cuMemsetD8(IntPtr dstDevice, byte value, UIntPtr count); - [DllImport(LibName, EntryPoint = "cuMemsetD8Async")] - public static extern int cuMemsetD8Async(IntPtr dstDevice, byte value, UIntPtr count, IntPtr stream); + [LibraryImport(LibName, EntryPoint = "cuMemsetD8Async")] + public static partial int cuMemsetD8Async(IntPtr dstDevice, byte value, UIntPtr count, IntPtr stream); - [DllImport(LibName)] - public static extern int cuModuleLoadData(out IntPtr module, IntPtr image); + [LibraryImport(LibName)] + public static partial int cuModuleLoadData(out IntPtr module, IntPtr image); - [DllImport(LibName)] - public static extern int cuModuleGetFunction(out IntPtr function, IntPtr module, string name); + [LibraryImport(LibName, StringMarshalling = StringMarshalling.Utf8)] + public static partial int cuModuleGetFunction(out IntPtr function, IntPtr module, string name); - [DllImport(LibName)] - public static extern int cuModuleUnload(IntPtr module); + [LibraryImport(LibName)] + public static partial int cuModuleUnload(IntPtr module); - [DllImport(LibName)] - public static extern int cuFuncSetAttribute(IntPtr function, int attribute, int value); + [LibraryImport(LibName)] + public static partial int cuFuncSetAttribute(IntPtr function, int attribute, int value); // CUfunction_attribute: opt-in cap for dynamic shared memory per block // (default launches are limited to 48 KB without it). public const int CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES = 8; - [DllImport(LibName)] - public static extern int cuLaunchKernel( + [LibraryImport(LibName)] + public static partial int cuLaunchKernel( IntPtr function, uint gridDimX, uint gridDimY, @@ -116,33 +116,33 @@ public static extern int cuLaunchKernel( IntPtr kernelParams, IntPtr extra); - [DllImport(LibName)] - public static extern int cuStreamCreate(out IntPtr stream, uint flags); + [LibraryImport(LibName)] + public static partial int cuStreamCreate(out IntPtr stream, uint flags); - [DllImport(LibName, EntryPoint = "cuStreamDestroy_v2")] - public static extern int cuStreamDestroy(IntPtr stream); + [LibraryImport(LibName, EntryPoint = "cuStreamDestroy_v2")] + public static partial int cuStreamDestroy(IntPtr stream); - [DllImport(LibName)] - public static extern int cuStreamSynchronize(IntPtr stream); + [LibraryImport(LibName)] + public static partial int cuStreamSynchronize(IntPtr stream); // ---- CUDA Graphs (stream capture) ---- - [DllImport(LibName, EntryPoint = "cuStreamBeginCapture_v2")] - public static extern int cuStreamBeginCapture(IntPtr stream, int mode); + [LibraryImport(LibName, EntryPoint = "cuStreamBeginCapture_v2")] + public static partial int cuStreamBeginCapture(IntPtr stream, int mode); - [DllImport(LibName)] - public static extern int cuStreamEndCapture(IntPtr stream, out IntPtr graph); + [LibraryImport(LibName)] + public static partial int cuStreamEndCapture(IntPtr stream, out IntPtr graph); - [DllImport(LibName)] - public static extern int cuGraphInstantiateWithFlags(out IntPtr graphExec, IntPtr graph, ulong flags); + [LibraryImport(LibName)] + public static partial int cuGraphInstantiateWithFlags(out IntPtr graphExec, IntPtr graph, ulong flags); - [DllImport(LibName)] - public static extern int cuGraphLaunch(IntPtr graphExec, IntPtr stream); + [LibraryImport(LibName)] + public static partial int cuGraphLaunch(IntPtr graphExec, IntPtr stream); - [DllImport(LibName)] - public static extern int cuGraphExecDestroy(IntPtr graphExec); + [LibraryImport(LibName)] + public static partial int cuGraphExecDestroy(IntPtr graphExec); - [DllImport(LibName)] - public static extern int cuGraphDestroy(IntPtr graph); + [LibraryImport(LibName)] + public static partial int cuGraphDestroy(IntPtr graph); // CUstreamCaptureMode: 0 = GLOBAL, 1 = THREAD_LOCAL, 2 = RELAXED. public const int CU_STREAM_CAPTURE_MODE_THREAD_LOCAL = 1; @@ -151,34 +151,34 @@ public static extern int cuLaunchKernel( // responsibility for the referenced memory's lifetime. public const int CU_STREAM_CAPTURE_MODE_RELAXED = 2; - [DllImport(LibName)] - public static extern int cuGetErrorString(int error, out IntPtr str); + [LibraryImport(LibName)] + public static partial int cuGetErrorString(int error, out IntPtr str); // ---- CUDA Events ---- - [DllImport(LibName)] - public static extern int cuEventCreate(out IntPtr phEvent, uint flags); + [LibraryImport(LibName)] + public static partial int cuEventCreate(out IntPtr phEvent, uint flags); - [DllImport(LibName, EntryPoint = "cuEventDestroy_v2")] - public static extern int cuEventDestroy(IntPtr hEvent); + [LibraryImport(LibName, EntryPoint = "cuEventDestroy_v2")] + public static partial int cuEventDestroy(IntPtr hEvent); - [DllImport(LibName, EntryPoint = "cuEventRecord")] - public static extern int cuEventRecord(IntPtr hEvent, IntPtr hStream); + [LibraryImport(LibName, EntryPoint = "cuEventRecord")] + public static partial int cuEventRecord(IntPtr hEvent, IntPtr hStream); - [DllImport(LibName)] - public static extern int cuEventSynchronize(IntPtr hEvent); + [LibraryImport(LibName)] + public static partial int cuEventSynchronize(IntPtr hEvent); - [DllImport(LibName)] - public static extern int cuStreamWaitEvent(IntPtr hStream, IntPtr hEvent, uint flags); + [LibraryImport(LibName)] + public static partial int cuStreamWaitEvent(IntPtr hStream, IntPtr hEvent, uint flags); // ---- Peer-to-Peer (multi-GPU) ---- - [DllImport(LibName)] - public static extern int cuDeviceCanAccessPeer(out int canAccessPeer, int dev, int peerDev); + [LibraryImport(LibName)] + public static partial int cuDeviceCanAccessPeer(out int canAccessPeer, int dev, int peerDev); - [DllImport(LibName)] - public static extern int cuCtxEnablePeerAccess(IntPtr peerContext, uint flags); + [LibraryImport(LibName)] + public static partial int cuCtxEnablePeerAccess(IntPtr peerContext, uint flags); - [DllImport(LibName)] - public static extern int cuMemcpyPeerAsync( + [LibraryImport(LibName)] + public static partial int cuMemcpyPeerAsync( IntPtr dstDevice, IntPtr dstContext, IntPtr srcDevice, IntPtr srcContext, UIntPtr byteCount, IntPtr hStream); From d3e9c11682990fa6c433b3850c649d953b2f515a Mon Sep 17 00:00:00 2001 From: craig-b <3112503+craig-b@users.noreply.github.com> Date: Sat, 15 Aug 2026 14:31:40 +0100 Subject: [PATCH 5/5] Migrate remaining libc P/Invokes to LibraryImport Covers _exit (CLI), madvise (ModelBase), and mlock/munlock (GgufReader). The Qwen35 decode contract test now asserts the LibraryImport/UnmanagedCallConv attributes instead of DllImport. --- InferenceWeb.Tests/Qwen35TokenDecodeContractTests.cs | 8 +++++--- TensorSharp.Cli/Program.cs | 6 +++--- TensorSharp.Models/ModelBase.cs | 6 +++--- TensorSharp.Runtime/GgufReader.cs | 10 +++++----- 4 files changed, 16 insertions(+), 14 deletions(-) diff --git a/InferenceWeb.Tests/Qwen35TokenDecodeContractTests.cs b/InferenceWeb.Tests/Qwen35TokenDecodeContractTests.cs index e68b7f82..cffdabd3 100644 --- a/InferenceWeb.Tests/Qwen35TokenDecodeContractTests.cs +++ b/InferenceWeb.Tests/Qwen35TokenDecodeContractTests.cs @@ -167,10 +167,12 @@ public void DecodePInvoke_MatchesPublicManagedParameterContract( publicMethod.GetParameters().Select(parameter => parameter.ParameterType), nativeMethod.GetParameters().Select(parameter => parameter.ParameterType)); - DllImportAttribute import = nativeMethod.GetCustomAttribute()!; + LibraryImportAttribute import = nativeMethod.GetCustomAttribute()!; Assert.NotNull(import); - Assert.Equal("GgmlOps", import.Value); - Assert.Equal(CallingConvention.Cdecl, import.CallingConvention); + Assert.Equal("GgmlOps", import.LibraryName); + UnmanagedCallConvAttribute callConv = nativeMethod.GetCustomAttribute()!; + Assert.NotNull(callConv); + Assert.Contains(typeof(System.Runtime.CompilerServices.CallConvCdecl), callConv.CallConvs!); Assert.Equal(nativeMethodName, nativeMethod.Name); ParameterInfo reseedParameter = nativeMethod.GetParameters()[2]; diff --git a/TensorSharp.Cli/Program.cs b/TensorSharp.Cli/Program.cs index 12dfae86..a2a790b0 100644 --- a/TensorSharp.Cli/Program.cs +++ b/TensorSharp.Cli/Program.cs @@ -27,7 +27,7 @@ namespace TensorSharp.Cli { - class Program + partial class Program { private static readonly IPromptRenderer PromptRenderer = new GgufPromptRenderer(); private static ILogger _log = NullLogger.Instance; @@ -124,8 +124,8 @@ private static string SelectedBackend(string[] args) return backend; } - [System.Runtime.InteropServices.DllImport("libc", EntryPoint = "_exit")] - private static extern void LibcExit(int status); + [System.Runtime.InteropServices.LibraryImport("libc", EntryPoint = "_exit")] + private static partial void LibcExit(int status); static void MainCore(string[] args) { diff --git a/TensorSharp.Models/ModelBase.cs b/TensorSharp.Models/ModelBase.cs index f0e6281f..d04d9cbc 100644 --- a/TensorSharp.Models/ModelBase.cs +++ b/TensorSharp.Models/ModelBase.cs @@ -25,7 +25,7 @@ namespace TensorSharp.Models { - public class QuantizedWeight : IDisposable + public partial class QuantizedWeight : IDisposable { private IntPtr _data; private GCHandle _cacheKeyHandle; @@ -273,8 +273,8 @@ private static unsafe void AdviseExternalViewCanBePagedOut(IntPtr data, long byt private const int MadvDontNeed = 4; - [DllImport("libc", SetLastError = true, EntryPoint = "madvise")] - private static extern unsafe int madvise(void* addr, nuint len, int advice); + [LibraryImport("libc", EntryPoint = "madvise", SetLastError = true)] + private static unsafe partial int madvise(void* addr, nuint len, int advice); } /// diff --git a/TensorSharp.Runtime/GgufReader.cs b/TensorSharp.Runtime/GgufReader.cs index 332c0844..c32c88a6 100644 --- a/TensorSharp.Runtime/GgufReader.cs +++ b/TensorSharp.Runtime/GgufReader.cs @@ -56,7 +56,7 @@ public long NumElements } } - public class GgufFile : IDisposable + public partial class GgufFile : IDisposable { public uint Version { get; private set; } public Dictionary Metadata { get; } = new(); @@ -234,11 +234,11 @@ public void PrefaultFileCache() } } - [DllImport("libc", SetLastError = true, EntryPoint = "mlock")] - private static extern unsafe int mlock(void* addr, nuint len); + [LibraryImport("libc", EntryPoint = "mlock", SetLastError = true)] + private static unsafe partial int mlock(void* addr, nuint len); - [DllImport("libc", SetLastError = true, EntryPoint = "munlock")] - private static extern unsafe int munlock(void* addr, nuint len); + [LibraryImport("libc", EntryPoint = "munlock", SetLastError = true)] + private static unsafe partial int munlock(void* addr, nuint len); private void Parse() {