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()
{