Added new binaries.

- Built by this run: https://github.com/SciSharp/LLamaSharp/actions/runs/6921572568
 - commit: `e937066420b79a757bf80e9836eb12b88420a218`
 - Rearranged paths

LLama/LLamaSharp.Runtime.targets

@@ -3,41 +3,69 @@
         <IncludeBuiltInRuntimes Condition="'$(IncludeBuiltInRuntimes)' == ''">true</IncludeBuiltInRuntimes>
     </PropertyGroup>
     <ItemGroup Condition="'$(IncludeBuiltInRuntimes)' == 'true'">
-        <None Include="$(MSBuildThisFileDirectory)runtimes/libllama.dll">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/win-x64/native/libllama.dll</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/libllama-cuda11.dll">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/win-x64/native/cuda11/libllama.dll</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/libllama-cuda12.dll">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/win-x64/native/cuda12/libllama.dll</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/libllama.so">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/linux-x64/native/libllama.so</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/libllama-cuda11.so">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/linux-x64/native/cuda11/libllama.so</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/libllama-cuda12.so">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/linux-x64/native/cuda12/libllama.so</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/osx-arm64/libllama.dylib">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/osx-arm64/native/libllama.dylib</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/osx-arm64/ggml-metal.metal">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/osx-arm64/native/ggml-metal.metal</Link>
-        </None>
-        <None Include="$(MSBuildThisFileDirectory)runtimes/osx-x64/libllama.dylib">
-            <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
-            <Link>runtimes/osx-x64/native/libllama.dylib</Link>
-        </None>
+
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/libllama.dll">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/win-x64/noavx/libllama.dll</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/avx/libllama.dll">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/win-x64/avx/libllama.dll</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/avx2/libllama.dll">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/win-x64/avx2/libllama.dll</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/avx512/libllama.dll">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/win-x64/avx512/libllama.dll</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/cu11.7.1/libllama.dll">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/win-x64/cuda11/libllama.dll</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/cu12.1.0/libllama.dll">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/win-x64/cuda12/libllama.dll</Link>
+      </None>
+
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/libllama.so">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/linux-x64/noavx/libllama.so</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/avx/libllama.so">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/linux-x64/avx/libllama.so</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/avx2/libllama.so">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/linux-x64/avx2/libllama.so</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/avx512/libllama.so">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/linux-x64/avx512/libllama.so</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/cu11.7.1/libllama.so">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/linux-x64/cuda11/libllama.so</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/cu12.1.0/libllama.so">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/linux-x64/cuda12/libllama.so</Link>
+      </None>
+
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/osx-arm64/libllama.dylib">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/osx-arm64/libllama.dylib</Link>
+      </None>
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/osx-arm64/ggml-metal.metal">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/osx-arm64/ggml-metal.metal</Link>
+      </None>
+
+      <None Include="$(MSBuildThisFileDirectory)runtimes/deps/osx-x64/libllama.dylib">
+        <CopyToOutputDirectory>PreserveNewest</CopyToOutputDirectory>
+        <Link>runtimes/osx-x64/libllama.dylib</Link>
+      </None>
     </ItemGroup>
 </Project>

LLama/Native/NativeApi.Load.cs

@@ -155,13 +155,13 @@ namespace LLama.Native
             if (RuntimeInformation.IsOSPlatform(OSPlatform.Windows))
             {
                 platform = OSPlatform.Windows;
-                prefix = "runtimes/win-x64/native/";
+                prefix = "runtimes/win-x64/";
                 suffix = ".dll";
             }
             else if (RuntimeInformation.IsOSPlatform(OSPlatform.Linux))
             {
                 platform = OSPlatform.Linux;
-                prefix = "runtimes/linux-x64/native/";
+                prefix = "runtimes/linux-x64/";
                 suffix = ".so";
             }
             else if (RuntimeInformation.IsOSPlatform(OSPlatform.OSX))
@@ -170,11 +170,11 @@ namespace LLama.Native
                 suffix = ".dylib";
                 if (System.Runtime.Intrinsics.Arm.ArmBase.Arm64.IsSupported)
                 {
-                    prefix = "runtimes/osx-arm64/native/";
+                    prefix = "runtimes/osx-arm64/";
                 }
                 else
                 {
-                    prefix = "runtimes/osx-x64/native/";
+                    prefix = "runtimes/osx-x64/";
                 }
             }
             else

LLama/Native/NativeApi.cs

@@ -348,18 +348,77 @@ namespace LLama.Native
         public static extern int llama_n_embd(SafeLlamaModelHandle model);
 
         /// <summary>
