Commit 669204ee2 for llama.cpp
commit 669204ee2b9004845b34f26466a91d003edc98aa
Author: Mark Todorovich <mtodorov@qti.qualcomm.com>
Date: Fri Nov 6 11:51:57 2026 -0800
hexagon: conv_2d optimizations and support for conv_3d (#31814)
* hexagon: Generalization of CONV_2D implementation for CONV_3D
- New op HTP_OP_CONV_3D is wired through htp-ops.h, main.c (shares op_conv_2d) and the op remap in htp-kparams.cpp.
- conv_3d runs on the existing HMX tiled (implicit GEMM) conv2d path. KD and KH are folded into one virtual kernel row (KD*KH), and output rows are counted as N*OD*OH.
- htp_kparams_solve_conv3d checks the conv_3d layout and calls the tiled solver htp_kparams_conv2d_tiled, which takes an is3d flag.
- The tiled kernel takes the sizes and strides that differ between conv_2d and conv_3d from a struct htp_conv_geom. Only the input stage unfolds depth.
- 256-entry DMA queue can no longer overflow
Assisted-by: Claude Code
* hexagon: modify checking to eliminate divide-by-zero hazards
* hexagon: conv2d tiled: compute the conv_3d unfold divisors on the host
The KH and OH fastdiv values used to unfold conv_3d kh into (kd, kh) and
oh into (od, oh) were built per op on the device, each with a 64-bit
divide. Pass them in the kernel params instead; they fit in the space the
window variant already reserves, so the struct size does not change.
Assisted-by: Claude Code
* hexagon : preserve fold scheduling in conv2d for small size
hexagon : support arbitrary positive conv2d strides
Support positive horizontal strides for WHCN and CWHN inputs. Preserve
load/shuffle paths for strides 1/2 and use gather for larger strides.
hexagon : improve conv2d tile size selection
Prefer the largest feasible output-channel block and derive spatial
chunk capacity from VTCM usage. Compare capacity and thread-aligned
candidates using per-thread work and chunk overhead.
Co-authored-by: Yiwei Shao <yiwei@aizip.ai>
* hex-conv: fix compiler warning
---------
Co-authored-by: Yiwei Shao <yiwei@aizip.ai>
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
diff --git a/ggml/src/ggml-hexagon/htp-kparams.cpp b/ggml/src/ggml-hexagon/htp-kparams.cpp
index 4fa900bed..022532328 100644
--- a/ggml/src/ggml-hexagon/htp-kparams.cpp
+++ b/ggml/src/ggml-hexagon/htp-kparams.cpp
@@ -62,6 +62,7 @@ htp_op_code htp_op_remap(const struct ggml_tensor * t) {
case GGML_OP_IM2COL_3D: return HTP_OP_IM2COL_3D;
case GGML_OP_CONV_2D: return HTP_OP_CONV_2D;
case GGML_OP_CONV_2D_DW: return HTP_OP_CONV_2D_DW;
+ case GGML_OP_CONV_3D: return HTP_OP_CONV_3D;
case GGML_OP_ROLL: return HTP_OP_ROLL;
case GGML_OP_UPSCALE: return HTP_OP_UPSCALE;
case GGML_OP_LIGHTNING_INDEXER: return HTP_OP_LIGHTNING_INDEXER;
@@ -3944,6 +3945,84 @@ static bool htp_kparams_conv2d_win(
}
}
+ uint32_t n_threads = nth;
+ // Keep the fold schedule, where output traffic can dominate input packing.
+ if (!fold && rt_block < n_rt && ct > 0 && nth > 0) {
+ const uint32_t old_rt = rt_block;
+ const uint32_t ct_max = (std::min)((std::min)((std::max)(4 * halo, (n_tiles + 3) / 4), 512u), n_tiles);
+ const size_t budget = htp_session_vtcm_size(sess);
+ // Approximate chunk overhead in input crouton work units.
+ const size_t c_chunk = 8;
+ const uint32_t tail = n_tiles % ct;
+ const size_t work =
+ (size_t) (n_tiles / ct) * ((ct + halo + nth - 1) / nth) + (tail ? (tail + halo + nth - 1) / nth : 0);
+ size_t best_cost =
+ (size_t) ((n_rt + rt_block - 1) / rt_block) * (n_cc * work + c_chunk * (n_tiles / ct + (tail != 0)));
+ bool found = false;
+
+ // Keep output DMA segments. Try larger kernel blocks to avoid rebuilding input croutons.
+ for (uint32_t rb = n_rt; rb > old_rt && !found; rb = (rb + 1) / 2) {
+ for (uint32_t nt = nth; nt > 0; nt /= 2) {
+ for (uint32_t sb = 8; sb > 0; sb /= 2) {
+ const size_t input_stage = hex_round_up(cwhn ? (size_t) sb * 32 * Cin * sizeof(float) + 256 : (size_t) n_ch * sb * 32 * sizeof(float), 128);
+ const size_t scratch = hex_round_up(
+ (std::max)(2 * input_stage + (size_t) HTP_CONV2D_WIN_OBUF * 32 * seg_t * 32 * sizeof(float),
+ wscr),
+ tile);
+ if (n_cc == 0 || n_wt > budget / tile / n_cc / rb || halo > budget / tile / n_cc ||
+ scratch > budget / nt) {
+ continue;
+ }
+ const size_t weights = (size_t) rb * n_wt * n_cc * tile;
+ htp_conv2d_vtcm_layout fixed;
+ htp_conv2d_vtcm_layout_build(&fixed, rb, weights, (size_t) n_cc * halo * tile, 0, scratch, nt);
+ if (fixed.total_bytes >= budget) {
+ continue;
+ }
+ const size_t per_ct = 2 * ((size_t) n_cc + rb) * tile;
+ const uint32_t cap = (uint32_t) (std::min)((size_t) ct_max, (budget - fixed.total_bytes) / per_ct);
+ const uint32_t aligned = (cap + halo) / nt * nt;
+ const uint32_t candidates[2] = { cap, aligned > halo ? aligned - halo : 0 };
+ for (uint32_t i = 0; i < 2; i++) {
+ const uint32_t candidate = candidates[i];
+ if (candidate == 0 || (i > 0 && candidate == cap)) {
+ continue;
+ }
+ if (hex_round_up((size_t) T * sizeof(uint32_t), 128) + 2 * hex_round_up((size_t) candidate, 128) > HTP_DDR_SPAD_SIZE) {
+ continue;
+ }
+ const uint32_t rem = n_tiles % candidate;
+ const size_t input_work = (size_t) (n_tiles / candidate) * ((candidate + halo + nt - 1) / nt) + (rem ? (rem + halo + nt - 1) / nt : 0);
+ const size_t chunks = n_tiles / candidate + (rem != 0);
+ const size_t cost = (size_t) ((n_rt + rb - 1) / rb) * (n_cc * input_work + c_chunk * chunks);
+ if (cost < best_cost) {
+ const size_t input = (size_t) n_cc * (candidate + halo) * tile;
+ const size_t output = (size_t) rb * candidate * tile;
+ htp_conv2d_vtcm_layout L;
+ htp_conv2d_vtcm_layout_build(&L, rb, weights, input, output, scratch, nt);
+ if (L.total_bytes > budget) {
+ continue;
+ }
+ found = true;
+ best_cost = cost;
+ rt_block = rb;
+ ct = candidate;
+ seg_b = sb;
+ n_threads = nt;
+ stage = input_stage;
+ thr = scratch;
+ act = input;
+ out = output;
+ w_size = weights;
+ total = L.total_bytes;
+ n_bc = candidate + halo;
+ }
+ }
+ }
+ }
+ }
+ }
+
if (hex_round_up((size_t) T * sizeof(uint32_t), 128) + 2 * hex_round_up((size_t) ct, 128) > HTP_DDR_SPAD_SIZE) {
return false;
}
@@ -3954,7 +4033,7 @@ static bool htp_kparams_conv2d_win(
const uint32_t n_cs = (32 + CS - 1) / CS;
kp->kernel_type = (uint8_t) HTP_CONV2D_KERNEL_HMX_WINDOW;
- kp->n_threads = (uint8_t) nth;
+ kp->n_threads = (uint8_t) n_threads;
kp->cwhn = cwhn ? 1 : 0;
kp->w_kmajor = !w_rows ? 1 : 0;
kp->n_row_tiles = n_rt;
@@ -3982,6 +4061,50 @@ static bool htp_kparams_conv2d_win(
return true;
}
+struct htp_conv_norm_geom {
+ uint32_t KW, KH, Cin, Cout;
+ uint32_t OW, OH, N;
+ int32_t s0, d0;
+ size_t es;
+ bool is3d;
+};
+
+static bool htp_conv_normalize_geom(const struct ggml_tensor * dst, struct htp_conv_norm_geom * g) {
+ const struct ggml_tensor * src0 = dst->src[0];
+ const bool is3d = dst->op == GGML_OP_CONV_3D;
+ const int32_t s0 = ggml_get_op_params_i32(dst, 0);
+ const int32_t d0 = ggml_get_op_params_i32(dst, is3d ? 6 : 4);
+
+ const int64_t kh_w = is3d ? src0->ne[2] * src0->ne[1] : src0->ne[1];
+ const int64_t cin_w = is3d ? ggml_get_op_params_i32(dst, 9) : src0->ne[2];
+ const int64_t oh_w = is3d ? dst->ne[2] * dst->ne[1] : dst->ne[1];
+ const int64_t n_w = is3d ? ggml_get_op_params_i32(dst, 10) : dst->ne[3];
+ const int64_t k_w = src0->ne[0] * kh_w * cin_w;
+ const int64_t ct_w = n_w * oh_w * ((dst->ne[0] + 31) / 32);
+
+ if (kh_w < 1 || kh_w > 255 || k_w < 1 || k_w > (1 << 24) || ct_w < 1 || ct_w > UINT32_MAX) {
+ return false;
+ }
+ if (s0 <= 0 || d0 <= 0 || dst->nb[0] != sizeof(float)) {
+ return false;
+ }
+
+ *g = {
+ /* .KW = */ (uint32_t) src0->ne[0],
+ /* .KH = */ (uint32_t) kh_w,
+ /* .Cin = */ (uint32_t) cin_w,
+ /* .Cout = */ (uint32_t) (is3d ? ggml_get_op_params_i32(dst, 11) : src0->ne[3]),
+ /* .OW = */ (uint32_t) dst->ne[0],
+ /* .OH = */ (uint32_t) oh_w,
+ /* .N = */ (uint32_t) n_w,
+ /* .s0 = */ s0,
+ /* .d0 = */ d0,
+ /* .es = */ ggml_type_size(src0->type),
+ /* .is3d = */ is3d,
+ };
+ return true;
+}
+
// Tiled implicit GEMM, MMA on HMX if available, else HVX
static bool htp_kparams_conv2d_tiled(
const struct ggml_hexagon_session * sess,
@@ -3992,92 +4115,129 @@ static bool htp_kparams_conv2d_tiled(
return false;
}
- const struct ggml_tensor * src0 = dst->src[0];
- const struct ggml_tensor * src1 = dst->src[1];
- const int32_t s0 = ggml_get_op_params_i32(dst, 0);
- const int32_t d0 = ggml_get_op_params_i32(dst, 4);
-
- const uint32_t KW = src0->ne[0];
- const uint32_t KH = src0->ne[1];
- const uint32_t Cin = src0->ne[2];
- const uint32_t Cout = src0->ne[3];
- const uint32_t OW = dst->ne[0];
- const uint32_t OH = dst->ne[1];
- const uint32_t N = dst->ne[3];
- const size_t es = ggml_type_size(src0->type);
-
- if (s0 != 1 || dst->nb[0] != sizeof(float)) {
+ struct htp_conv_norm_geom g;
+ if (!htp_conv_normalize_geom(dst, &g)) {
return false;
}
+ const struct ggml_tensor * src0 = dst->src[0];
+ const struct ggml_tensor * src1 = dst->src[1];
+
// kernel k order must match the staged input: (cin, kh, kw) for WHCN input, (kh, kw, cin) for CWHN input.
// Kernel is either k-contiguous rows per cout (WHCN only) or cout-contiguous with a uniform k stride.
- const bool whcn = src1->nb[0] == sizeof(float);
- const bool cwhn = !whcn && src1->nb[2] == sizeof(float) && src1->nb[0] == Cin * sizeof(float);
- const bool co_c = Cout == 1 || src0->nb[3] == es;
- const bool w_rows = whcn && src0->nb[0] == es && (KH == 1 || src0->nb[1] == KW * es) && (Cin == 1 || src0->nb[2] == KW * KH * es);
+ const bool whcn = g.is3d || src1->nb[0] == sizeof(float);
+ const bool cwhn = !whcn && src1->nb[2] == sizeof(float) && src1->nb[0] == g.Cin * sizeof(float);
+ const bool co_c = g.Cout == 1 || src0->nb[3] == g.es;
+ const bool w_rows = g.is3d || (whcn && src0->nb[0] == g.es && (g.KH == 1 || src0->nb[1] == g.KW * g.es) && (g.Cin == 1 || src0->nb[2] == g.KW * g.KH * g.es));
const size_t ks = whcn ? src0->nb[0] : src0->nb[2];
bool kmajor = false;
if (whcn) {
- kmajor = co_c && ks >= Cout * es && (KH == 1 || src0->nb[1] == KW * ks) && (Cin == 1 || src0->nb[2] == KW * KH * ks);
+ kmajor = co_c && ks >= g.Cout * g.es && (g.KH == 1 || src0->nb[1] == g.KW * ks) && (g.Cin == 1 || src0->nb[2] == g.KW * g.KH * ks);
} else if (cwhn) {
- kmajor = co_c && ks >= Cout * es && src0->nb[0] == Cin * ks && (KH == 1 || src0->nb[1] == KW * Cin * ks);
+ kmajor = co_c && ks >= g.Cout * g.es && src0->nb[0] == g.Cin * ks && (g.KH == 1 || src0->nb[1] == g.KW * g.Cin * ks);
}
if (!w_rows && !kmajor) {
return false;
}
kmajor = kmajor && !w_rows;
- const uint32_t K = KW * KH * Cin;
- const uint32_t kg = !cwhn ? K : d0 == 1 ? KW * Cin : Cin;
+ const uint32_t K = g.KW * g.KH * g.Cin;
+ const uint32_t kg = !cwhn ? K : g.d0 == 1 ? g.KW * g.Cin : g.Cin;
const uint32_t n_grp = K / kg;
const uint32_t kg_p = hex_round_up(kg, 32);
const uint32_t k_pad = hex_round_up(n_grp * kg_p, 64);
const uint32_t n_kt = k_pad / 32;
- const uint32_t n_rt = (Cout + 31) / 32;
- const uint32_t n_owt = (OW + 31) / 32;
- const uint32_t n_ct = N * OH * n_owt;
+ const uint32_t n_rt = (g.Cout + 31) / 32;
+ const uint32_t n_owt = (g.OW + 31) / 32;
+ const uint32_t n_ct = g.N * g.OH * n_owt;
const uint32_t wk = (std::min)(k_pad, 512u);
- const uint32_t nth = htp_session_n_threads(sess);
+ const uint32_t max_threads = htp_session_n_threads(sess);
const size_t tile = 2048;
uint32_t rt_block = n_rt;
- uint32_t ct = (std::min)(n_ct, 64u);
- uint32_t seg = (std::min)(8u, n_owt);
+ uint32_t ct = 0;
+ uint32_t seg = 0;
+ uint32_t nth = 0;
uint32_t n_stage = 1;
uint32_t sr = 0;
- size_t stage = 0, ostage = 0, prow = 0, wpack = 0;
- size_t w_size = 0, act_size = 0, out_size = 0, thr_size = 0, total = 0;
+ size_t stage = 0, thr_size = 0;
+ size_t best_cost = SIZE_MAX;
+ const size_t wpack = kmajor ? 2 * (size_t) wk * 64 * sizeof(float) + 96 * 128 : 2 * (size_t) 32 * wk * (sizeof(float) + sizeof(uint16_t));
+ const size_t budget = htp_session_vtcm_size(sess);
+ // Approximate chunk overhead in patch K-tile work units.
+ const size_t c_chunk = 8;
+
+ // Keep the kernel resident to avoid regenerating the input for each rt block.
+ // Balance patch work per thread and chunk overhead before reducing rt.
+ for (uint32_t rb = n_rt; rb > 0 && ct == 0; rb = rb > 1 ? (rb + 1) / 2 : 0) {
+ for (uint32_t nt = max_threads; nt > 0; nt /= 2) {
+ for (uint32_t sg = (std::min)(8u, n_owt); sg > 0; --sg) {
+ // Keep the unused final stride lane for full-vector loads during packing.
+ const uint64_t wb = (uint64_t) 32 * sg * g.s0 + (uint64_t) (g.KW - 1) * g.d0;
+ if (wb > budget / sizeof(float) || wb > INT32_MAX) {
+ continue;
+ }
+ const uint64_t width = cwhn ? wb * g.Cin : wb;
+ if (width > budget / sizeof(float) || width > UINT32_MAX - 63) {
+ continue;
+ }
+ const uint32_t row = hex_round_up((uint32_t) width, cwhn ? 64 : 32);
+ const uint64_t rows = cwhn ? g.KH : (uint64_t) g.Cin * g.KH;
+ if (rows > budget / ((uint64_t) row * sizeof(float))) {
+ continue;
+ }
+ const size_t input_stage = cwhn ? (size_t) g.KH * row * sizeof(float) : (size_t) g.Cin * g.KH * row * sizeof(float);
+ const size_t output_stage = 2 * 32 * (size_t) sg * 32 * sizeof(float);
+ const size_t transpose = cwhn ? input_stage / 2 + 256 + 64 * 128 : 0;
+ const size_t scratch = hex_round_up((std::max)(input_stage + output_stage + transpose, wpack), tile);
+ htp_conv2d_vtcm_layout fixed;
+ htp_conv2d_vtcm_layout_build(&fixed, rb, (size_t) rb * n_kt * tile, 0, 0, scratch, nt);
+ if (fixed.total_bytes >= budget) {
+ continue;
+ }
- while (true) {
- const uint32_t wb = 32 * seg + (KW - 1) * d0;
- sr = cwhn ? hex_round_up(wb * Cin, 64) : hex_round_up(wb, 32);
- stage = cwhn ? (size_t) KH * sr * sizeof(float) : (size_t) Cin * KH * sr * sizeof(float);
- ostage = 2 * 32 * (size_t) seg * 32 * sizeof(float);
- prow = cwhn ? stage / 2 + 256 + 64 * 128 : 0;
- wpack = kmajor ? 2 * (size_t) wk * 64 * sizeof(float) + 96 * 128 : 2 * (size_t) 32 * wk * (sizeof(float) + sizeof(uint16_t));
- thr_size = hex_round_up((std::max)(stage + ostage + prow, wpack), tile);
- w_size = (size_t) rt_block * n_kt * tile;
- act_size = (size_t) ct * n_kt * tile;
- out_size = (size_t) rt_block * ct * tile;
- struct htp_conv2d_vtcm_layout L;
- htp_conv2d_vtcm_layout_build(&L, rt_block, w_size, act_size, out_size, thr_size, nth);
- total = L.total_bytes;
- if (total <= htp_session_vtcm_size(sess)) {
- break;
- }
- if (ct > 8) {
- ct /= 2;
- } else if (seg > 1) {
- seg--;
- } else if (rt_block > 1) {
- rt_block = (rt_block + 1) / 2;
- } else if (ct > 1) {
- ct /= 2;
- } else {
- return false;
+ // Each spatial tile adds two activation tiles and two output tiles per rt.
+ const size_t per_ct = 2 * ((size_t) n_kt + rb) * tile;
+ const uint32_t nc = (uint32_t) (std::min)((size_t) (std::min)(n_ct, 64u), (budget - fixed.total_bytes) / per_ct);
+ const uint32_t candidates[2] = { nc, nc / nt * nt };
+ for (uint32_t i = 0; i < 2; i++) {
+ const uint32_t candidate = candidates[i];
+ if (candidate == 0 || (i > 0 && candidate == candidates[0])) {
+ continue;
+ }
+ const uint32_t per_thread = (candidate + nt - 1) / nt;
+ if (sg > per_thread) {
+ continue;
+ }
+ const size_t work = (size_t) (n_ct / candidate) * per_thread + (n_ct % candidate + nt - 1) / nt;
+ const size_t chunks = n_ct / candidate + (n_ct % candidate != 0);
+ const size_t cost = work * n_kt + c_chunk * chunks;
+ if (cost < best_cost) {
+ best_cost = cost;
+ rt_block = rb;
+ ct = candidate;
+ seg = sg;
+ nth = nt;
+ sr = row;
+ stage = input_stage;
+ thr_size = scratch;
+ }
+ }
+ }
}
}
+ if (ct == 0) {
+ return false;
+ }
+
+ const size_t w_size = (size_t) rt_block * n_kt * tile;
+ const size_t act_size = (size_t) ct * n_kt * tile;
+ const size_t out_size = (size_t) rt_block * ct * tile;
+ const size_t ostage = 2 * 32 * (size_t) seg * 32 * sizeof(float);
+ const size_t prow = cwhn ? stage / 2 + 256 + 64 * 128 : 0;
+ struct htp_conv2d_vtcm_layout L;
+ htp_conv2d_vtcm_layout_build(&L, rt_block, w_size, act_size, out_size, thr_size, nth);
+ size_t total = L.total_bytes;
// Second stage buffer: DMA of the next segment overlaps HVX pack of this one.
// Skip large stages. They already fill VTCM, and the extra DMA slows HMX.
@@ -4085,7 +4245,7 @@ static bool htp_kparams_conv2d_tiled(
const size_t thr2 = hex_round_up((std::max)(2 * stage + ostage + prow, wpack), tile);
struct htp_conv2d_vtcm_layout L2;
htp_conv2d_vtcm_layout_build(&L2, rt_block, w_size, act_size, out_size, thr2, nth);
- if (stage <= 128 * 1024 && L2.total_bytes <= htp_session_vtcm_size(sess)) {
+ if (stage <= 128 * 1024 && 2 * g.KH <= 255 && L2.total_bytes <= htp_session_vtcm_size(sess)) {
n_stage = 2;
thr_size = thr2;
total = L2.total_bytes;
@@ -4114,7 +4274,9 @@ static bool htp_kparams_conv2d_tiled(
kp->tiled.n_stage = n_stage;
kp->tiled.wk_chunk = wk;
kp->tiled.div_ow_tiles = init_fastdiv_values(n_owt);
- kp->tiled.div_oh = init_fastdiv_values(OH);
+ kp->tiled.div_oh = init_fastdiv_values(g.OH);
+ kp->tiled.div_khr = init_fastdiv_values(src0->ne[1]);
+ kp->tiled.div_ohr = init_fastdiv_values(dst->ne[1]);
return true;
}
@@ -4163,6 +4325,10 @@ static bool htp_kparams_solve_conv2d(
const int32_t d0 = ggml_get_op_params_i32(op, 4);
const int32_t d1 = ggml_get_op_params_i32(op, 5);
+ if (s0 <= 0 || s1 <= 0 || d0 <= 0 || d1 <= 0) {
+ return false;
+ }
+
const int64_t ow = (src1->ne[0] + 2 * p0 - d0 * (src0->ne[0] - 1) - 1) / s0 + 1;
const int64_t oh = (src1->ne[1] + 2 * p1 - d1 * (src0->ne[1] - 1) - 1) / s1 + 1;
if (ow != dst->ne[0] || oh != dst->ne[1] || src0->ne[3] != dst->ne[2] || src1->ne[3] != dst->ne[3]) {
@@ -4337,6 +4503,60 @@ static bool htp_kparams_solve_conv2d_dw(
return true;
}
+static bool htp_kparams_solve_conv3d(
+ const struct ggml_hexagon_session * sess,
+ const struct ggml_tensor * op,
+ struct htp_conv2d_kernel_params * kparams
+) {
+ memset(kparams, 0, sizeof(*kparams));
+
+ const struct ggml_tensor * src0 = op->src[0];
+ const struct ggml_tensor * src1 = op->src[1];
+ const struct ggml_tensor * dst = op;
+
+ if ((src0->type != GGML_TYPE_F16 && src0->type != GGML_TYPE_F32) ||
+ src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
+ return false;
+ }
+
+ const int32_t IC = ggml_get_op_params_i32(op, 9);
+ const int32_t N = ggml_get_op_params_i32(op, 10);
+ const int32_t OC = ggml_get_op_params_i32(op, 11);
+ if (IC <= 0 || N <= 0 || OC <= 0 || src0->ne[3] != (int64_t) IC * OC || src1->ne[3] != (int64_t) IC * N ||
+ dst->ne[3] != (int64_t) OC * N) {
+ return false;
+ }
+
+ for (int i = 0; i < 3; i++) {
+ const int32_t s = ggml_get_op_params_i32(op, i);
+ const int32_t p = ggml_get_op_params_i32(op, 3 + i);
+ const int32_t d = ggml_get_op_params_i32(op, 6 + i);
+ if (s <= 0 || d <= 0 || p < 0 || src0->ne[i] <= 0) {
+ return false;
+ }
+ const int64_t o = (src1->ne[i] + 2 * p - d * (src0->ne[i] - 1) - 1) / s + 1;
+ if (o != dst->ne[i]) {
+ return false;
+ }
+ }
+
+ const size_t es = ggml_type_size(src0->type);
+ const bool w_rows = src0->nb[0] == es && src0->nb[1] == src0->ne[0] * es && src0->nb[2] == src0->ne[1] * src0->nb[1] &&
+ ((int64_t) IC * OC == 1 || src0->nb[3] == src0->ne[2] * src0->nb[2]);
+ if (!w_rows || src1->nb[0] != sizeof(float) || dst->nb[2] != dst->ne[1] * dst->nb[1]) {
+ return false;
+ }
+
+ const int64_t total_rows = dst->ne[1] * dst->ne[2] * dst->ne[3];
+ if (total_rows > UINT32_MAX) {
+ return false;
+ }
+ kparams->total_rows = (uint32_t) total_rows;
+ kparams->n_threads = total_rows == 0 ? 1 : (std::min)((uint32_t) htp_session_n_threads(sess), (uint32_t) total_rows);
+
+ return htp_kparams_conv2d_tiled(sess, op, kparams);
+}
+
static bool htp_dsv4_hc_pre_can_collapse(
const struct ggml_tensor * x,
const struct ggml_tensor * weights,
@@ -5241,6 +5461,9 @@ bool htp_kparams_solve(
case HTP_OP_CONV_2D_DW:
return htp_kparams_invoke<htp_kparams_solve_conv2d_dw>(sess, op, kparams);
+ case HTP_OP_CONV_3D:
+ return htp_kparams_invoke<htp_kparams_solve_conv3d>(sess, op, kparams);
+
case HTP_OP_SUM_ROWS:
case HTP_OP_SUM:
case HTP_OP_ARGMAX:
diff --git a/ggml/src/ggml-hexagon/htp/conv2d-ops.c b/ggml/src/ggml-hexagon/htp/conv2d-ops.c
index 42e70e089..6f35cb19b 100644
--- a/ggml/src/ggml-hexagon/htp/conv2d-ops.c
+++ b/ggml/src/ggml-hexagon/htp/conv2d-ops.c
@@ -87,11 +87,11 @@ static inline HVX_Vector conv2d_hmx_rt_restore(const __fp16 * scales, uint32_t r
return Q6_V_vsplat_R((127 + e) << 23);
}
-// Output rows [*row_base, *row_base + *n_rows) of N * OH for this device.
+// Output rows [*row_base, *row_base + *n_rows) of total for this device.
// The output is written by DMA only, so device ranges need no cache line alignment.
-static void conv2d_mdev_rows(const struct htp_ops_context * octx, uint32_t * row_base, uint32_t * n_rows) {
+static void conv2d_mdev_rows(const struct htp_ops_context * octx, uint32_t total, uint32_t * row_base, uint32_t * n_rows) {
*row_base = 0;
- *n_rows = octx->dst->ne[3] * octx->dst->ne[1];
+ *n_rows = total;
if (octx->ctx->mdev.count > 1) {
const struct htp_tensor_mdev_range range = htp_tensor_mdev_partition(
*n_rows, 1, octx->ctx->mdev.idx, octx->ctx->mdev.count, &octx->ctx->mdev.count_div);
@@ -134,9 +134,51 @@ static void conv2d_hmx_worker(void * data) {
}
}
+// Tiled path values whose source differs between conv_2d and conv_3d.
+// conv_3d runs as conv_2d with virtual KH = KD * KHr and OH = OD * OHr, only the input stage unfolds them.
+struct htp_conv_geom {
+ uint32_t KH, KHr, Cin, Cout;
+ uint32_t OH, OHr, N, ID;
+ int32_t s0, s1, s2, p0, p1, p2, d0, d1, d2;
+ size_t x_nb_d, x_nb_c, x_nb_n; // input strides for d, cin, n
+ size_t y_nb_c, y_nb_n; // dst strides for cout, n
+ size_t w_nb_co; // kernel cout stride
+};
+
+// conv_3d tensors are [W, H, D, C * N] with the channel as the inner index of ne[3]
+static void conv_geom_init(struct htp_conv_geom * g, const struct htp_ops_context * octx) {
+ const struct htp_tensor * kernel = octx->src[0];
+ const struct htp_tensor * input = octx->src[1];
+ const struct htp_tensor * dst = octx->dst;
+ const int32_t * op = octx->op_params;
+
+ if (octx->op == HTP_OP_CONV_3D) {
+ const uint32_t IC = op[9];
+ const uint32_t OC = op[11];
+ *g = (struct htp_conv_geom) {
+ .KH = kernel->ne[2] * kernel->ne[1], .KHr = kernel->ne[1], .Cin = IC, .Cout = OC,
+ .OH = dst->ne[2] * dst->ne[1], .OHr = dst->ne[1], .N = op[10], .ID = input->ne[2],
+ .s0 = op[0], .s1 = op[1], .s2 = op[2], .p0 = op[3], .p1 = op[4], .p2 = op[5], .d0 = op[6], .d1 = op[7], .d2 = op[8],
+ .x_nb_d = input->nb[2], .x_nb_c = input->nb[3], .x_nb_n = (size_t) IC * input->nb[3],
+ .y_nb_c = dst->nb[3], .y_nb_n = (size_t) OC * dst->nb[3],
+ .w_nb_co = (size_t) IC * kernel->nb[3],
+ };
+ } else {
+ *g = (struct htp_conv_geom) {
+ .KH = kernel->ne[1], .KHr = kernel->ne[1], .Cin = kernel->ne[2], .Cout = kernel->ne[3],
+ .OH = dst->ne[1], .OHr = dst->ne[1], .N = dst->ne[3], .ID = 1,
+ .s0 = op[0], .s1 = op[1], .s2 = 1, .p0 = op[2], .p1 = op[3], .p2 = 0, .d0 = op[4], .d1 = op[5], .d2 = 1,
+ .x_nb_d = 0, .x_nb_c = input->nb[2], .x_nb_n = input->nb[3],
+ .y_nb_c = dst->nb[2], .y_nb_n = dst->nb[3],
+ .w_nb_co = kernel->nb[3],
+ };
+ }
+}
+
struct htp_conv2d_hmx_context {
struct htp_ops_context * octx;
const struct htp_conv2d_kernel_params * kp;
+ const struct htp_conv_geom * g;
struct htp_conv2d_vtcm_layout layout;
__fp16 * scales; // [rt] HMX scales
__fp16 * w_tiles;
@@ -169,7 +211,7 @@ static void conv2d_hmx_wpack_thread(unsigned int nth, unsigned int ith, void * d
const struct htp_conv2d_kernel_params * kp = hc->kp;
const struct htp_tensor * kernel = octx->src[0];
- const uint32_t Cout = kernel->ne[3];
+ const uint32_t Cout = hc->g->Cout;
const uint32_t K = kp->tiled.k_size;
const uint32_t k_pad = kp->tiled.n_k_tiles * HMX_FP16_TILE_N_COLS;
const uint32_t WK = kp->tiled.wk_chunk;
@@ -200,8 +242,8 @@ static void conv2d_hmx_wpack_thread(unsigned int nth, unsigned int ith, void * d
for (uint32_t i = 0; i < 2 && i * WK < k_pad; i++) {
const uint32_t kc = i * WK;
const uint32_t kv = MIN(WK, K - kc);
- dma_queue_push(dmaq, dma_make_data(wdma[i], kernel->data + (size_t) co0 * kernel->nb[3] + (size_t) kc * es),
- WK * es, kernel->nb[3], kv * es, nco);
+ dma_queue_push(dmaq, dma_make_data(wdma[i], kernel->data + (size_t) co0 * hc->g->w_nb_co + (size_t) kc * es),
+ WK * es, hc->g->w_nb_co, kv * es, nco);
}
for (uint32_t kc = 0; kc < k_pad; kc += WK) {
@@ -234,42 +276,50 @@ static void conv2d_hmx_wpack_thread(unsigned int nth, unsigned int ith, void * d
const uint32_t next_kc = kc + 2 * WK;
if (next_kc < k_pad) {
const uint32_t next_kv = MIN(WK, K - next_kc);
- dma_queue_push(dmaq, dma_make_data(wdma[bi], kernel->data + (size_t) co0 * kernel->nb[3] + (size_t) next_kc * es),
- WK * es, kernel->nb[3], next_kv * es, nco);
+ dma_queue_push(dmaq, dma_make_data(wdma[bi], kernel->data + (size_t) co0 * hc->g->w_nb_co + (size_t) next_kc * es),
+ WK * es, hc->g->w_nb_co, next_kv * es, nco);
}
bi ^= 1;
}
}
}
-// Stage Cin*KH input rows for G col tiles: stage[(cin*KH + kh)*SR + j] = x[n][cin][ih(kh)][ow0 - p0 + j], zero outside
+// Stage Cin*KH input rows for G col tiles: stage[(cin*KH + kh)*SR + j] = x[n][cin][ih(kh)][ow0*s0 - p0 + j], zero outside
+// For conv_3d, kh = kd * KHr + kh' and oh = od * OHr + oh' are unfolded here
static void conv2d_hmx_stage(const struct htp_ops_context * octx, const struct htp_conv2d_kernel_params * kp,
- dma_queue * dmaq, float * stage, uint32_t n, uint32_t oh, uint32_t ow0, uint32_t G) {
+ const struct htp_conv_geom * g, dma_queue * dmaq, float * stage, uint32_t n, uint32_t oh,
+ uint32_t ow0, uint32_t G) {
const struct htp_tensor * kernel = octx->src[0];
const struct htp_tensor * input = octx->src[1];
const uint32_t KW = kernel->ne[0];
- const uint32_t KH = kernel->ne[1];
- const uint32_t Cin = kernel->ne[2];
+ const uint32_t KH = g->KH;
+ const uint32_t Cin = g->Cin;
const int32_t IW = input->ne[0];
const int32_t IH = input->ne[1];
- const int32_t s1 = octx->op_params[1];
- const int32_t p0 = octx->op_params[2];
- const int32_t p1 = octx->op_params[3];
- const int32_t d0 = octx->op_params[4];
- const int32_t d1 = octx->op_params[5];
+ const int32_t s0 = g->s0;
+ const int32_t s1 = g->s1;
+ const int32_t p0 = g->p0;
+ const int32_t p1 = g->p1;
+ const int32_t d0 = g->d0;
+ const int32_t d1 = g->d1;
const uint32_t SR = kp->tiled.stage_row;
- const int32_t Wb = 32 * G + (KW - 1) * d0;
- const int32_t iw0 = (int32_t) ow0 - p0;
+ const int32_t Wb = 32 * G * s0 + (KW - 1) * d0;
+ const int32_t iw0 = (int32_t) ow0 * s0 - p0;
const int32_t lo = MAX(0, -iw0);
const int32_t hi = MIN(Wb, IW - iw0);
+ const uint32_t od = fastdiv(oh, &kp->tiled.div_ohr);
+ oh -= od * g->OHr;
+
for (uint32_t kh = 0; kh < KH; kh++) {
- const int32_t ih = (int32_t) oh * s1 + (int32_t) kh * d1 - p1;
+ const int32_t kd = fastdiv(kh, &kp->tiled.div_khr);
+ const int32_t id = (int32_t) od * g->s2 + kd * g->d2 - g->p2;
+ const int32_t ih = (int32_t) oh * s1 + (int32_t) (kh - kd * g->KHr) * d1 - p1;
float * rows = stage + (size_t) kh * SR;
- if (ih < 0 || ih >= IH || hi <= lo) {
+ if (id < 0 || id >= (int32_t) g->ID || ih < 0 || ih >= IH || hi <= lo) {
for (uint32_t c = 0; c < Cin; c++) {
hvx_splat_f32_u(rows + (size_t) c * KH * SR, 0.0f, Wb);
}
@@ -277,8 +327,8 @@ static void conv2d_hmx_stage(const struct htp_ops_context * octx, const struct h
continue;
}
- const dma_addr_t src = input->data + (size_t) n * input->nb[3] + (size_t) ih * input->nb[1] + (size_t) (iw0 + lo) * sizeof(float);
- dma_queue_push(dmaq, dma_make_data(rows + lo, src), (size_t) KH * SR * sizeof(float), input->nb[2],
+ const dma_addr_t src = input->data + (size_t) n * g->x_nb_n + (size_t) id * g->x_nb_d + (size_t) ih * input->nb[1] + (size_t) (iw0 + lo) * sizeof(float);
+ dma_queue_push(dmaq, dma_make_data(rows + lo, src), (size_t) KH * SR * sizeof(float), g->x_nb_c,
(size_t) (hi - lo) * sizeof(float), Cin);
if (lo > 0 || hi < Wb) {
@@ -295,9 +345,24 @@ static void conv2d_hmx_stage(const struct htp_ops_context * octx, const struct h
}
}
+static inline HVX_Vector conv2d_hmx_load_strided_f32(const float * src, uint32_t s0, HVX_Vector offsets, HVX_Vector * scratch) {
+ if (s0 == 1) {
+ return hvx_vmemu(src);
+ }
+ if (s0 == 2) {
+ const HVX_Vector lo = hvx_vmemu(src);
+ const HVX_Vector hi = hvx_vmemu(src + 32);
+ return Q6_V_lo_W(Q6_W_vdeal_VVR(hi, lo, -4));
+ }
+ Q6_vgather_ARMVw(scratch, (size_t) src, 32 * s0 * sizeof(float) - 1, offsets);
+ hvx_gather_sync(scratch);
+ return *scratch;
+}
+
// One weight-side tile column: vector i holds k rows 2i and 2i+1 interleaved over 32 pixels
static inline void conv2d_hmx_pack_tile(__fp16 * restrict dst, const float * restrict src, uint32_t K, uint32_t KW,
- uint32_t d0, uint32_t SR, uint32_t n_vec) {
+ uint32_t s0, uint32_t d0, uint32_t SR, uint32_t n_vec, HVX_Vector offsets,
+ HVX_Vector * scratch) {
HVX_Vector * restrict out = (HVX_Vector *) dst;
const HVX_Vector zero = Q6_V_vzero();
const uint32_t wrap = SR - KW * d0;
@@ -307,13 +372,13 @@ static inline void conv2d_hmx_pack_tile(__fp16 * restrict dst, const float * res
uint32_t i = 0;
for (; i < K / 2; i++) {
- const HVX_Vector a = hvx_vmemu(p);
+ const HVX_Vector a = conv2d_hmx_load_strided_f32(p, s0, offsets, scratch);
p += d0;
if (++kw == KW) {
kw = 0;
p += wrap;
}
- const HVX_Vector b = hvx_vmemu(p);
+ const HVX_Vector b = conv2d_hmx_load_strided_f32(p, s0, offsets, scratch);
p += d0;
if (++kw == KW) {
kw = 0;
@@ -322,14 +387,14 @@ static inline void conv2d_hmx_pack_tile(__fp16 * restrict dst, const float * res
out[i] = conv2d_f32_to_f16_shuff(a, b);
}
if (K & 1) {
- out[i++] = conv2d_f32_to_f16_shuff(hvx_vmemu(p), zero);
+ out[i++] = conv2d_f32_to_f16_shuff(conv2d_hmx_load_strided_f32(p, s0, offsets, scratch), zero);
}
for (; i < n_vec; i++) {
out[i] = zero;
}
}
-// CWHN: stage[kh*SR + j*Cin + cin] = x[n][ih(kh)][ow0 - p0 + j][cin], zero outside
+// CWHN: stage[kh*SR + j*Cin + cin] = x[n][ih(kh)][ow0*s0 - p0 + j][cin], zero outside
static void conv2d_hmx_stage_cwhn(const struct htp_ops_context * octx, const struct htp_conv2d_kernel_params * kp,
dma_queue * dmaq, float * stage, uint32_t n, uint32_t oh, uint32_t ow0, uint32_t G) {
const struct htp_tensor * kernel = octx->src[0];
@@ -340,6 +405,7 @@ static void conv2d_hmx_stage_cwhn(const struct htp_ops_context * octx, const str
const uint32_t Cin = kernel->ne[2];
const int32_t IW = input->ne[0];
const int32_t IH = input->ne[1];
+ const int32_t s0 = octx->op_params[0];
const int32_t s1 = octx->op_params[1];
const int32_t p0 = octx->op_params[2];
const int32_t p1 = octx->op_params[3];
@@ -347,8 +413,8 @@ static void conv2d_hmx_stage_cwhn(const struct htp_ops_context * octx, const str
const int32_t d1 = octx->op_params[5];
const uint32_t SR = kp->tiled.stage_row;
- const int32_t Wb = 32 * G + (KW - 1) * d0;
- const int32_t iw0 = (int32_t) ow0 - p0;
+ const int32_t Wb = 32 * G * s0 + (KW - 1) * d0;
+ const int32_t iw0 = (int32_t) ow0 * s0 - p0;
const int32_t lo = MAX(0, -iw0);
const int32_t hi = MIN(Wb, IW - iw0);
@@ -383,8 +449,8 @@ static void conv2d_hmx_cwhn_f16(__fp16 * restrict dst, const float * restrict sr
// Stage the segment of col tiles starting at t, returns its number of col tiles
static uint32_t conv2d_hmx_issue_seg(const struct htp_ops_context * octx, const struct htp_conv2d_kernel_params * kp,
- dma_queue * dmaq, float * stage, uint32_t t, uint32_t t1) {
- const uint32_t OH = octx->dst->ne[1];
+ const struct htp_conv_geom * g, dma_queue * dmaq, float * stage, uint32_t t, uint32_t t1) {
+ const uint32_t OH = g->OH;
const uint32_t n_owt = kp->tiled.n_ow_tiles;
const uint32_t row = fastdiv(t, &kp->tiled.div_ow_tiles);
const uint32_t owt = t - row * n_owt;
@@ -395,7 +461,7 @@ static uint32_t conv2d_hmx_issue_seg(const struct htp_ops_context * octx, const
if (kp->cwhn) {
conv2d_hmx_stage_cwhn(octx, kp, dmaq, stage, n, oh, owt * 32, G);
} else {
- conv2d_hmx_stage(octx, kp, dmaq, stage, n, oh, owt * 32, G);
+ conv2d_hmx_stage(octx, kp, g, dmaq, stage, n, oh, owt * 32, G);
}
return G;
}
@@ -487,7 +553,7 @@ static inline void conv2d_transpose_strided(const __fp16 * restrict base, uint32
// the second one rotated into the upper half.
static void conv2d_hmx_pack_tile_cwhn(__fp16 * restrict dst, const __fp16 * restrict src, HVX_Vector * restrict vin,
HVX_Vector * restrict scratch, const struct htp_conv2d_kernel_params * kp,
- uint32_t gpr, uint32_t Cin, uint32_t d0) {
+ uint32_t gpr, uint32_t Cin, uint32_t s0, uint32_t d0) {
const uint32_t k_pad = kp->tiled.n_k_tiles * HMX_FP16_TILE_N_COLS;
const uint32_t KG = kp->tiled.k_group;
const uint32_t KGp = kp->tiled.k_group_pad;
@@ -523,15 +589,15 @@ static void conv2d_hmx_pack_tile_cwhn(__fp16 * restrict dst, const __fp16 * rest
}
if (va[0] == 32 && va[1] == 32 && pa[1] == pa[0] + 32) {
- conv2d_transpose_strided(pa[0], Cin, scratch, out + 32 * b, out + 32 * b + 16);
+ conv2d_transpose_strided(pa[0], Cin * s0, scratch, out + 32 * b, out + 32 * b + 16);
continue;
}
const HVX_VectorPred q0 = Q6_Q_vsetq_R(va[0] * sizeof(__fp16));
const HVX_VectorPred q1 = Q6_Q_vsetq_R(va[1] * sizeof(__fp16));
for (uint32_t j = 0; j < 32; j++) {
- const HVX_Vector v0 = Q6_V_vand_QV(q0, hvx_vmemu(pa[0] + j * Cin));
- const HVX_Vector v1 = Q6_V_vand_QV(q1, hvx_vmemu(pa[1] + j * Cin));
+ const HVX_Vector v0 = Q6_V_vand_QV(q0, hvx_vmemu(pa[0] + j * Cin * s0));
+ const HVX_Vector v1 = Q6_V_vand_QV(q1, hvx_vmemu(pa[1] + j * Cin * s0));
vin[j] = Q6_V_vor_VV(v0, Q6_V_vror_VR(v1, 64));
}
@@ -718,9 +784,10 @@ static void conv2d_hmx_pack(const struct htp_conv2d_hmx_context * hc, uint32_t i
const struct htp_tensor * kernel = octx->src[0];
const uint32_t KW = kernel->ne[0];
- const uint32_t KH = kernel->ne[1];
- const uint32_t Cin = kernel->ne[2];
- const uint32_t d0 = octx->op_params[4];
+ const uint32_t KH = hc->g->KH;
+ const uint32_t Cin = hc->g->Cin;
+ const uint32_t s0 = (uint32_t) hc->g->s0;
+ const uint32_t d0 = (uint32_t) hc->g->d0;
const uint32_t n_vec = kp->tiled.n_k_tiles * (HMX_FP16_TILE_N_ELMS / 64);
dma_queue * dmaq = octx->ctx->dma[ith];
@@ -747,10 +814,14 @@ static void conv2d_hmx_pack(const struct htp_conv2d_hmx_context * hc, uint32_t i
return;
}
+ // Reuse output staging after unpack has drained its DMA.
+ HVX_Vector * gather_tmp = (HVX_Vector *) (thr + (size_t) n_stage * kp->stage_size);
+ const HVX_Vector offsets = kp->cwhn || s0 <= 2 ? Q6_V_vzero() : hvx_vec_gather_offsets_w(s0 * sizeof(float));
+
uint32_t ng[2] = { 0, 0 };
uint32_t tq = t0;
for (uint32_t i = 0; i < n_stage && tq < t1; i++) {
- ng[i] = conv2d_hmx_issue_seg(octx, kp, dmaq, stages[i], tq, t1);
+ ng[i] = conv2d_hmx_issue_seg(octx, kp, hc->g, dmaq, stages[i], tq, t1);
tq += ng[i];
}
@@ -765,22 +836,29 @@ static void conv2d_hmx_pack(const struct htp_conv2d_hmx_context * hc, uint32_t i
htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_A_PREP, (uint16_t) t);
if (kp->cwhn) {
- const uint32_t Wb = 32 * G + (KW - 1) * d0;
+ const uint32_t Wb = 32 * G * s0 + (KW - 1) * d0;
conv2d_hmx_cwhn_f16(stage16, stage, KH, kp->tiled.stage_row, Wb * Cin);
for (uint32_t g = 0; g < G; g++) {
__fp16 * tile = act + (size_t) (t - cs + g) * n_vec * 64;
- conv2d_hmx_pack_tile_cwhn(tile, stage16 + (size_t) g * 32 * Cin, vin, scratch, kp, gpr, Cin, d0);
+ conv2d_hmx_pack_tile_cwhn(tile, stage16 + (size_t) g * 32 * Cin * s0, vin, scratch, kp, gpr, Cin, s0, d0);
}
} else {
for (uint32_t g = 0; g < G; g++) {
- __fp16 * tile = act + (size_t) (t - cs + g) * n_vec * 64;
- conv2d_hmx_pack_tile(tile, stage + g * 32, kp->tiled.k_size, KW, d0, kp->tiled.stage_row, n_vec);
+ __fp16 * tile = act + (size_t) (t - cs + g) * n_vec * 64;
+ const float * pixels = stage + (size_t) g * 32 * s0;
+ if (s0 == 1) {
+ conv2d_hmx_pack_tile(tile, pixels, kp->tiled.k_size, KW, 1, d0, kp->tiled.stage_row, n_vec, offsets, gather_tmp);
+ } else if (s0 == 2) {
+ conv2d_hmx_pack_tile(tile, pixels, kp->tiled.k_size, KW, 2, d0, kp->tiled.stage_row, n_vec, offsets, gather_tmp);
+ } else {
+ conv2d_hmx_pack_tile(tile, pixels, kp->tiled.k_size, KW, s0, d0, kp->tiled.stage_row, n_vec, offsets, gather_tmp);
+ }
}
}
htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_A_PREP, (uint16_t) t);
if (tq < t1) {
- ng[bi] = conv2d_hmx_issue_seg(octx, kp, dmaq, stage, tq, t1);
+ ng[bi] = conv2d_hmx_issue_seg(octx, kp, hc->g, dmaq, stage, tq, t1);
tq += ng[bi];
}
t += G;
@@ -794,8 +872,8 @@ static void conv2d_hmx_unpack(const struct htp_conv2d_hmx_context * hc, uint32_t
const struct htp_tensor * dst = octx->dst;
const uint32_t OW = dst->ne[0];
- const uint32_t OH = dst->ne[1];
- const uint32_t Cout = dst->ne[2];
+ const uint32_t OH = hc->g->OH;
+ const uint32_t Cout = hc->g->Cout;
const uint32_t n_owt = kp->tiled.n_ow_tiles;
const uint32_t OS = kp->tiled.seg_tiles * 32;
@@ -850,8 +928,8 @@ static void conv2d_hmx_unpack(const struct htp_conv2d_hmx_context * hc, uint32_t
}
htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_O_PROC, (uint16_t) t);
- const dma_addr_t ddr = dst->data + (size_t) n * dst->nb[3] + (size_t) oh * dst->nb[1] + (size_t) co0 * dst->nb[2] + (size_t) ow0 * sizeof(float);
- dma_queue_push(dmaq, dma_make_data(ddr, ob), dst->nb[2], (size_t) OS * sizeof(float), (size_t) nv * sizeof(float), nco);
+ const dma_addr_t ddr = dst->data + (size_t) n * hc->g->y_nb_n + (size_t) oh * dst->nb[1] + (size_t) co0 * hc->g->y_nb_c + (size_t) ow0 * sizeof(float);
+ dma_queue_push(dmaq, dma_make_data(ddr, ob), hc->g->y_nb_c, (size_t) OS * sizeof(float), (size_t) nv * sizeof(float), nco);
n_pushed++;
b ^= 1;
}
@@ -904,8 +982,11 @@ static void conv2d_hmx_run_phase(struct htp_conv2d_hmx_context * hc, int32_t pac
}
static int conv2d_hmx_tiled(struct htp_ops_context * octx, const struct htp_conv2d_kernel_params * kp) {
+ struct htp_conv_geom g;
+ conv_geom_init(&g, octx);
+
uint32_t row_base, n_rows;
- conv2d_mdev_rows(octx, &row_base, &n_rows);
+ conv2d_mdev_rows(octx, g.N * g.OH, &row_base, &n_rows);
if (n_rows == 0) {
return HTP_STATUS_OK;
}
@@ -923,6 +1004,7 @@ static int conv2d_hmx_tiled(struct htp_ops_context * octx, const struct htp_conv
struct htp_conv2d_hmx_context hc = {
.octx = octx,
.kp = kp,
+ .g = &g,
.layout = layout,
.scales = VTCM_LAYOUT_PTR(__fp16, base, layout.off_scales),
.w_tiles = VTCM_LAYOUT_PTR(__fp16, base, layout.off_w),
@@ -2153,7 +2235,7 @@ static int conv2d_hmx_win(struct htp_ops_context * octx, const struct htp_conv2d
const uint32_t d1 = octx->op_params[5];
uint32_t row_base, n_rows;
- conv2d_mdev_rows(octx, &row_base, &n_rows);
+ conv2d_mdev_rows(octx, dst->ne[3] * dst->ne[1], &row_base, &n_rows);
if (n_rows == 0) {
return HTP_STATUS_OK;
}
diff --git a/ggml/src/ggml-hexagon/htp/conv2d-ops.h b/ggml/src/ggml-hexagon/htp/conv2d-ops.h
index a48bf3e83..aefed269e 100644
--- a/ggml/src/ggml-hexagon/htp/conv2d-ops.h
+++ b/ggml/src/ggml-hexagon/htp/conv2d-ops.h
@@ -35,6 +35,8 @@ struct htp_conv2d_tiled_params {
uint32_t wk_chunk;
struct fastdiv_values div_ow_tiles;
struct fastdiv_values div_oh;
+ struct fastdiv_values div_khr; // kernel ne[1], unfolds conv_3d kh = kd * KHr + kh'
+ struct fastdiv_values div_ohr; // dst ne[1], unfolds conv_3d oh = od * OHr + oh'
};
struct htp_conv2d_window_params {
diff --git a/ggml/src/ggml-hexagon/htp/htp-ops.h b/ggml/src/ggml-hexagon/htp/htp-ops.h
index 96257c241..bd7812486 100644
--- a/ggml/src/ggml-hexagon/htp/htp-ops.h
+++ b/ggml/src/ggml-hexagon/htp/htp-ops.h
@@ -117,6 +117,7 @@ enum htp_op_code {
HTP_OP_UNARY_CEIL,
HTP_OP_UNARY_SGN,
HTP_OP_UPSCALE,
+ HTP_OP_CONV_3D,
HTP_OP_DSV4_HC_PRE,
HTP_OP_DSV4_HC_POST,
HTP_OP_TOP_K_MASK,
diff --git a/ggml/src/ggml-hexagon/htp/main.c b/ggml/src/ggml-hexagon/htp/main.c
index ac8820beb..565673212 100644
--- a/ggml/src/ggml-hexagon/htp/main.c
+++ b/ggml/src/ggml-hexagon/htp/main.c
@@ -1030,6 +1030,7 @@ static int execute_op(struct htp_ops_context * octx) {
return op_im2col_3d(octx);
case HTP_OP_CONV_2D:
+ case HTP_OP_CONV_3D:
return op_conv_2d(octx);
case HTP_OP_CONV_2D_DW:
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 96954a5d7..91b0135c2 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -10907,6 +10907,13 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_conv_3d(1, 1, 2, 2, 2, 1, 1, 1, 0, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
test_cases.emplace_back(new test_conv_3d(1, 1, 2, 2, 2, 1, 1, 0, 1, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
test_cases.emplace_back(new test_conv_3d(1, 1, 2, 2, 2, 1, 0, 1, 1, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
+ // input and output channel strides just over 16 MB, beyond the 24-bit DMA stride field
+ test_cases.emplace_back(new test_conv_3d(1, 2, 65, 256, 256, 2, 1, 1, 1, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
+ // KD * KH = 127, 128, 255, 256 around the HTP DMA queue limits
+ test_cases.emplace_back(new test_conv_3d(1, 2, 127, 1, 40, 3, 127, 1, 1, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
+ test_cases.emplace_back(new test_conv_3d(1, 2, 64, 2, 40, 3, 64, 2, 1, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
+ test_cases.emplace_back(new test_conv_3d(1, 2, 85, 3, 40, 3, 85, 3, 1, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
+ test_cases.emplace_back(new test_conv_3d(1, 2, 16, 16, 40, 3, 16, 16, 1, 1, 1, 1, 0, 0, 0, 1, 1, 1, kernel_type));
}
for(uint32_t Cout : {1, 9}){