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,
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;
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
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;
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);
}
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) {
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];
}
}
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;
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];
}
}
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 {
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) {
assert(op->op == GGML_OP_CONV_3D);
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, 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);
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:
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;
{
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);
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) {
ggml_tensor * op = ctx->node(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);
uint3 tpitg[[thread_position_in_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)(
constant ggml_metal_kargs_conv_transpose_1d & args,
device const float * src0,
struct test_conv_2d_dw : public test_case {
const std::array<int64_t, 4> ne_input;
const std::array<int64_t, 4> 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<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},
+ 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) {
// 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 {
}
}
- 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({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));
overview / pkg / ggml / sources / llama.cpp / commitdiff