From d7ee80cd393fd8bafa24616c0edf168a1262651e Mon Sep 17 00:00:00 2001 From: Skyrion9 Date: Wed, 15 Jul 2026 05:44:23 +0300 Subject: [PATCH 1/5] chore: update ggml submodule to latest master. - Many fixes & perf updates in general. - Seems to have some important updates for vulkan including fp16: 1 -> fp16: dot2 which should be more efficient. "ggml_vulkan: 0 = AMD Radeon RX 6700 (RADV NAVI22) (radv) | uma: 0 | fp16: dot2 | bf16: 0 | warp size: 32 | shared memory: 65536 | int dot: 1 | matrix cores: none" --- ggml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml b/ggml index 57ea0bc..af97976 160000 --- a/ggml +++ b/ggml @@ -1 +1 @@ -Subproject commit 57ea0bc119d722d74594196cc5b494a34dd87be4 +Subproject commit af97976c7810cdabb1863172f31c432dab767de7 From b0b4b251412ee51bdb79990a33c04d2865cd261d Mon Sep 17 00:00:00 2001 From: Skyrion9 Date: Wed, 15 Jul 2026 05:45:46 +0300 Subject: [PATCH 2/5] chore: remove obsolete CUDA conv patches and disable auto patching by default. --- CMakeLists.txt | 2 +- ...-conv-transpose-1d-grid-fix-loop-fix.patch | 84 ------------------- patches/ggml-conv2d-dw-grid-fix.patch | 39 --------- 3 files changed, 1 insertion(+), 124 deletions(-) delete mode 100644 patches/ggml-conv-transpose-1d-grid-fix-loop-fix.patch delete mode 100644 patches/ggml-conv2d-dw-grid-fix.patch diff --git a/CMakeLists.txt b/CMakeLists.txt index f6d8a51..67c14e0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -13,7 +13,7 @@ option(S2_VULKAN "Build with Vulkan backend" OFF) option(S2_CUDA "Build with CUDA backend" OFF) option(S2_METAL "Build with Metal backend" OFF) option(S2_BUILD_SHARED_LIBRARIES "Build S2 shared and static library targets" OFF) -option(S2_AUTO_APPLY_LOCAL_PATCHES "Automatically apply local patches from patches/*.patch during build" ON) +option(S2_AUTO_APPLY_LOCAL_PATCHES "Automatically apply local patches from patches/*.patch during build" OFF) if(S2_METAL AND NOT APPLE) message(FATAL_ERROR "S2_METAL is only supported on Apple platforms.") diff --git a/patches/ggml-conv-transpose-1d-grid-fix-loop-fix.patch b/patches/ggml-conv-transpose-1d-grid-fix-loop-fix.patch deleted file mode 100644 index 5846840..0000000 --- a/patches/ggml-conv-transpose-1d-grid-fix-loop-fix.patch +++ /dev/null @@ -1,84 +0,0 @@ ---- a/ggml/src/ggml-cuda/conv-transpose-1d.cu -+++ b/ggml/src/ggml-cuda/conv-transpose-1d.cu -@@ -1,51 +1,53 @@ - #include "conv-transpose-1d.cuh" -+#include - - static __global__ void conv_transpose_1d_kernel( -- const int s0, const int p0, const int d0, const int output_size, -+ const int s0, const int p0, const int d0, const int64_t output_size, - const int src0_ne0, const int src0_ne1, const int src0_ne2, const int src0_ne3, - const int src1_ne0, const int src1_ne1, const int src1_ne2, const int src1_ne3, - const int dst_ne0, const int dst_ne1, const int dst_ne2, const int dst_ne3, - const float * src0, const float * src1, float * dst) { -- int global_index = threadIdx.x + blockIdx.x * blockDim.x; -- if (global_index >= output_size) { -- return; -- } - -- int out_index = global_index / dst_ne0; -+ for (int64_t global_index = threadIdx.x + blockIdx.x * blockDim.x; global_index < output_size; global_index += blockDim.x * gridDim.x) { -+ int out_index = global_index / dst_ne0; - -- float accumulator = 0; -+ float accumulator = 0; - -- for (int c = 0; c < src0_ne2; c++) { -- int idx = global_index % dst_ne0; -+ for (int c = 0; c < src0_ne2; c++) { -+ int idx = global_index % dst_ne0; - -- int kernel_offset = (src0_ne0 * src0_ne1 * c) + (out_index * src0_ne0); -- int input_offset = src1_ne0 * c; -+ int kernel_offset = (src0_ne0 * src0_ne1 * c) + (out_index * src0_ne0); -+ int input_offset = src1_ne0 * c; - -- for (int i = 0; i < src1_ne0; i++) { -- if (!(idx >= i*s0 && idx < i*s0 + src0_ne0)) { -- continue; -- } -- int weight_idx = idx - i*s0; - -- float kernel_weight = src0[kernel_offset + weight_idx]; -- float input_value = src1[input_offset+i]; -+ int start_i = (idx >= src0_ne0) ? (idx - src0_ne0 + s0) / s0 : 0; -+ int end_i = idx / s0; -+ -+ if (end_i >= src1_ne0) end_i = src1_ne0 - 1; - -- accumulator += kernel_weight * input_value; -+ for (int i = start_i; i <= end_i; i++) { -+ int weight_idx = idx - i*s0; -+ float kernel_weight = src0[kernel_offset + weight_idx]; -+ float input_value = src1[input_offset+i]; -+ -+ accumulator += kernel_weight * input_value; -+ } - } -+ dst[global_index] = accumulator; - } -- dst[global_index] = accumulator; - GGML_UNUSED_VARS(p0, d0, src0_ne3, src1_ne3, dst_ne3, src1_ne1, dst_ne1, src1_ne2, dst_ne2); - } - - static void conv_transpose_1d_f32_f32_cuda( -- const int s0, const int p0, const int d0, const int output_size, -+ const int s0, const int p0, const int d0, const int64_t output_size, - const int src0_ne0, const int src0_ne1, const int src0_ne2, const int src0_ne3, - const int src1_ne0, const int src1_ne1, const int src1_ne2, const int src1_ne3, - const int dst_ne0, const int dst_ne1, const int dst_ne2, const int dst_ne3, - const float * src0, const float * src1, float * dst, - cudaStream_t stream) { - -- const int num_blocks = (output_size + CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE - 1) / CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE; -+ const int num_blocks = std::min((int64_t)65535, (output_size + CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE - 1) / CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE); -+ - conv_transpose_1d_kernel<<>>( - s0,p0,d0,output_size, - src0_ne0, src0_ne1, src0_ne2, src0_ne3, -@@ -83,4 +85,4 @@ void ggml_cuda_op_conv_transpose_1d(ggml_backend_cuda_context & ctx, ggml_tensor - src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3], - dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3], - src0_d, src1_d, dst_d, stream); --} -+} diff --git a/patches/ggml-conv2d-dw-grid-fix.patch b/patches/ggml-conv2d-dw-grid-fix.patch deleted file mode 100644 index 2494247..0000000 --- a/patches/ggml-conv2d-dw-grid-fix.patch +++ /dev/null @@ -1,39 +0,0 @@ ---- a/ggml/src/ggml-cuda/conv2d-dw.cu -+++ b/ggml/src/ggml-cuda/conv2d-dw.cu -@@ -1,4 +1,5 @@ - #include "conv2d-dw.cuh" -+#include - - struct conv_params { - int in_w, in_h; -@@ -84,12 +85,9 @@ __global__ void conv2d_dw_kernel(const T * __restrict__ input, const T * __restr - const int kernel_w, const int kernel_h, const int stride_x, const int stride_y, - const int padding_x, const int padding_y, const int dilation_x, const int dilation_y, - const int channels, const int batches) { -- const int global_idx = blockIdx.x * blockDim.x + threadIdx.x; - const int total_elements = batches * channels * out_h * out_w; - -- if (global_idx >= total_elements) { -- return; -- } -+ for (int global_idx = blockIdx.x * blockDim.x + threadIdx.x; global_idx < total_elements; global_idx += blockDim.x * gridDim.x) { - - conv_params params = { in_w, in_h, out_w, out_h, kernel_w, kernel_h, stride_x, - stride_y, padding_x, padding_y, dilation_x, dilation_y, channels, batches }; -@@ -114,6 +112,7 @@ __global__ void conv2d_dw_kernel(const T * __restrict__ input, const T * __restr - } - - output[Layout::output_index(batch_idx, channel_idx, out_y_idx, out_x_idx, params)] = accumulator; -+ } - } - - void ggml_cuda_op_conv2d_dw(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { -@@ -145,7 +144,7 @@ void ggml_cuda_op_conv2d_dw(ggml_backend_cuda_context & ctx, ggml_tensor * dst) - cudaStream_t st = ctx.stream(); - - const int total = batches * channels * out_h * out_w; -- const int blocks = (total + CUDA_CONV2D_DW_BLOCK_SIZE - 1) / CUDA_CONV2D_DW_BLOCK_SIZE; -+ const int blocks = std::min(65535, (total + CUDA_CONV2D_DW_BLOCK_SIZE - 1) / CUDA_CONV2D_DW_BLOCK_SIZE); - - if (ggml_is_contiguous(input)) { - conv2d_dw_kernel<<>>( From 584ad64d4a5b30890020559d2b11e400f4654dbf Mon Sep 17 00:00:00 2001 From: Skyrion9 Date: Sat, 18 Jul 2026 03:36:44 +0300 Subject: [PATCH 3/5] chore: Update ggml to latest master (0.17.0) --- ggml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml b/ggml index af97976..9be3133 160000 --- a/ggml +++ b/ggml @@ -1 +1 @@ -Subproject commit af97976c7810cdabb1863172f31c432dab767de7 +Subproject commit 9be313313c8ecb9488911bd64550190e3ed80f38 From 71955cf7816c1d6aad92d5d390da98c96a0b3d97 Mon Sep 17 00:00:00 2001 From: Skyrion9 Date: Tue, 4 Aug 2026 20:22:39 +0300 Subject: [PATCH 4/5] chore: Update ggml to latest master (0.18.0) --- ggml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml b/ggml index 9be3133..78de606 160000 --- a/ggml +++ b/ggml @@ -1 +1 @@ -Subproject commit 9be313313c8ecb9488911bd64550190e3ed80f38 +Subproject commit 78de606907c364b76601f734021f76bd29ca5637 From f95fc3e94f4deebdf7a737c104363c099e5fb8ae Mon Sep 17 00:00:00 2001 From: Skyrion9 Date: Tue, 4 Aug 2026 20:49:42 +0300 Subject: [PATCH 5/5] chore: Update ggml to latest master (0.18.1) --- ggml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml b/ggml index 78de606..90951f9 160000 --- a/ggml +++ b/ggml @@ -1 +1 @@ -Subproject commit 78de606907c364b76601f734021f76bd29ca5637 +Subproject commit 90951f99af1fbebef3fbdd58ff5b8715b0bb9c43