Lumpiasty/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

metal : add CONV_2D_DW (depthwise convolution) support (#21565)

Browse Source 
* metal : add CONV_2D_DW (depthwise 2D convolution) support

* test : add perf cases for CONV_2D_DW

* metal : use 3D dispatch for CONV_2D_DW kernel

* metal : add channel-tiled CONV_2D_DW kernel for non-contiguous layouts

* metal : simplify CONV_2D_DW dispatch and trim comments

* metal : merge duplicate CONV_2D_DW pipeline getters

* tests : add F16 CONV2D_DW tests

* cpu : fix F16 kernel support for CONV_2D_DW

* tests : remove commented-out CONV_2D_DW test block

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
This commit is contained in:
Sou-ly
2026-07-09 12:29:15 +03:00
committed by GitHub
co-authored by Georgi Gerganov
parent ccb0c34223
commit 92b187c97e
9 changed files with 381 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cpu/ops.cpp
+22 -11
View File
@@ -7299,6 +7299,13 @@ struct ggml_conv_2d_dw_params {
int dilation_y; 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( static void ggml_compute_forward_conv_2d_dw_cwhn(
const ggml_compute_params * params, const ggml_compute_params * params,
const ggml_tensor * src, 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 ggml_conv_2d_dw_params & p) {
const int64_t c = p.channels; 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_total = p.dst_h * p.batch;
const int64_t rows_per_thread = (rows_total + params->nth - 1) / params->nth; 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); const int64_t row_end = MIN(row_start + rows_per_thread, rows_total);
#ifdef GGML_SIMD #ifdef GGML_SIMD
int64_t c_pkg_end = 0;
if (knl_type == GGML_TYPE_F32) {
#if defined(__ARM_FEATURE_SVE) #if defined(__ARM_FEATURE_SVE)
const int64_t pkg_size = svcntw(); const int64_t pkg_size = svcntw();
#else #else
const int64_t pkg_size = GGML_F32_EPR; const int64_t pkg_size = GGML_F32_EPR;
#endif #endif
const int64_t pkg_count = c / pkg_size; c_pkg_end = (c / pkg_size) * pkg_size;
const int64_t c_pkg_end = pkg_count * pkg_size; }
#else #else
const int64_t c_pkg_end = 0; const int64_t c_pkg_end = 0;
#endif #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; const int64_t src_x_base = dst_x * p.stride_x - p.pad_x;
#ifdef GGML_SIMD #ifdef GGML_SIMD
// Vectorized loop for (int64_t c_i = 0; c_i < c_pkg_end; c_i += GGML_F32_EPR) {
for (int64_t c_i = 0; c_i < c_pkg_end; c_i += pkg_size) {
GGML_F32_VEC sum = GGML_F32_VEC_ZERO; GGML_F32_VEC sum = GGML_F32_VEC_ZERO;
for (int64_t knl_y = 0; knl_y < p.knl_h; ++knl_y) { 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; 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) { if (src_x < 0 || src_x >= p.src_w) {
continue; 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); 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); 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); GGML_F32_VEC_STORE(dst_data + c_i, sum);
} }
#endif #endif
// Scalar loop
for (int64_t c_i = c_pkg_end; c_i < c; ++c_i) { for (int64_t c_i = c_pkg_end; c_i < c; ++c_i) {
float sum = 0.0f; float sum = 0.0f;
for (int64_t knl_y = 0; knl_y < p.knl_h; ++knl_y) { 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) { if (src_x < 0 || src_x >= p.src_w) {
continue; 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]; * 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 per_thread = (n + params->nth - 1) / params->nth;
const int64_t start = params->ith * per_thread; const int64_t start = params->ith * per_thread;
const int64_t end = MIN(start + per_thread, n); 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) { 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; 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; 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) { if (src_x < 0 || src_x >= p.src_w) {
continue; 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]; * 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_x = dst->op_params[4];
p.dilation_y = dst->op_params[5]; 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(kernel->ne[3] == p.channels);
GGML_ASSERT(dst->ne[3] == p.batch); GGML_ASSERT(dst->ne[3] == p.batch);
if (ggml_is_contiguous(src)) { if (ggml_is_contiguous(src)) {
ggml_compute_forward_conv_2d_dw_whcn(params, src, kernel, dst, p); ggml_compute_forward_conv_2d_dw_whcn(params, src, kernel, dst, p);
} else if (ggml_is_contiguous_channels(src)) { } 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_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); ggml_compute_forward_conv_2d_dw_cwhn(params, src, kernel, dst, p);
} else { } else {
ggml/src/ggml-metal/ggml-metal-device.cpp
+23
View File
@@ -1869,6 +1869,29 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d(ggml_met
return res; return res;
} }
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);
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_%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);
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) { 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); assert(op->op == GGML_OP_CONV_3D);
ggml/src/ggml-metal/ggml-metal-device.h
+1
View File
@@ -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_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_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 (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_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_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); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pad (ggml_metal_library_t lib, const struct ggml_tensor * op);
ggml/src/ggml-metal/ggml-metal-device.m
+4
View File
@@ -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->src[1]->type == GGML_TYPE_F32 &&
op->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&
(op->src[0]->type == GGML_TYPE_F16 || op->src[0]->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: case GGML_OP_UPSCALE:
return op->src[0]->type == GGML_TYPE_F32; return op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_POOL_1D: case GGML_OP_POOL_1D:
ggml/src/ggml-metal/ggml-metal-impl.h
+28
View File
@@ -656,6 +656,34 @@ typedef struct {
int32_t d1; int32_t d1;
} ggml_metal_kargs_conv_2d; } 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 { typedef struct {
uint64_t ofs0; uint64_t ofs0;
uint64_t ofs1; uint64_t ofs1;
ggml/src/ggml-metal/ggml-metal-ops.cpp
+84
View File
@@ -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); n_fuse = ggml_metal_op_conv_2d(ctx, idx);
} break; } break;
case GGML_OP_CONV_2D_DW:
{
n_fuse = ggml_metal_op_conv_2d_dw(ctx, idx);
} break;
case GGML_OP_CONV_TRANSPOSE_1D: case GGML_OP_CONV_TRANSPOSE_1D:
{ {
n_fuse = ggml_metal_op_conv_transpose_1d(ctx, idx); n_fuse = ggml_metal_op_conv_transpose_1d(ctx, idx);
@@ -3742,6 +3746,86 @@ int ggml_metal_op_conv_2d(ggml_metal_op_t ctx, int idx) {
return 1; 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];
ggml_metal_kargs_conv_2d_dw args = {
/*.nb00 =*/ nb00,
/*.nb01 =*/ nb01,
/*.nb02 =*/ nb03,
/*.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,
/*.OW =*/ ne0,
/*.OH =*/ ne1,
/*.N =*/ ne13,
/*.s0 =*/ s0,
/*.s1 =*/ s1,
/*.p0 =*/ p0,
/*.p1 =*/ p1,
/*.d0 =*/ d0,
/*.d1 =*/ d1,
};
const bool use_tiled = (nb12 < nb10);
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);
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 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_dispatch_threadgroups(enc, tg_x, tg_y, tg_z, nth, 1, 1);
return 1;
}
int ggml_metal_op_conv_3d(ggml_metal_op_t ctx, int idx) { int ggml_metal_op_conv_3d(ggml_metal_op_t ctx, int idx) {
ggml_tensor * op = ctx->node(idx); ggml_tensor * op = ctx->node(idx);
ggml/src/ggml-metal/ggml-metal-ops.h
+1
View File
@@ -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_rope (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_im2col (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 (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_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_1d (ggml_metal_op_t ctx, int idx);
int ggml_metal_op_conv_transpose_2d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_transpose_2d (ggml_metal_op_t ctx, int idx);
ggml/src/ggml-metal/ggml-metal.metal
+196
View File
@@ -4908,6 +4908,202 @@ kernel void kernel_conv_2d<half>(
uint3 tpitg[[thread_position_in_threadgroup]], uint3 tpitg[[thread_position_in_threadgroup]],
uint3 ntg[[threads_per_threadgroup]]); uint3 ntg[[threads_per_threadgroup]]);
// grid: x = C tile, y = OH, z = OW * N (for channel-contiguous layouts)
template <typename TK>
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;
}
// grid: x = OW tile, y = OH, z = C * N (for spatially-contiguous layouts)
template <typename TK>
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 tpitg[[thread_position_in_threadgroup]],
uint3 ntg[[threads_per_threadgroup]]) {
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;
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 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 ow = (int32_t)(tgpig.x * ntg.x + tpitg.x);
if (ow >= args.OW) {
return;
}
float acc = 0.0f;
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);
}
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;
*(device float *)(dst + dst_offs) = acc;
}
template [[host_name("kernel_conv_2d_dw_f32_f32")]]
kernel void kernel_conv_2d_dw<float>(
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_f16_f32")]]
kernel void kernel_conv_2d_dw<half>(
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_f32_f32")]]
kernel void kernel_conv_2d_dw_tiled<float>(
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<half>(
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)( typedef void (conv_transpose_1d_t)(
constant ggml_metal_kargs_conv_transpose_1d & args, constant ggml_metal_kargs_conv_transpose_1d & args,
device const float * src0, device const float * src0,
tests/test-backend-ops.cpp
+22 -10
View File
@@ -5451,25 +5451,28 @@ struct test_conv_2d : public test_case {
struct test_conv_2d_dw : public test_case { struct test_conv_2d_dw : public test_case {
const std::array<int64_t, 4> ne_input; const std::array<int64_t, 4> ne_input;
const std::array<int64_t, 4> ne_kernel; const std::array<int64_t, 4> ne_kernel;
const ggml_type type_kernel;
const int stride; const int stride;
const int padding; const int padding;
const int dilation; const int dilation;
const bool cwhn; const bool cwhn;
std::string vars() override { 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<int64_t, 4> ne_input = {64, 64, 16, 1}, test_conv_2d_dw(
std::array<int64_t, 4> ne_input = {64, 64, 16, 1},
std::array<int64_t, 4> ne_kernel = {3, 3, 1, 16}, std::array<int64_t, 4> 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) 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 * build_graph(ggml_context * ctx) override {
ggml_tensor * input = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne_input.data()); ggml_tensor * input = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne_input.data());
ggml_set_name(input, "input"); 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"); ggml_set_name(kernel, "kernel");
if (cwhn) { if (cwhn) {
@@ -8114,10 +8117,15 @@ static std::vector<std::unique_ptr<test_case>> 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_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_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}, GGML_TYPE_F32, 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({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}, 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, 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({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 // CONV_3D
auto calc_conv_output_size_3d = [](int64_t ins, int64_t ks, int s, int p, int d) -> int64_t { auto calc_conv_output_size_3d = [](int64_t ins, int64_t ks, int s, int p, int d) -> int64_t {
@@ -9621,8 +9629,12 @@ static std::vector<std::unique_ptr<test_case>> 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}, GGML_TYPE_F32, 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({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}) { 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)); test_cases.emplace_back(new test_conv_transpose_2d({256, 256, 256, 1}, {3, 3, 16, 256}, 1, kernel_type));