-        /// Get the size of the model in bytes
+        /// Get the model's RoPE frequency scaling factor
+        /// </summary>
+        /// <param name="model"></param>
+        /// <returns></returns>
+        [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
+        public static extern float llama_rope_freq_scale_train(SafeLlamaModelHandle model);
+
+        /// <summary>
+        /// Get metadata value as a string by key name
+        /// </summary>
+        /// <param name="model"></param>
+        /// <param name="key"></param>
+        /// <param name="buf"></param>
+        /// <param name="buf_size"></param>
+        /// <returns></returns>
+        [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
+        public static extern int llama_model_meta_val_str(SafeLlamaModelHandle model, char* key, char* buf, long buf_size);
+
+        /// <summary>
+        /// Get the number of metadata key/value pairs
+        /// </summary>
+        /// <param name="model"></param>
+        /// <returns></returns>
+        [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
+        public static extern int llama_model_meta_count(SafeLlamaModelHandle model);
+
+        /// <summary>
+        /// Get metadata key name by index
         /// </summary>
         /// <param name="model"></param>
+        /// <param name="i"></param>
+        /// <param name="buf"></param>
+        /// <param name="buf_size"></param>
         /// <returns></returns>
         [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
+        public static extern int llama_model_meta_key_by_index(SafeLlamaModelHandle model, int i, char* buf, long buf_size);
+
+        /// <summary>
+        /// Get metadata value as a string by index
+        /// </summary>
+        /// <param name="model"></param>
+        /// <param name="i"></param>
+        /// <param name="buf"></param>
+        /// <param name="buf_size"></param>
+        /// <returns>The functions return the length of the string on success, or -1 on failure</returns>
+        [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
+        public static extern int llama_model_meta_val_str_by_index(SafeLlamaModelHandle model, int i, char* buf, long buf_size);
+
+        /// <summary>
+        /// Get a string describing the model type
+        /// </summary>
+        /// <param name="model"></param>
+        /// <param name="buf"></param>
+        /// <param name="buf_size"></param>
+        /// <returns>The functions return the length of the string on success, or -1 on failure</returns>
+        [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
+        public static extern int llama_model_desc(SafeLlamaModelHandle model, char* buf, long buf_size);
+
+        /// <summary>
+        /// Get the size of the model in bytes
+        /// </summary>
+        /// <param name="model"></param>
+        /// <returns>The functions return the length of the string on success, or -1 on failure</returns>
+        [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
         public static extern ulong llama_model_size(SafeLlamaModelHandle model);
 
         /// <summary>
         /// Get the number of parameters in this model
         /// </summary>
         /// <param name="model"></param>
-        /// <returns></returns>
+        /// <returns>The functions return the length of the string on success, or -1 on failure</returns>
         [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
         public static extern ulong llama_model_n_params(SafeLlamaModelHandle model);
 
@@ -370,7 +429,7 @@ namespace LLama.Native
         /// <param name="llamaToken"></param>
         /// <param name="buffer">buffer to write string into</param>
         /// <param name="length">size of the buffer</param>
-        /// <returns>The length writte, or if the buffer is too small a negative that indicates the length required</returns>
+        /// <returns>The length written, or if the buffer is too small a negative that indicates the length required</returns>
         [DllImport(libraryName, CallingConvention = CallingConvention.Cdecl)]
         public static extern int llama_token_to_piece(SafeLlamaModelHandle model, int llamaToken, byte* buffer, int length);
 

LLama/runtimes/osx-arm64/ggml-metal.metal → LLama/runtimes/deps/osx-arm64/ggml-metal.metal

@@ -792,7 +792,7 @@ kernel void kernel_mul_mv_f32_f32(
         constant   int64_t & ne0,
         constant   int64_t & ne1,
         uint3 tgpig[[threadgroup_position_in_grid]],
-        uint tiisg[[thread_index_in_simdgroup]]) {
+        uint  tiisg[[thread_index_in_simdgroup]]) {
 
     const int64_t r0 = tgpig.x;
     const int64_t rb = tgpig.y*N_F32_F32;
@@ -844,6 +844,79 @@ kernel void kernel_mul_mv_f32_f32(
     }
 }
 
+#define N_F16_F16 4
+
+kernel void kernel_mul_mv_f16_f16(
+        device const  char * src0,
+        device const  char * src1,
+        device       float * dst,
+        constant   int64_t & ne00,
+        constant   int64_t & ne01,
+        constant   int64_t & ne02,
+        constant  uint64_t & nb00,
+        constant  uint64_t & nb01,
+        constant  uint64_t & nb02,
+        constant   int64_t & ne10,
+        constant   int64_t & ne11,
+        constant   int64_t & ne12,
+        constant  uint64_t & nb10,
+        constant  uint64_t & nb11,
+        constant  uint64_t & nb12,
+        constant   int64_t & ne0,
+        constant   int64_t & ne1,
+        uint3 tgpig[[threadgroup_position_in_grid]],
+        uint  tiisg[[thread_index_in_simdgroup]]) {
+
+    const int64_t r0 = tgpig.x;
+    const int64_t rb = tgpig.y*N_F16_F16;
+    const int64_t im = tgpig.z;
+
+    device const half * x = (device const half *) (src0 + r0*nb01 + im/(ne12/ne02)*nb02);
+
+    if (ne00 < 128) {
+        for (int row = 0; row < N_F16_F16; ++row) {
+            int r1 = rb + row;
+            if (r1 >= ne11) {
+                break;
+            }
+
+            device const half * y = (device const half *) (src1 + r1*nb11 + im*nb12);
+
+            float sumf = 0;
+            for (int i = tiisg; i < ne00; i += 32) {
+                sumf += (half) x[i] * (half) y[i];
+            }
+
+            float all_sum = simd_sum(sumf);
+            if (tiisg == 0) {
+                dst[im*ne1*ne0 + r1*ne0 + r0] = all_sum;
+            }
+        }
+    } else {
+        device const half4 * x4 = (device const half4 *)x;
+        for (int row = 0; row < N_F16_F16; ++row) {
+            int r1 = rb + row;
+            if (r1 >= ne11) {
+                break;
+            }
+
+            device const half  * y  = (device const half  *) (src1 + r1*nb11 + im*nb12);
+            device const half4 * y4 = (device const half4 *) y;
+
+            float sumf = 0;
+            for (int i = tiisg; i < ne00/4; i += 32) {
+                for (int k = 0; k < 4; ++k) sumf += (half) x4[i][k] * y4[i][k];
+            }
+
+            float all_sum = simd_sum(sumf);
+            if (tiisg == 0) {
+                for (int i = 4*(ne00/4); i < ne00; ++i) all_sum += (half) x[i] * y[i];
+                dst[im*ne1*ne0 + r1*ne0 + r0] = all_sum;
+            }
+        }
+    }
+}
+
 kernel void kernel_mul_mv_f16_f32_1row(
         device const  char * src0,
         device const  char * src1,
@@ -1229,6 +1302,39 @@ kernel void kernel_rope(
 template [[host_name("kernel_rope_f32")]] kernel rope_t kernel_rope<float>;
 template [[host_name("kernel_rope_f16")]] kernel rope_t kernel_rope<half>;
 
+kernel void kernel_im2col_f16(
+        device const float * x,
+        device       half * dst,
+        constant   int32_t & ofs0,
+        constant   int32_t & ofs1,
+        constant   int32_t & IW,
+        constant   int32_t & IH,
+        constant   int32_t & CHW,
+        constant   int32_t & s0,
+        constant   int32_t & s1,
+        constant   int32_t & p0,
+        constant   int32_t & p1,
+        constant   int32_t & d0,
+        constant   int32_t & d1,
+        uint3 tgpig[[threadgroup_position_in_grid]],
+        uint3  tgpg[[threadgroups_per_grid]],
+        uint3 tpitg[[thread_position_in_threadgroup]],
+        uint3   ntg[[threads_per_threadgroup]]) {
+    const int32_t iiw = tgpig[2] * s0 + tpitg[2] * d0 - p0;
+    const int32_t iih = tgpig[1] * s1 + tpitg[1] * d1 - p1;
+
+    const int32_t offset_dst =
+        (tpitg[0] * tgpg[1] * tgpg[2] + tgpig[1] * tgpg[2] + tgpig[2]) * CHW +
+        (tgpig[0] * (ntg[1] * ntg[2]) + tpitg[1] * ntg[2] + tpitg[2]);
+
+    if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW) {
+        dst[offset_dst] = 0.0f;
+    } else {
+        const int32_t offset_src = tpitg[0] * ofs0 + tgpig[0] * ofs1;
+        dst[offset_dst] = x[offset_src + iih * IW + iiw];
+    }
+}
+
 kernel void kernel_cpy_f16_f16(
         device const half * src0,
         device       half * dst,