From 69cb03d57d46b22e9f64bb9e76242b193011fea5 Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Tue, 7 Apr 2026 23:49:03 +0900 Subject: [PATCH 1/9] metal : add CONV_2D_DW (depthwise 2D convolution) support --- ggml/src/ggml-metal/ggml-metal-device.cpp | 21 +++++ ggml/src/ggml-metal/ggml-metal-device.h | 1 + ggml/src/ggml-metal/ggml-metal-device.m | 4 + ggml/src/ggml-metal/ggml-metal-impl.h | 28 ++++++ ggml/src/ggml-metal/ggml-metal-ops.cpp | 81 +++++++++++++++++ ggml/src/ggml-metal/ggml-metal-ops.h | 1 + ggml/src/ggml-metal/ggml-metal.metal | 106 ++++++++++++++++++++++ 7 files changed, 242 insertions(+) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index e7ac21a2be29..0868d6078e65 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1869,6 +1869,27 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d(ggml_met return res; } +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw(ggml_metal_library_t lib, const ggml_tensor * op) { + assert(op->op == GGML_OP_CONV_2D_DW); + + GGML_ASSERT(op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32); + GGML_ASSERT(op->src[1]->type == GGML_TYPE_F32); + GGML_ASSERT(op->type == GGML_TYPE_F32); + + char base[256]; + char name[256]; + + snprintf(base, 256, "kernel_conv_2d_dw_%s_%s", ggml_type_name(op->src[0]->type), ggml_type_name(op->src[1]->type)); + snprintf(name, 256, "%s", base); + + ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); + if (!res.pipeline) { + res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); + } + + return res; +} + ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d(ggml_metal_library_t lib, const ggml_tensor * op) { assert(op->op == GGML_OP_CONV_3D); diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index dc75a34b2b77..c2484f82f9b5 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -152,6 +152,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_tran struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_upscale (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pad (ggml_metal_library_t lib, const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 962a001f39f8..5d29250f654b 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1198,6 +1198,10 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 && (op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32); + case GGML_OP_CONV_2D_DW: + return op->src[1]->type == GGML_TYPE_F32 && + op->type == GGML_TYPE_F32 && + (op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32); case GGML_OP_UPSCALE: return op->src[0]->type == GGML_TYPE_F32; case GGML_OP_POOL_1D: diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 6bf61423ce73..d6761023b76c 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -656,6 +656,34 @@ typedef struct { int32_t d1; } ggml_metal_kargs_conv_2d; +typedef struct { + uint64_t nb00; // kernel strides + uint64_t nb01; + uint64_t nb02; + uint64_t nb10; // input strides + uint64_t nb11; + uint64_t nb12; + uint64_t nb13; + uint64_t nb0; // output strides + uint64_t nb1; + uint64_t nb2; + uint64_t nb3; + int32_t IW; // input width + int32_t IH; // input height + int32_t KW; // kernel width + int32_t KH; // kernel height + int32_t C; // channels (IC == OC for depthwise) + int32_t OW; // output width + int32_t OH; // output height + int32_t N; // batch size + int32_t s0; // stride x + int32_t s1; // stride y + int32_t p0; // padding x + int32_t p1; // padding y + int32_t d0; // dilation x + int32_t d1; // dilation y +} ggml_metal_kargs_conv_2d_dw; + typedef struct { uint64_t ofs0; uint64_t ofs1; diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index d196bae4faf2..288ef1821433 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -387,6 +387,10 @@ static int ggml_metal_op_encode_impl(ggml_metal_op_t ctx, int idx) { { n_fuse = ggml_metal_op_conv_2d(ctx, idx); } break; + case GGML_OP_CONV_2D_DW: + { + n_fuse = ggml_metal_op_conv_2d_dw(ctx, idx); + } break; case GGML_OP_CONV_TRANSPOSE_1D: { n_fuse = ggml_metal_op_conv_transpose_1d(ctx, idx); @@ -3742,6 +3746,83 @@ int ggml_metal_op_conv_2d(ggml_metal_op_t ctx, int idx) { return 1; } +int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { + ggml_tensor * op = ctx->node(idx); + + ggml_metal_library_t lib = ctx->lib; + ggml_metal_encoder_t enc = ctx->enc; + + GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne); + GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb); + GGML_TENSOR_LOCALS( int32_t, ne1, op->src[1], ne); + GGML_TENSOR_LOCALS(uint64_t, nb1, op->src[1], nb); + GGML_TENSOR_LOCALS( int32_t, ne, op, ne); + GGML_TENSOR_LOCALS(uint64_t, nb, op, nb); + + GGML_ASSERT(op->src[1]->type == GGML_TYPE_F32); + GGML_ASSERT(op->type == GGML_TYPE_F32); + GGML_ASSERT(op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32); + + const int32_t s0 = ((const int32_t *) op->op_params)[0]; + const int32_t s1 = ((const int32_t *) op->op_params)[1]; + const int32_t p0 = ((const int32_t *) op->op_params)[2]; + const int32_t p1 = ((const int32_t *) op->op_params)[3]; + const int32_t d0 = ((const int32_t *) op->op_params)[4]; + const int32_t d1 = ((const int32_t *) op->op_params)[5]; + + // Kernel shape: [KW, KH, 1, C] — nb03 is the per-channel stride + // Since ne02 == 1, nb03 == nb02, but we use nb03 explicitly for clarity + ggml_metal_kargs_conv_2d_dw args = { + /*.nb00 =*/ nb00, + /*.nb01 =*/ nb01, + /*.nb02 =*/ nb03, // per-channel kernel stride (dim 3, since dim 2 == 1) + /*.nb10 =*/ nb10, + /*.nb11 =*/ nb11, + /*.nb12 =*/ nb12, + /*.nb13 =*/ nb13, + /*.nb0 =*/ nb0, + /*.nb1 =*/ nb1, + /*.nb2 =*/ nb2, + /*.nb3 =*/ nb3, + /*.IW =*/ ne10, + /*.IH =*/ ne11, + /*.KW =*/ ne00, + /*.KH =*/ ne01, + /*.C =*/ ne12, // input channels == output channels for depthwise + /*.OW =*/ ne0, + /*.OH =*/ ne1, + /*.N =*/ ne13, + /*.s0 =*/ s0, + /*.s1 =*/ s1, + /*.p0 =*/ p0, + /*.p1 =*/ p1, + /*.d0 =*/ d0, + /*.d1 =*/ d1, + }; + + auto pipeline = ggml_metal_library_get_pipeline_conv_2d_dw(lib, op); + + int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); + nth = std::min(nth, 256); + nth = std::max(nth, 1); + + const uint64_t n_out = ggml_nelements(op); + + uint64_t tg = (n_out + nth - 1)/nth; + tg = std::max(tg, 1); + tg = std::min(tg, (uint64_t) std::numeric_limits::max()); + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); + + ggml_metal_encoder_dispatch_threadgroups(enc, tg, 1, 1, nth, 1, 1); + + return 1; +} + int ggml_metal_op_conv_3d(ggml_metal_op_t ctx, int idx) { ggml_tensor * op = ctx->node(idx); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index 13f5274c7953..0bebd836a185 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -75,6 +75,7 @@ int ggml_metal_op_norm (ggml_metal_op_t ctx, int idx); int ggml_metal_op_rope (ggml_metal_op_t ctx, int idx); int ggml_metal_op_im2col (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_2d (ggml_metal_op_t ctx, int idx); +int ggml_metal_op_conv_2d_dw (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_3d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_transpose_1d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_transpose_2d (ggml_metal_op_t ctx, int idx); diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index dcb6803f5f6a..64dca44e6e39 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -4908,6 +4908,112 @@ kernel void kernel_conv_2d( uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]); +template +kernel void kernel_conv_2d_dw( + constant ggml_metal_kargs_conv_2d_dw & args, + device const char * weights, + device const char * src, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + uint3 tgpg[[threadgroups_per_grid]], + uint3 tpitg[[thread_position_in_threadgroup]], + uint3 ntg[[threads_per_threadgroup]]) { + + const uint threads_per_tg = ntg.x * ntg.y * ntg.z; + const uint tg_index = (tgpig.z * tgpg.y + tgpig.y) * tgpg.x + tgpig.x; + const uint local_thread = tpitg.z * (ntg.x * ntg.y) + tpitg.y * ntg.x + tpitg.x; + const uint thread_index = tg_index * threads_per_tg + local_thread; + const uint64_t total_threads = (uint64_t) threads_per_tg * tgpg.x * tgpg.y * tgpg.z; + const uint64_t total_outputs = (uint64_t) args.N * args.C * args.OH * args.OW; + + for (uint64_t index = thread_index; index < total_outputs; index += total_threads) { + uint64_t tmp = index; + + const int32_t ow = tmp % args.OW; tmp /= args.OW; + const int32_t oh = tmp % args.OH; tmp /= args.OH; + const int32_t c = tmp % args.C; tmp /= args.C; + const int32_t n = tmp; + + float acc = 0.0f; + + const int32_t base_x = ow*args.s0 - args.p0; + const int32_t base_y = oh*args.s1 - args.p1; + + int32_t ky_start = 0; + if (base_y < 0) { + ky_start = (-base_y + args.d1 - 1)/args.d1; + } + int32_t ky_end = args.KH; + const int32_t y_max = args.IH - 1 - base_y; + if (y_max < 0) { + ky_end = ky_start; + } else if (base_y + (args.KH - 1)*args.d1 >= args.IH) { + ky_end = min(ky_end, y_max/args.d1 + 1); + } + + int32_t kx_start = 0; + if (base_x < 0) { + kx_start = (-base_x + args.d0 - 1)/args.d0; + } + int32_t kx_end = args.KW; + const int32_t x_max = args.IW - 1 - base_x; + if (x_max < 0) { + kx_end = kx_start; + } else if (base_x + (args.KW - 1)*args.d0 >= args.IW) { + kx_end = min(kx_end, x_max/args.d0 + 1); + } + + if (ky_start < ky_end && kx_start < kx_end) { + // Depthwise: kernel indexed by channel c (via ne[3] dimension), no IC loop + const uint64_t w_base = (uint64_t) c * args.nb02; + const uint64_t src_base = (uint64_t) n * args.nb13 + (uint64_t) c * args.nb12; + + for (int32_t ky = ky_start; ky < ky_end; ++ky) { + const int32_t iy = base_y + ky*args.d1; + const uint64_t src_row = src_base + (uint64_t) iy * args.nb11; + const uint64_t w_row = w_base + (uint64_t) ky * args.nb01; + + for (int32_t kx = kx_start; kx < kx_end; ++kx) { + const int32_t ix = base_x + kx*args.d0; + const float x = *(device const float *)(src + src_row + (uint64_t) ix * args.nb10); + const float w = (float)(*(device const TK *)(weights + w_row + (uint64_t) kx * args.nb00)); + acc += x * w; + } + } + } + + const uint64_t dst_offs = + (uint64_t) n * args.nb3 + + (uint64_t) c * args.nb2 + + (uint64_t) oh * args.nb1 + + (uint64_t) ow * args.nb0; + + *(device float *)(dst + dst_offs) = acc; + } +} + +template [[host_name("kernel_conv_2d_dw_f32_f32")]] +kernel void kernel_conv_2d_dw( + constant ggml_metal_kargs_conv_2d_dw & args, + device const char * weights, + device const char * src, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + uint3 tgpg[[threadgroups_per_grid]], + uint3 tpitg[[thread_position_in_threadgroup]], + uint3 ntg[[threads_per_threadgroup]]); + +template [[host_name("kernel_conv_2d_dw_f16_f32")]] +kernel void kernel_conv_2d_dw( + constant ggml_metal_kargs_conv_2d_dw & args, + device const char * weights, + device const char * src, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + uint3 tgpg[[threadgroups_per_grid]], + uint3 tpitg[[thread_position_in_threadgroup]], + uint3 ntg[[threads_per_threadgroup]]); + typedef void (conv_transpose_1d_t)( constant ggml_metal_kargs_conv_transpose_1d & args, device const float * src0, From dc964ac2dd5ec16eed421b1481eee0bd3c50b1a1 Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Wed, 8 Apr 2026 00:15:20 +0900 Subject: [PATCH 2/9] test : add perf cases for CONV_2D_DW --- tests/test-backend-ops.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 0c0aac2bba76..6f9d3ce21043 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9623,6 +9623,10 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, 1, 1, 1, false)); test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, 1, 1, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({112, 112, 32, 1}, {3, 3, 1, 32}, 1, 1, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({112, 112, 32, 1}, {3, 3, 1, 32}, 1, 1, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({56, 56, 128, 1}, {5, 5, 1, 128}, 2, 2, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({56, 56, 128, 1}, {5, 5, 1, 128}, 2, 2, 1, true)); for (ggml_type kernel_type : {GGML_TYPE_F32, GGML_TYPE_F16}) { test_cases.emplace_back(new test_conv_transpose_2d({256, 256, 256, 1}, {3, 3, 16, 256}, 1, kernel_type)); From b897362427f9705d6ed6b12acfa4eac2a32ea3b5 Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Sat, 27 Jun 2026 22:59:28 +0900 Subject: [PATCH 3/9] metal : use 3D dispatch for CONV_2D_DW kernel --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 14 +-- ggml/src/ggml-metal/ggml-metal.metal | 120 ++++++++++++------------- 2 files changed, 66 insertions(+), 68 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 288ef1821433..91720bea648a 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3806,11 +3806,15 @@ int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { nth = std::min(nth, 256); nth = std::max(nth, 1); - const uint64_t n_out = ggml_nelements(op); + // 3D grid: x = OW (tiled by nth), y = OH, z = C * N + const int32_t OW = ne0; + const int32_t OH = ne1; + const int32_t C = ne12; + const int32_t N = ne13; - uint64_t tg = (n_out + nth - 1)/nth; - tg = std::max(tg, 1); - tg = std::min(tg, (uint64_t) std::numeric_limits::max()); + const int tg_x = (OW + nth - 1) / nth; + const int tg_y = OH; + const int tg_z = C * N; ggml_metal_encoder_set_pipeline(enc, pipeline); ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); @@ -3818,7 +3822,7 @@ int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); - ggml_metal_encoder_dispatch_threadgroups(enc, tg, 1, 1, nth, 1, 1); + ggml_metal_encoder_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1); return 1; } diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 64dca44e6e39..7651eb02478f 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -4915,81 +4915,77 @@ kernel void kernel_conv_2d_dw( device const char * src, device char * dst, uint3 tgpig[[threadgroup_position_in_grid]], - uint3 tgpg[[threadgroups_per_grid]], uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]) { - const uint threads_per_tg = ntg.x * ntg.y * ntg.z; - const uint tg_index = (tgpig.z * tgpg.y + tgpig.y) * tgpg.x + tgpig.x; - const uint local_thread = tpitg.z * (ntg.x * ntg.y) + tpitg.y * ntg.x + tpitg.x; - const uint thread_index = tg_index * threads_per_tg + local_thread; - const uint64_t total_threads = (uint64_t) threads_per_tg * tgpg.x * tgpg.y * tgpg.z; - const uint64_t total_outputs = (uint64_t) args.N * args.C * args.OH * args.OW; + // 3D grid: x = OW tile, y = OH, z = C * N + const int32_t oh = tgpig.y; + const int32_t cn = tgpig.z; + const int32_t c = cn % args.C; + const int32_t n = cn / args.C; - for (uint64_t index = thread_index; index < total_outputs; index += total_threads) { - uint64_t tmp = index; + // Precompute vertical bounds (shared by all threads in this row) + const int32_t base_y = oh*args.s1 - args.p1; - const int32_t ow = tmp % args.OW; tmp /= args.OW; - const int32_t oh = tmp % args.OH; tmp /= args.OH; - const int32_t c = tmp % args.C; tmp /= args.C; - const int32_t n = tmp; + int32_t ky_start = 0; + if (base_y < 0) { + ky_start = (-base_y + args.d1 - 1)/args.d1; + } + int32_t ky_end = args.KH; + const int32_t y_max = args.IH - 1 - base_y; + if (y_max < 0) { + ky_end = ky_start; + } else if (base_y + (args.KH - 1)*args.d1 >= args.IH) { + ky_end = min(ky_end, y_max/args.d1 + 1); + } - float acc = 0.0f; + const uint64_t w_base = (uint64_t) c * args.nb02; + const uint64_t src_base = (uint64_t) n * args.nb13 + (uint64_t) c * args.nb12; - const int32_t base_x = ow*args.s0 - args.p0; - const int32_t base_y = oh*args.s1 - args.p1; + // Each thread handles one OW position + const int32_t ow = (int32_t)(tgpig.x * ntg.x + tpitg.x); + if (ow >= args.OW) { + return; + } - int32_t ky_start = 0; - if (base_y < 0) { - ky_start = (-base_y + args.d1 - 1)/args.d1; - } - int32_t ky_end = args.KH; - const int32_t y_max = args.IH - 1 - base_y; - if (y_max < 0) { - ky_end = ky_start; - } else if (base_y + (args.KH - 1)*args.d1 >= args.IH) { - ky_end = min(ky_end, y_max/args.d1 + 1); - } + float acc = 0.0f; - int32_t kx_start = 0; - if (base_x < 0) { - kx_start = (-base_x + args.d0 - 1)/args.d0; - } - int32_t kx_end = args.KW; - const int32_t x_max = args.IW - 1 - base_x; - if (x_max < 0) { - kx_end = kx_start; - } else if (base_x + (args.KW - 1)*args.d0 >= args.IW) { - kx_end = min(kx_end, x_max/args.d0 + 1); - } + const int32_t base_x = ow*args.s0 - args.p0; - if (ky_start < ky_end && kx_start < kx_end) { - // Depthwise: kernel indexed by channel c (via ne[3] dimension), no IC loop - const uint64_t w_base = (uint64_t) c * args.nb02; - const uint64_t src_base = (uint64_t) n * args.nb13 + (uint64_t) c * args.nb12; - - for (int32_t ky = ky_start; ky < ky_end; ++ky) { - const int32_t iy = base_y + ky*args.d1; - const uint64_t src_row = src_base + (uint64_t) iy * args.nb11; - const uint64_t w_row = w_base + (uint64_t) ky * args.nb01; - - for (int32_t kx = kx_start; kx < kx_end; ++kx) { - const int32_t ix = base_x + kx*args.d0; - const float x = *(device const float *)(src + src_row + (uint64_t) ix * args.nb10); - const float w = (float)(*(device const TK *)(weights + w_row + (uint64_t) kx * args.nb00)); - acc += x * w; - } + int32_t kx_start = 0; + if (base_x < 0) { + kx_start = (-base_x + args.d0 - 1)/args.d0; + } + int32_t kx_end = args.KW; + const int32_t x_max = args.IW - 1 - base_x; + if (x_max < 0) { + kx_end = kx_start; + } else if (base_x + (args.KW - 1)*args.d0 >= args.IW) { + kx_end = min(kx_end, x_max/args.d0 + 1); + } + + if (ky_start < ky_end && kx_start < kx_end) { + for (int32_t ky = ky_start; ky < ky_end; ++ky) { + const int32_t iy = base_y + ky*args.d1; + const uint64_t src_row = src_base + (uint64_t) iy * args.nb11; + const uint64_t w_row = w_base + (uint64_t) ky * args.nb01; + + for (int32_t kx = kx_start; kx < kx_end; ++kx) { + const int32_t ix = base_x + kx*args.d0; + const float x = *(device const float *)(src + src_row + (uint64_t) ix * args.nb10); + const float w = (float)(*(device const TK *)(weights + w_row + (uint64_t) kx * args.nb00)); + acc += x * w; } } + } - const uint64_t dst_offs = - (uint64_t) n * args.nb3 + - (uint64_t) c * args.nb2 + - (uint64_t) oh * args.nb1 + - (uint64_t) ow * args.nb0; + const uint64_t dst_offs = + (uint64_t) n * args.nb3 + + (uint64_t) c * args.nb2 + + (uint64_t) oh * args.nb1 + + (uint64_t) ow * args.nb0; - *(device float *)(dst + dst_offs) = acc; - } + *(device float *)(dst + dst_offs) = acc; } template [[host_name("kernel_conv_2d_dw_f32_f32")]] @@ -4999,7 +4995,6 @@ kernel void kernel_conv_2d_dw( device const char * src, device char * dst, uint3 tgpig[[threadgroup_position_in_grid]], - uint3 tgpg[[threadgroups_per_grid]], uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]); @@ -5010,7 +5005,6 @@ kernel void kernel_conv_2d_dw( device const char * src, device char * dst, uint3 tgpig[[threadgroup_position_in_grid]], - uint3 tgpg[[threadgroups_per_grid]], uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]); From 309d67caba6d2403f54d9e083115affa837a7135 Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Mon, 29 Jun 2026 00:11:17 +0900 Subject: [PATCH 4/9] metal : add channel-tiled CONV_2D_DW kernel for non-contiguous layouts --- ggml/src/ggml-metal/ggml-metal-device.cpp | 21 +++++ ggml/src/ggml-metal/ggml-metal-device.h | 1 + ggml/src/ggml-metal/ggml-metal-ops.cpp | 65 ++++++++++---- ggml/src/ggml-metal/ggml-metal.metal | 102 ++++++++++++++++++++++ 4 files changed, 171 insertions(+), 18 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 0868d6078e65..3ee0fc2901c4 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1890,6 +1890,27 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw(ggml_ return res; } +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw_tiled(ggml_metal_library_t lib, const ggml_tensor * op) { + assert(op->op == GGML_OP_CONV_2D_DW); + + GGML_ASSERT(op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32); + GGML_ASSERT(op->src[1]->type == GGML_TYPE_F32); + GGML_ASSERT(op->type == GGML_TYPE_F32); + + char base[256]; + char name[256]; + + snprintf(base, 256, "kernel_conv_2d_dw_tiled_%s_%s", ggml_type_name(op->src[0]->type), ggml_type_name(op->src[1]->type)); + snprintf(name, 256, "%s", base); + + ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); + if (!res.pipeline) { + res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); + } + + return res; +} + ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d(ggml_metal_library_t lib, const ggml_tensor * op) { assert(op->op == GGML_OP_CONV_3D); diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index c2484f82f9b5..5e02496b925f 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -153,6 +153,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_tran struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw (ggml_metal_library_t lib, const struct ggml_tensor * op); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw_tiled (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_upscale (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pad (ggml_metal_library_t lib, const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 91720bea648a..d5c568aa3c60 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3800,29 +3800,58 @@ int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { /*.d1 =*/ d1, }; - auto pipeline = ggml_metal_library_get_pipeline_conv_2d_dw(lib, op); + // Use channel-tiled kernel when input has non-contiguous spatial layout (e.g. CWHN) + const bool use_tiled = (nb12 < nb10); - int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); - nth = std::min(nth, 256); - nth = std::max(nth, 1); + if (use_tiled) { + auto pipeline = ggml_metal_library_get_pipeline_conv_2d_dw_tiled(lib, op); - // 3D grid: x = OW (tiled by nth), y = OH, z = C * N - const int32_t OW = ne0; - const int32_t OH = ne1; - const int32_t C = ne12; - const int32_t N = ne13; + int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); + nth = std::min(nth, 256); + nth = std::max(nth, 1); - const int tg_x = (OW + nth - 1) / nth; - const int tg_y = OH; - const int tg_z = C * N; + const int32_t OW = ne0; + const int32_t OH = ne1; + const int32_t C = ne12; + const int32_t N = ne13; - ggml_metal_encoder_set_pipeline(enc, pipeline); - ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); + // Grid: x = C (tiled by nth), y = OH, z = OW * N + const int tg_x = (C + nth - 1) / nth; + const int tg_y = OH; + const int tg_z = OW * N; + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); + + ggml_metal_encoder_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1); + } else { + auto pipeline = ggml_metal_library_get_pipeline_conv_2d_dw(lib, op); + + int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); + nth = std::min(nth, 256); + nth = std::max(nth, 1); - ggml_metal_encoder_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1); + const int32_t OW = ne0; + const int32_t OH = ne1; + const int32_t C = ne12; + const int32_t N = ne13; + + // Grid: x = OW (tiled by nth), y = OH, z = C * N + const int tg_x = (OW + nth - 1) / nth; + const int tg_y = OH; + const int tg_z = C * N; + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); + + ggml_metal_encoder_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1); + } return 1; } diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 7651eb02478f..075aa8f6eb4e 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -4908,6 +4908,88 @@ kernel void kernel_conv_2d( uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]); +// Channel-tiled kernel for non-contiguous spatial layouts (e.g. CWHN). +// Adjacent threads process adjacent channels (which are contiguous in memory) +// instead of adjacent OW positions, giving coalesced reads. +// Grid: x = C tile (nth channels), y = OH, z = OW * N +template +kernel void kernel_conv_2d_dw_tiled( + constant ggml_metal_kargs_conv_2d_dw & args, + device const char * weights, + device const char * src, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + uint3 tpitg[[thread_position_in_threadgroup]], + uint3 ntg[[threads_per_threadgroup]]) { + + const int32_t c = (int32_t)(tgpig.x * ntg.x + tpitg.x); + if (c >= args.C) { + return; + } + + const int32_t oh = tgpig.y; + const int32_t own = tgpig.z; + const int32_t ow = own % args.OW; + const int32_t n = own / args.OW; + + const int32_t base_y = oh*args.s1 - args.p1; + + int32_t ky_start = 0; + if (base_y < 0) { + ky_start = (-base_y + args.d1 - 1)/args.d1; + } + int32_t ky_end = args.KH; + const int32_t y_max = args.IH - 1 - base_y; + if (y_max < 0) { + ky_end = ky_start; + } else if (base_y + (args.KH - 1)*args.d1 >= args.IH) { + ky_end = min(ky_end, y_max/args.d1 + 1); + } + + const int32_t base_x = ow*args.s0 - args.p0; + + int32_t kx_start = 0; + if (base_x < 0) { + kx_start = (-base_x + args.d0 - 1)/args.d0; + } + int32_t kx_end = args.KW; + const int32_t x_max = args.IW - 1 - base_x; + if (x_max < 0) { + kx_end = kx_start; + } else if (base_x + (args.KW - 1)*args.d0 >= args.IW) { + kx_end = min(kx_end, x_max/args.d0 + 1); + } + + float acc = 0.0f; + + if (ky_start < ky_end && kx_start < kx_end) { + const uint64_t w_base = (uint64_t) c * args.nb02; + const uint64_t src_base = (uint64_t) n * args.nb13 + (uint64_t) c * args.nb12; + + for (int32_t ky = ky_start; ky < ky_end; ++ky) { + const int32_t iy = base_y + ky*args.d1; + const uint64_t src_row = src_base + (uint64_t) iy * args.nb11; + const uint64_t w_row = w_base + (uint64_t) ky * args.nb01; + + for (int32_t kx = kx_start; kx < kx_end; ++kx) { + const int32_t ix = base_x + kx*args.d0; + const float x = *(device const float *)(src + src_row + (uint64_t) ix * args.nb10); + const float w = (float)(*(device const TK *)(weights + w_row + (uint64_t) kx * args.nb00)); + acc += x * w; + } + } + } + + const uint64_t dst_offs = + (uint64_t) n * args.nb3 + + (uint64_t) c * args.nb2 + + (uint64_t) oh * args.nb1 + + (uint64_t) ow * args.nb0; + + *(device float *)(dst + dst_offs) = acc; +} + +// Direct kernel for contiguous (WHCN) layouts — no shared memory needed. template kernel void kernel_conv_2d_dw( constant ggml_metal_kargs_conv_2d_dw & args, @@ -5008,6 +5090,26 @@ kernel void kernel_conv_2d_dw( uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]); +template [[host_name("kernel_conv_2d_dw_tiled_f32_f32")]] +kernel void kernel_conv_2d_dw_tiled( + constant ggml_metal_kargs_conv_2d_dw & args, + device const char * weights, + device const char * src, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + uint3 tpitg[[thread_position_in_threadgroup]], + uint3 ntg[[threads_per_threadgroup]]); + +template [[host_name("kernel_conv_2d_dw_tiled_f16_f32")]] +kernel void kernel_conv_2d_dw_tiled( + constant ggml_metal_kargs_conv_2d_dw & args, + device const char * weights, + device const char * src, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + uint3 tpitg[[thread_position_in_threadgroup]], + uint3 ntg[[threads_per_threadgroup]]); + typedef void (conv_transpose_1d_t)( constant ggml_metal_kargs_conv_transpose_1d & args, device const float * src0, From 12dbfda057846300a0472c10c990a5dcf433025e Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Mon, 29 Jun 2026 00:22:31 +0900 Subject: [PATCH 5/9] metal : simplify CONV_2D_DW dispatch and trim comments --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 70 ++++++++------------------ ggml/src/ggml-metal/ggml-metal.metal | 10 +--- 2 files changed, 23 insertions(+), 57 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index d5c568aa3c60..5a1872872383 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3770,12 +3770,10 @@ int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { const int32_t d0 = ((const int32_t *) op->op_params)[4]; const int32_t d1 = ((const int32_t *) op->op_params)[5]; - // Kernel shape: [KW, KH, 1, C] — nb03 is the per-channel stride - // Since ne02 == 1, nb03 == nb02, but we use nb03 explicitly for clarity ggml_metal_kargs_conv_2d_dw args = { /*.nb00 =*/ nb00, /*.nb01 =*/ nb01, - /*.nb02 =*/ nb03, // per-channel kernel stride (dim 3, since dim 2 == 1) + /*.nb02 =*/ nb03, /*.nb10 =*/ nb10, /*.nb11 =*/ nb11, /*.nb12 =*/ nb12, @@ -3788,7 +3786,7 @@ int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { /*.IH =*/ ne11, /*.KW =*/ ne00, /*.KH =*/ ne01, - /*.C =*/ ne12, // input channels == output channels for depthwise + /*.C =*/ ne12, /*.OW =*/ ne0, /*.OH =*/ ne1, /*.N =*/ ne13, @@ -3800,58 +3798,32 @@ int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { /*.d1 =*/ d1, }; - // Use channel-tiled kernel when input has non-contiguous spatial layout (e.g. CWHN) const bool use_tiled = (nb12 < nb10); - if (use_tiled) { - auto pipeline = ggml_metal_library_get_pipeline_conv_2d_dw_tiled(lib, op); + auto pipeline = use_tiled + ? ggml_metal_library_get_pipeline_conv_2d_dw_tiled(lib, op) + : ggml_metal_library_get_pipeline_conv_2d_dw(lib, op); - int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); - nth = std::min(nth, 256); - nth = std::max(nth, 1); - - const int32_t OW = ne0; - const int32_t OH = ne1; - const int32_t C = ne12; - const int32_t N = ne13; - - // Grid: x = C (tiled by nth), y = OH, z = OW * N - const int tg_x = (C + nth - 1) / nth; - const int tg_y = OH; - const int tg_z = OW * N; - - ggml_metal_encoder_set_pipeline(enc, pipeline); - ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); - - ggml_metal_encoder_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1); - } else { - auto pipeline = ggml_metal_library_get_pipeline_conv_2d_dw(lib, op); - - int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); - nth = std::min(nth, 256); - nth = std::max(nth, 1); + int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); + nth = std::min(nth, 256); + nth = std::max(nth, 1); - const int32_t OW = ne0; - const int32_t OH = ne1; - const int32_t C = ne12; - const int32_t N = ne13; + const int32_t OW = ne0; + const int32_t OH = ne1; + const int32_t C = ne12; + const int32_t N = ne13; - // Grid: x = OW (tiled by nth), y = OH, z = C * N - const int tg_x = (OW + nth - 1) / nth; - const int tg_y = OH; - const int tg_z = C * N; + const int tg_x = use_tiled ? (C + nth - 1) / nth : (OW + nth - 1) / nth; + const int tg_y = OH; + const int tg_z = use_tiled ? OW * N : C * N; - ggml_metal_encoder_set_pipeline(enc, pipeline); - ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); - ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); - ggml_metal_encoder_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1); - } + ggml_metal_encoder_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1); return 1; } diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 075aa8f6eb4e..6b6f9fd870c2 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -4908,10 +4908,7 @@ kernel void kernel_conv_2d( uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]); -// Channel-tiled kernel for non-contiguous spatial layouts (e.g. CWHN). -// Adjacent threads process adjacent channels (which are contiguous in memory) -// instead of adjacent OW positions, giving coalesced reads. -// Grid: x = C tile (nth channels), y = OH, z = OW * N +// grid: x = C tile, y = OH, z = OW * N (for channel-contiguous layouts) template kernel void kernel_conv_2d_dw_tiled( constant ggml_metal_kargs_conv_2d_dw & args, @@ -4989,7 +4986,7 @@ kernel void kernel_conv_2d_dw_tiled( *(device float *)(dst + dst_offs) = acc; } -// Direct kernel for contiguous (WHCN) layouts — no shared memory needed. +// grid: x = OW tile, y = OH, z = C * N (for spatially-contiguous layouts) template kernel void kernel_conv_2d_dw( constant ggml_metal_kargs_conv_2d_dw & args, @@ -5000,13 +4997,11 @@ kernel void kernel_conv_2d_dw( uint3 tpitg[[thread_position_in_threadgroup]], uint3 ntg[[threads_per_threadgroup]]) { - // 3D grid: x = OW tile, y = OH, z = C * N const int32_t oh = tgpig.y; const int32_t cn = tgpig.z; const int32_t c = cn % args.C; const int32_t n = cn / args.C; - // Precompute vertical bounds (shared by all threads in this row) const int32_t base_y = oh*args.s1 - args.p1; int32_t ky_start = 0; @@ -5024,7 +5019,6 @@ kernel void kernel_conv_2d_dw( const uint64_t w_base = (uint64_t) c * args.nb02; const uint64_t src_base = (uint64_t) n * args.nb13 + (uint64_t) c * args.nb12; - // Each thread handles one OW position const int32_t ow = (int32_t)(tgpig.x * ntg.x + tpitg.x); if (ow >= args.OW) { return; From eb6456ac9e71602e32965caf8cebeaf609568ba0 Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Mon, 29 Jun 2026 00:27:40 +0900 Subject: [PATCH 6/9] metal : merge duplicate CONV_2D_DW pipeline getters --- ggml/src/ggml-metal/ggml-metal-device.cpp | 27 ++++------------------- ggml/src/ggml-metal/ggml-metal-device.h | 3 +-- ggml/src/ggml-metal/ggml-metal-ops.cpp | 4 +--- 3 files changed, 6 insertions(+), 28 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 3ee0fc2901c4..15290c3d1091 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1869,7 +1869,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d(ggml_met return res; } -ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw(ggml_metal_library_t lib, const ggml_tensor * op) { +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw(ggml_metal_library_t lib, const ggml_tensor * op, bool tiled) { assert(op->op == GGML_OP_CONV_2D_DW); GGML_ASSERT(op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32); @@ -1879,28 +1879,9 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw(ggml_ char base[256]; char name[256]; - snprintf(base, 256, "kernel_conv_2d_dw_%s_%s", ggml_type_name(op->src[0]->type), ggml_type_name(op->src[1]->type)); - snprintf(name, 256, "%s", base); - - ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); - if (!res.pipeline) { - res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); - } - - return res; -} - -ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw_tiled(ggml_metal_library_t lib, const ggml_tensor * op) { - assert(op->op == GGML_OP_CONV_2D_DW); - - GGML_ASSERT(op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32); - GGML_ASSERT(op->src[1]->type == GGML_TYPE_F32); - GGML_ASSERT(op->type == GGML_TYPE_F32); - - char base[256]; - char name[256]; - - snprintf(base, 256, "kernel_conv_2d_dw_tiled_%s_%s", ggml_type_name(op->src[0]->type), ggml_type_name(op->src[1]->type)); + snprintf(base, 256, "kernel_conv_2d_dw%s_%s_%s", + tiled ? "_tiled" : "", + ggml_type_name(op->src[0]->type), ggml_type_name(op->src[1]->type)); snprintf(name, 256, "%s", base); ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index 5e02496b925f..9d4aca121595 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -152,8 +152,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_tran struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); -struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw (ggml_metal_library_t lib, const struct ggml_tensor * op); -struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw_tiled (ggml_metal_library_t lib, const struct ggml_tensor * op); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw (ggml_metal_library_t lib, const struct ggml_tensor * op, bool tiled); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_upscale (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pad (ggml_metal_library_t lib, const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 5a1872872383..45909c4777b5 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3800,9 +3800,7 @@ int ggml_metal_op_conv_2d_dw(ggml_metal_op_t ctx, int idx) { const bool use_tiled = (nb12 < nb10); - auto pipeline = use_tiled - ? ggml_metal_library_get_pipeline_conv_2d_dw_tiled(lib, op) - : ggml_metal_library_get_pipeline_conv_2d_dw(lib, op); + auto pipeline = ggml_metal_library_get_pipeline_conv_2d_dw(lib, op, use_tiled); int nth = ggml_metal_pipeline_max_theads_per_threadgroup(pipeline); nth = std::min(nth, 256); From 8ce0cb4138652b4f20a978a6783c515de7bed723 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Wed, 8 Jul 2026 11:54:52 +0300 Subject: [PATCH 7/9] tests : add F16 CONV2D_DW tests --- tests/test-backend-ops.cpp | 43 +++++++++++++++++++++++++------------- 1 file changed, 29 insertions(+), 14 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 6f9d3ce21043..44c0fc3b3832 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5451,25 +5451,28 @@ struct test_conv_2d : public test_case { struct test_conv_2d_dw : public test_case { const std::array ne_input; const std::array ne_kernel; + const ggml_type type_kernel; const int stride; const int padding; const int dilation; const bool cwhn; std::string vars() override { - return VARS_TO_STR6(ne_input, ne_kernel, stride, padding, dilation, cwhn); + return VARS_TO_STR7(ne_input, ne_kernel, type_kernel, stride, padding, dilation, cwhn); } - test_conv_2d_dw(std::array ne_input = {64, 64, 16, 1}, + test_conv_2d_dw( + std::array ne_input = {64, 64, 16, 1}, std::array ne_kernel = {3, 3, 1, 16}, + ggml_type type_kernel = GGML_TYPE_F32, int stride = 1, int padding = 0, int dilation = 1, bool cwhn = false) - : ne_input(ne_input), ne_kernel(ne_kernel), stride(stride), padding(padding), dilation(dilation), cwhn(cwhn) {} + : ne_input(ne_input), ne_kernel(ne_kernel), type_kernel(type_kernel), stride(stride), padding(padding), dilation(dilation), cwhn(cwhn) {} ggml_tensor * build_graph(ggml_context * ctx) override { ggml_tensor * input = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne_input.data()); ggml_set_name(input, "input"); - ggml_tensor * kernel = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne_kernel.data()); + ggml_tensor * kernel = ggml_new_tensor(ctx, type_kernel, 4, ne_kernel.data()); ggml_set_name(kernel, "kernel"); if (cwhn) { @@ -8093,6 +8096,13 @@ static std::vector> make_test_cases_eval() { for (auto kernel_type : {GGML_TYPE_F32, GGML_TYPE_F16}) { test_cases.emplace_back(new test_conv_2d( { W, H, Cin, 2 }, { KW, KH, Cin, Cout }, kernel_type, s0, s1, p0, p1, d0, d1, false)); + + //if (Cin == Cout) { + // test_cases.emplace_back(new test_conv_2d_dw( + // { W, H, Cin, 2 }, { KW, KH, Cin, Cout }, s0, p0, d0, false)); + // test_cases.emplace_back(new test_conv_2d_dw( + // { W, H, Cin, 2 }, { KW, KH, Cin, Cout }, s0, p0, d0, true)); + //} } } } @@ -8114,10 +8124,15 @@ static std::vector> make_test_cases_eval() { // test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_F16, {1024, 1024, 256, 1}, {3, 3, 256, 1}, 1, 1, 1, 1, 1, 1, true)); // test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_F32, {1024, 1024, 256, 1}, {3, 3, 256, 1}, 1, 1, 1, 1, 1, 1, true)); - test_cases.emplace_back(new test_conv_2d_dw({17, 34, 9, 1}, {3, 3, 1, 9}, 1, 0, 1, false)); - test_cases.emplace_back(new test_conv_2d_dw({17, 34, 9, 1}, {3, 3, 1, 9}, 1, 0, 1, true)); - test_cases.emplace_back(new test_conv_2d_dw({32, 8, 64, 1}, {3, 3, 1, 64}, 2, 1, 1, false)); - test_cases.emplace_back(new test_conv_2d_dw({32, 8, 64, 1}, {3, 3, 1, 64}, 2, 1, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({17, 34, 9, 1}, {3, 3, 1, 9}, GGML_TYPE_F32, 1, 0, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({17, 34, 9, 1}, {3, 3, 1, 9}, GGML_TYPE_F32, 1, 0, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({32, 8, 64, 1}, {3, 3, 1, 64}, GGML_TYPE_F32, 2, 1, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({32, 8, 64, 1}, {3, 3, 1, 64}, GGML_TYPE_F32, 2, 1, 1, true)); + + test_cases.emplace_back(new test_conv_2d_dw({17, 34, 9, 1}, {3, 3, 1, 9}, GGML_TYPE_F16, 1, 0, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({17, 34, 9, 1}, {3, 3, 1, 9}, GGML_TYPE_F16, 1, 0, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({32, 8, 64, 1}, {3, 3, 1, 64}, GGML_TYPE_F16, 2, 1, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({32, 8, 64, 1}, {3, 3, 1, 64}, GGML_TYPE_F16, 2, 1, 1, true)); // CONV_3D auto calc_conv_output_size_3d = [](int64_t ins, int64_t ks, int s, int p, int d) -> int64_t { @@ -9621,12 +9636,12 @@ static std::vector> make_test_cases_perf() { } } - test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, 1, 1, 1, false)); - test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, 1, 1, 1, true)); - test_cases.emplace_back(new test_conv_2d_dw({112, 112, 32, 1}, {3, 3, 1, 32}, 1, 1, 1, false)); - test_cases.emplace_back(new test_conv_2d_dw({112, 112, 32, 1}, {3, 3, 1, 32}, 1, 1, 1, true)); - test_cases.emplace_back(new test_conv_2d_dw({56, 56, 128, 1}, {5, 5, 1, 128}, 2, 2, 1, false)); - test_cases.emplace_back(new test_conv_2d_dw({56, 56, 128, 1}, {5, 5, 1, 128}, 2, 2, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, GGML_TYPE_F32, 1, 1, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, GGML_TYPE_F32, 1, 1, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({112, 112, 32, 1}, {3, 3, 1, 32}, GGML_TYPE_F32, 1, 1, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({112, 112, 32, 1}, {3, 3, 1, 32}, GGML_TYPE_F32, 1, 1, 1, true)); + test_cases.emplace_back(new test_conv_2d_dw({56, 56, 128, 1}, {5, 5, 1, 128}, GGML_TYPE_F32, 2, 2, 1, false)); + test_cases.emplace_back(new test_conv_2d_dw({56, 56, 128, 1}, {5, 5, 1, 128}, GGML_TYPE_F32, 2, 2, 1, true)); for (ggml_type kernel_type : {GGML_TYPE_F32, GGML_TYPE_F16}) { test_cases.emplace_back(new test_conv_transpose_2d({256, 256, 256, 1}, {3, 3, 16, 256}, 1, kernel_type)); From ce51dfc32cbb82c57c425d3deaf58b7d63cb5530 Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Wed, 8 Jul 2026 23:02:16 +0900 Subject: [PATCH 8/9] cpu : fix F16 kernel support for CONV_2D_DW --- ggml/src/ggml-cpu/ops.cpp | 33 ++++++++++++++++++++++----------- 1 file changed, 22 insertions(+), 11 deletions(-) diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index fde939b4ad2a..e1f47e5ab094 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -7299,6 +7299,13 @@ struct ggml_conv_2d_dw_params { int dilation_y; }; +static inline float ggml_conv_2d_dw_knl_f32(const char * data, int64_t i, ggml_type type) { + if (type == GGML_TYPE_F16) { + return GGML_FP16_TO_FP32(((const ggml_fp16_t *)data)[i]); + } + return ((const float *)data)[i]; +} + static void ggml_compute_forward_conv_2d_dw_cwhn( const ggml_compute_params * params, const ggml_tensor * src, @@ -7307,7 +7314,8 @@ static void ggml_compute_forward_conv_2d_dw_cwhn( const ggml_conv_2d_dw_params & p) { const int64_t c = p.channels; - const float * knl_data = (const float *)kernel->data; + const char * knl_data = (const char *)kernel->data; + const ggml_type knl_type = kernel->type; const int64_t rows_total = p.dst_h * p.batch; const int64_t rows_per_thread = (rows_total + params->nth - 1) / params->nth; @@ -7315,13 +7323,15 @@ static void ggml_compute_forward_conv_2d_dw_cwhn( const int64_t row_end = MIN(row_start + rows_per_thread, rows_total); #ifdef GGML_SIMD + int64_t c_pkg_end = 0; + if (knl_type == GGML_TYPE_F32) { #if defined(__ARM_FEATURE_SVE) const int64_t pkg_size = svcntw(); #else const int64_t pkg_size = GGML_F32_EPR; #endif - const int64_t pkg_count = c / pkg_size; - const int64_t c_pkg_end = pkg_count * pkg_size; + c_pkg_end = (c / pkg_size) * pkg_size; + } #else const int64_t c_pkg_end = 0; #endif @@ -7335,8 +7345,7 @@ static void ggml_compute_forward_conv_2d_dw_cwhn( const int64_t src_x_base = dst_x * p.stride_x - p.pad_x; #ifdef GGML_SIMD - // Vectorized loop - for (int64_t c_i = 0; c_i < c_pkg_end; c_i += pkg_size) { + for (int64_t c_i = 0; c_i < c_pkg_end; c_i += GGML_F32_EPR) { GGML_F32_VEC sum = GGML_F32_VEC_ZERO; for (int64_t knl_y = 0; knl_y < p.knl_h; ++knl_y) { const int64_t src_y = src_y_base + knl_y * p.dilation_y; @@ -7348,7 +7357,8 @@ static void ggml_compute_forward_conv_2d_dw_cwhn( if (src_x < 0 || src_x >= p.src_w) { continue; } - GGML_F32_VEC k = GGML_F32_VEC_LOAD(knl_data + (knl_y * p.knl_w + knl_x) * c + c_i); + const float * kp = (const float *)knl_data + (knl_y * p.knl_w + knl_x) * c + c_i; + GGML_F32_VEC k = GGML_F32_VEC_LOAD(kp); GGML_F32_VEC s = GGML_F32_VEC_LOAD(src_data + (src_y * p.src_w + src_x) * c + c_i); sum = GGML_F32_VEC_FMA(sum, k, s); } @@ -7356,7 +7366,6 @@ static void ggml_compute_forward_conv_2d_dw_cwhn( GGML_F32_VEC_STORE(dst_data + c_i, sum); } #endif - // Scalar loop for (int64_t c_i = c_pkg_end; c_i < c; ++c_i) { float sum = 0.0f; for (int64_t knl_y = 0; knl_y < p.knl_h; ++knl_y) { @@ -7369,7 +7378,7 @@ static void ggml_compute_forward_conv_2d_dw_cwhn( if (src_x < 0 || src_x >= p.src_w) { continue; } - sum += knl_data[(knl_y * p.knl_w + knl_x) * c + c_i] + sum += ggml_conv_2d_dw_knl_f32(knl_data, (knl_y * p.knl_w + knl_x) * c + c_i, knl_type) * src_data[(src_y * p.src_w + src_x) * c + c_i]; } } @@ -7390,9 +7399,11 @@ static void ggml_compute_forward_conv_2d_dw_whcn( const int64_t per_thread = (n + params->nth - 1) / params->nth; const int64_t start = params->ith * per_thread; const int64_t end = MIN(start + per_thread, n); + const char * knl_base = (const char *)kernel->data; + const ggml_type knl_type = kernel->type; for (int64_t i = start; i < end; ++i) { - const float * knl_data = (const float *)kernel->data + (i % p.channels) * p.knl_w * p.knl_h; + const int64_t knl_offset = (i % p.channels) * p.knl_w * p.knl_h; const float * src_data = (const float *)src->data + i * p.src_w * p.src_h; float * dst_data = (float *)dst->data + i * p.dst_w * p.dst_h; @@ -7410,7 +7421,7 @@ static void ggml_compute_forward_conv_2d_dw_whcn( if (src_x < 0 || src_x >= p.src_w) { continue; } - sum += knl_data[knl_y * p.knl_w + knl_x] + sum += ggml_conv_2d_dw_knl_f32(knl_base, knl_offset + knl_y * p.knl_w + knl_x, knl_type) * src_data[src_y * p.src_w + src_x]; } } @@ -7442,13 +7453,13 @@ void ggml_compute_forward_conv_2d_dw( p.dilation_x = dst->op_params[4]; p.dilation_y = dst->op_params[5]; + GGML_ASSERT(kernel->type == GGML_TYPE_F32 || kernel->type == GGML_TYPE_F16); GGML_ASSERT(kernel->ne[3] == p.channels); GGML_ASSERT(dst->ne[3] == p.batch); if (ggml_is_contiguous(src)) { ggml_compute_forward_conv_2d_dw_whcn(params, src, kernel, dst, p); } else if (ggml_is_contiguous_channels(src)) { - // kernel should also have channels most contiguous in memory GGML_ASSERT(kernel->nb[0] >= kernel->nb[2] && kernel->nb[1] >= kernel->nb[0]); ggml_compute_forward_conv_2d_dw_cwhn(params, src, kernel, dst, p); } else { From c8e6cb6eac8f009efe1a9ff9e37987ce376a67d1 Mon Sep 17 00:00:00 2001 From: Souleyman Boudouh Date: Thu, 9 Jul 2026 11:41:50 +0900 Subject: [PATCH 9/9] tests : remove commented-out CONV_2D_DW test block --- tests/test-backend-ops.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 44c0fc3b3832..09dc4a3e6d18 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -8096,13 +8096,6 @@ static std::vector> make_test_cases_eval() { for (auto kernel_type : {GGML_TYPE_F32, GGML_TYPE_F16}) { test_cases.emplace_back(new test_conv_2d( { W, H, Cin, 2 }, { KW, KH, Cin, Cout }, kernel_type, s0, s1, p0, p1, d0, d1, false)); - - //if (Cin == Cout) { - // test_cases.emplace_back(new test_conv_2d_dw( - // { W, H, Cin, 2 }, { KW, KH, Cin, Cout }, s0, p0, d0, false)); - // test_cases.emplace_back(new test_conv_2d_dw( - // { W, H, Cin, 2 }, { KW, KH, Cin, Cout }, s0, p0, d0, true)); - //} } } }