diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h
index 85a1ae7ae208..ab8269854750 100644
--- a/ggml/include/ggml.h
+++ b/ggml/include/ggml.h
@@ -430,7 +430,8 @@ extern "C" {
         GGML_TYPE_NVFP4   = 40, // NVFP4 (4 blocks, E4M3 scale)
         GGML_TYPE_Q1_0    = 41,
         GGML_TYPE_Q2_0    = 42,
-        GGML_TYPE_COUNT   = 43,
+        GGML_TYPE_F8_E4M3 = 43,
+        GGML_TYPE_COUNT   = 44,
     };
 
     // [TAG_GGML_PREC]
@@ -601,6 +602,12 @@ extern "C" {
 
         GGML_OP_GLU,
 
+        GGML_OP_MUL_MAT_SEGMENTED,
+        GGML_OP_QKNORM_ROPE,
+        GGML_OP_GROUP_NORM_AFFINE_SILU,
+        GGML_OP_CONV_2D_BIAS,
+        GGML_OP_CONV_2D_UPSCALE,
+
         GGML_OP_COUNT,
     };
 
@@ -1406,6 +1413,13 @@ extern "C" {
             struct ggml_tensor  * a,
             float                 eps);
 
+    GGML_API struct ggml_tensor * ggml_qknorm_rope(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * a,
+            struct ggml_tensor  * weight,
+            struct ggml_tensor  * theta,
+            float                 eps);
+
     // group normalize along ne0*ne1*n_groups
     // used in stable-diffusion
     GGML_API struct ggml_tensor * ggml_group_norm(
@@ -1420,6 +1434,22 @@ extern "C" {
             int                   n_groups,
             float                 eps);
 
+    GGML_API struct ggml_tensor * ggml_group_norm_affine_silu(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * a,
+            struct ggml_tensor  * weight,
+            struct ggml_tensor  * bias,
+            int                   n_groups,
+            float                 eps);
+
+    GGML_API struct ggml_tensor * ggml_group_norm_affine_silu_inplace(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * a,
+            struct ggml_tensor  * weight,
+            struct ggml_tensor  * bias,
+            int                   n_groups,
+            float                 eps);
+
     // l2 normalize along rows
     // used in rwkv v7
     GGML_API struct ggml_tensor * ggml_l2_norm(
@@ -1484,6 +1514,12 @@ extern "C" {
             struct ggml_tensor  * a,
             struct ggml_tensor  * b);
 
+    GGML_API struct ggml_tensor * ggml_mul_mat_segmented(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * a,
+            struct ggml_tensor  * b0,
+            struct ggml_tensor  * b1);
+
     // change the precision of a matrix multiplication
     // set to GGML_PREC_F32 for higher precision (useful for phi-2)
     GGML_DEPRECATED(GGML_API void ggml_mul_mat_set_prec(
@@ -2255,6 +2291,31 @@ extern "C" {
             int                   d0,  // dilation dimension 0
             int                   d1); // dilation dimension 1
 
+    GGML_API struct ggml_tensor * ggml_conv_2d_direct_bias(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * a,
+            struct ggml_tensor  * b,
+            struct ggml_tensor  * bias,
+            int                   s0,
+            int                   s1,
+            int                   p0,
+            int                   p1,
+            int                   d0,
+            int                   d1);
+
+    GGML_API struct ggml_tensor * ggml_conv_2d_direct_upscale(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * a,
+            struct ggml_tensor  * b,
+            struct ggml_tensor  * bias,
+            int                   upscale_factor,
+            int                   s0,
+            int                   s1,
+            int                   p0,
+            int                   p1,
+            int                   d0,
+            int                   d1);
+
     GGML_API struct ggml_tensor * ggml_conv_3d_direct(
             struct ggml_context * ctx,
             struct ggml_tensor  * a,   // kernel [KW, KH, KD, IC * OC]
diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c
index 87a329f26975..260d12d20ae9 100644
--- a/ggml/src/ggml-cpu/ggml-cpu.c
+++ b/ggml/src/ggml-cpu/ggml-cpu.c
@@ -2141,6 +2141,12 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
                 ggml_compute_forward_opt_step_sgd(params, tensor);
             }
             break;
+        case GGML_OP_MUL_MAT_SEGMENTED:
+        case GGML_OP_QKNORM_ROPE:
+        case GGML_OP_GROUP_NORM_AFFINE_SILU:
+        case GGML_OP_CONV_2D_BIAS:
+        case GGML_OP_CONV_2D_UPSCALE:
+            GGML_ABORT("operation is not supported by the CPU backend");
         case GGML_OP_NONE:
             {
                 // nop
@@ -2489,6 +2495,14 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) {
             {
                 n_tasks = n_threads;
             } break;
+        case GGML_OP_MUL_MAT_SEGMENTED:
+        case GGML_OP_QKNORM_ROPE:
+        case GGML_OP_GROUP_NORM_AFFINE_SILU:
+        case GGML_OP_CONV_2D_BIAS:
+        case GGML_OP_CONV_2D_UPSCALE:
+            {
+                n_tasks = 1;
+            } break;
         case GGML_OP_NONE:
             {
                 n_tasks = 1;
diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp
index 8cece71f186f..60f17c16828f 100644
--- a/ggml/src/ggml-cpu/ggml-cpu.cpp
+++ b/ggml/src/ggml-cpu/ggml-cpu.cpp
@@ -472,6 +472,12 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st
                 src1->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
         case GGML_OP_CONV_2D:
             return ggml_is_contiguous(op->src[0]);
+        case GGML_OP_MUL_MAT_SEGMENTED:
+        case GGML_OP_QKNORM_ROPE:
+        case GGML_OP_GROUP_NORM_AFFINE_SILU:
+        case GGML_OP_CONV_2D_BIAS:
+        case GGML_OP_CONV_2D_UPSCALE:
+            return false;
         case GGML_OP_SSM_SCAN:
             return ggml_get_op_params_i32(op, 0) == 1 || op->src[3]->ne[0] == 1;
         default:
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index 266261c5e5a4..e01273bf5ef5 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -5907,6 +5907,7 @@ void ggml_compute_forward_clamp(
                 ggml_compute_forward_clamp_f16(params, dst);
             } break;
         case GGML_TYPE_BF16:
+        case GGML_TYPE_F8_E4M3:
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
         case GGML_TYPE_Q4_0:
diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
index ec7801388689..e78a5b1c5ab4 100644
--- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp
+++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
@@ -57,6 +57,8 @@
 #include "htp/get-rows-ops.h"
 #include "htp/set-rows-ops.h"
 #include "htp/rope-ops.h"
+#include "htp/conv2d-ops.h"
+#include "htp/groupnorm-ops.h"
 #include "htp_iface.h"
 #include "htp-drv.h"
 
@@ -112,6 +114,8 @@ enum ggml_hexagon_fusion_flags {
     GGML_HEXAGON_FUSE_MUL_MAT_ADD   = (1 << 3), // 8
     GGML_HEXAGON_FUSE_MUL_MAT_NX    = (1 << 4), // 16
     GGML_HEXAGON_FUSE_MUL_MAT_ID_NX = (1 << 5), // 32
+    GGML_HEXAGON_FUSE_MUL_MAT_SCALE = (1 << 6), // 64
+    GGML_HEXAGON_FUSE_CONV2D_ADD     = (1 << 7), // 128
 };
 
 static inline bool ggml_hexagon_is_fusion_enabled(int flag) {
@@ -250,13 +254,17 @@ enum ggml_hexagon_tensor_flags {
 static inline bool ggml_hexagon_is_repack_type(enum ggml_type type) {
     return type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q4_1 ||
            type == GGML_TYPE_Q8_0 || type == GGML_TYPE_IQ4_NL ||
-           type == GGML_TYPE_MXFP4;
+           type == GGML_TYPE_MXFP4 || type == GGML_TYPE_F8_E4M3;
 }
 
 static inline bool ggml_hexagon_is_hmx_weight_type(enum ggml_type type) {
     return type == GGML_TYPE_F16 || type == GGML_TYPE_F32 || ggml_hexagon_is_repack_type(type);
 }
 
+static inline enum ggml_type ggml_hexagon_hmx_weight_storage_type(enum ggml_type type) {
+    return type == GGML_TYPE_BF16 ? GGML_TYPE_F16 : type;
+}
+
 struct ggml_hexagon_session;
 
 static void ggml_hexagon_precompute_matmul_params(
@@ -267,6 +275,15 @@ static void ggml_hexagon_precompute_matmul_params(
     struct htp_mm_kernel_params * kparams
 );
 
+static void ggml_hexagon_precompute_segmented_matmul_params(
+    const struct ggml_hexagon_session * sess,
+    const struct ggml_tensor * src0,
+    const struct ggml_tensor * src1,
+    const struct ggml_tensor * src2,
+    const struct ggml_tensor * dst,
+    struct htp_mm_kernel_params * kparams
+);
+
 static void ggml_hexagon_precompute_fused_matmul_add_params(
     const struct ggml_hexagon_session * sess,
     const struct ggml_tensor * src0,
@@ -1302,6 +1319,168 @@ static void repack_tiled_mxfp4(void * data, const ggml_tensor * t, size_t offset
     }
 }
 
+static inline size_t f8_e4m3_tile_offset(int row, int col) {
+    static const uint8_t lane[4] = { 0, 2, 1, 3 };
+    const int tile_row = (col / 4) * 4 + row / 8;
+    const int tile_col = (row % 8) * 4 + lane[col % 4];
+    return (size_t) tile_row * 32 + tile_col;
+}
+
+static void repack_f8_e4m3_tiled(ggml_tensor * t, const void * data, size_t size) {
+    GGML_ASSERT(size == ggml_nbytes(t));
+
+    const uint8_t * src = (const uint8_t *) data;
+    const int64_t ne0 = t->ne[0];
+    const int64_t ne1 = t->ne[1];
+    const int64_t ne2 = t->ne[2];
+    const int64_t ne3 = t->ne[3];
+    const int64_t ne0_padded = hex_round_up(ne0, 32);
+    const int64_t ne1_padded = hex_round_up(ne1, 32);
+    const int n_k_tiles = ne0_padded / 32;
+    const int n_col_tiles = ne1_padded / 32;
+    const size_t matrix_size = (size_t) n_col_tiles * n_k_tiles * HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3;
+
+    for (int64_t i3 = 0; i3 < ne3; ++i3) {
+        for (int64_t i2 = 0; i2 < ne2; ++i2) {
+            const uint8_t * matrix_src = src + (i3 * ne2 + i2) * ne1 * ne0;
+            uint8_t * matrix_dst = (uint8_t *) t->data + (i3 * ne2 + i2) * matrix_size;
+
+            for (int ct = 0; ct < n_col_tiles; ++ct) {
+                for (int kt = 0; kt < n_k_tiles; ++kt) {
+                    uint8_t * tile = matrix_dst + ((size_t) ct * n_k_tiles + kt) * HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3;
+                    for (int row = 0; row < 32; ++row) {
+                        const int64_t j = (int64_t) ct * 32 + row;
+                        for (int col = 0; col < 32; ++col) {
+                            const int64_t i = (int64_t) kt * 32 + col;
+                            const uint8_t value = i < ne0 && j < ne1 ? matrix_src[j * ne0 + i] : 0;
+                            tile[f8_e4m3_tile_offset(row, col)] = value == 0x80 ? 0 : value;
+                        }
+                    }
+                }
+            }
+        }
+    }
+}
+
+static void repack_tiled_f8_e4m3(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    GGML_ASSERT(offset + size <= ggml_nbytes(t));
+
+    uint8_t * dst = (uint8_t *) data;
+    const int64_t ne0 = t->ne[0];
+    const int64_t ne1 = t->ne[1];
+    const int64_t ne2 = t->ne[2];
+    const int64_t ne0_padded = hex_round_up(ne0, 32);
+    const int64_t ne1_padded = hex_round_up(ne1, 32);
+    const int n_k_tiles = ne0_padded / 32;
+    const int n_col_tiles = ne1_padded / 32;
+    const size_t matrix_size = (size_t) n_col_tiles * n_k_tiles * HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3;
+    const size_t slice_size = (size_t) ne0 * ne1;
+
+    for (size_t p = 0; p < size; ++p) {
+        const size_t index = offset + p;
+        const size_t slice = index / slice_size;
+        const size_t in_slice = index - slice * slice_size;
+        const int row = in_slice / ne0;
+        const int col = in_slice - (size_t) row * ne0;
+        const int ct = row / 32;
+        const int kt = col / 32;
+        const uint8_t * tile = (const uint8_t *) t->data + slice * matrix_size + ((size_t) ct * n_k_tiles + kt) * HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3;
+        dst[p] = tile[f8_e4m3_tile_offset(row % 32, col % 32)];
+    }
+
+    GGML_UNUSED(ne2);
+}
+
+static bool ggml_hexagon_is_conv2d_hmx_weight(const ggml_tensor * tensor) {
+    const int64_t kw = tensor->ne[0];
+    const int64_t kh = tensor->ne[1];
+    return tensor->type == GGML_TYPE_F16 &&
+           ((kw == 1 && kh == 1) || (kw == 3 && kh == 3)) &&
+           tensor->ne[2] % 32 == 0;
+}
+
+static void repack_conv2d_f16_hmx(ggml_tensor * tensor, const void * data, size_t size) {
+    GGML_ASSERT(tensor->type == GGML_TYPE_F16);
+    GGML_ASSERT(size == ggml_nbytes(tensor));
+
+    const uint32_t kw = (uint32_t) tensor->ne[0];
+    const uint32_t kh = (uint32_t) tensor->ne[1];
+    const uint32_t ic = (uint32_t) tensor->ne[2];
+    const uint32_t oc = (uint32_t) tensor->ne[3];
+    const uint32_t ic_blocks = ic / 32;
+    const uint32_t k_tiles = kh * kw * ic_blocks;
+    const uint32_t oc_padded = hex_round_up(oc, 32);
+    const ggml_fp16_t * src = (const ggml_fp16_t *) data;
+    ggml_fp16_t * dst = (ggml_fp16_t *) tensor->data;
+    memset(dst, 0, (size_t) kh * kw * ic * oc_padded * sizeof(*dst));
+
+    for (uint32_t o = 0; o < oc; ++o) {
+        const uint32_t nt = o / 32;
+        const uint32_t nr = o % 32;
+        for (uint32_t y = 0; y < kh; ++y) {
+            for (uint32_t x = 0; x < kw; ++x) {
+                for (uint32_t c = 0; c < ic; ++c) {
+                    const uint32_t kt = (y * kw + x) * ic_blocks + c / 32;
+                    const uint32_t kr = c % 32;
+                    const size_t tile = ((size_t) nt * k_tiles + kt) * HTP_CONV2D_TILE_BYTES / sizeof(*dst);
+                    const size_t tiled_index = tile + (nr / 2) * 64 + kr * 2 + (nr & 1);
+                    const size_t source_index = x + (size_t) kw * (y + (size_t) kh * (c + (size_t) ic * o));
+                    dst[tiled_index] = src[source_index];
+                }
+            }
+        }
+    }
+}
+
+static void unpack_conv2d_f16_hmx(void * data, const ggml_tensor * tensor, size_t size) {
+    GGML_ASSERT(tensor->type == GGML_TYPE_F16);
+    GGML_ASSERT(size == ggml_nbytes(tensor));
+
+    const uint32_t kw = (uint32_t) tensor->ne[0];
+    const uint32_t kh = (uint32_t) tensor->ne[1];
+    const uint32_t ic = (uint32_t) tensor->ne[2];
+    const uint32_t oc = (uint32_t) tensor->ne[3];
+    const uint32_t ic_blocks = ic / 32;
+    const uint32_t k_tiles = kh * kw * ic_blocks;
+    const ggml_fp16_t * src = (const ggml_fp16_t *) tensor->data;
+    ggml_fp16_t * dst = (ggml_fp16_t *) data;
+
+    for (uint32_t o = 0; o < oc; ++o) {
+        const uint32_t nt = o / 32;
+        const uint32_t nr = o % 32;
+        for (uint32_t y = 0; y < kh; ++y) {
+            for (uint32_t x = 0; x < kw; ++x) {
+                for (uint32_t c = 0; c < ic; ++c) {
+                    const uint32_t kt = (y * kw + x) * ic_blocks + c / 32;
+                    const uint32_t kr = c % 32;
+                    const size_t tile = ((size_t) nt * k_tiles + kt) * HTP_CONV2D_TILE_BYTES / sizeof(*src);
+                    const size_t tiled_index = tile + (nr / 2) * 64 + kr * 2 + (nr & 1);
+                    const size_t dest_index = x + (size_t) kw * (y + (size_t) kh * (c + (size_t) ic * o));
+                    dst[dest_index] = src[tiled_index];
+                }
+            }
+        }
+    }
+}
+
+static void convert_bf16_weights_to_f16(void * dst, const void * src, size_t size) {
+    GGML_ASSERT(size % sizeof(ggml_bf16_t) == 0);
+    const ggml_bf16_t * input = (const ggml_bf16_t *) src;
+    ggml_fp16_t * output = (ggml_fp16_t *) dst;
+    for (size_t i = 0; i < size / sizeof(*input); ++i) {
+        output[i] = ggml_fp32_to_fp16(ggml_bf16_to_fp32(input[i]));
+    }
+}
+
+static void convert_f16_weights_to_bf16(void * dst, const void * src, size_t size) {
+    GGML_ASSERT(size % sizeof(ggml_fp16_t) == 0);
+    const ggml_fp16_t * input = (const ggml_fp16_t *) src;
+    ggml_bf16_t * output = (ggml_bf16_t *) dst;
+    for (size_t i = 0; i < size / sizeof(*input); ++i) {
+        output[i] = ggml_fp32_to_bf16(ggml_fp16_to_fp32(input[i]));
+    }
+}
+
 static void repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t size) {
     switch (tensor->type) {
         case GGML_TYPE_Q4_0:
@@ -1324,6 +1503,10 @@ static void repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t
             repack_mxfp4_tiled(tensor, data, 0, size);
             break;
 
+        case GGML_TYPE_F8_E4M3:
+            repack_f8_e4m3_tiled(tensor, data, size);
+            break;
+
         default:
             break;
     }
@@ -1348,6 +1531,20 @@ static void ggml_backend_hexagon_buffer_set_tensor(ggml_backend_buffer_t buffer,
     HEX_VERBOSE("ggml-hex: %s set-tensor %s : data %p offset %zu size %zu usage %d flags 0x%x\n",
         sess->c_name(), tensor->name, data, offset, size, (int) buffer->usage, extra->flags);
 
+    if (ggml_backend_buffer_get_usage(buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS &&
+        ggml_hexagon_is_conv2d_hmx_weight(tensor)) {
+        GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor));
+        repack_conv2d_f16_hmx(tensor, data, size);
+        return;
+    }
+
+    if (ggml_backend_buffer_get_usage(buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS &&
+        tensor->type == GGML_TYPE_BF16) {
+        GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor));
+        convert_bf16_weights_to_f16(tensor->data, data, size);
+        return;
+    }
+
     if ((extra->flags & GGML_HEXAGON_TENSOR_REPACK) == 0) {
         memcpy((char *) tensor->data + offset, data, size);
         return;
@@ -1384,6 +1581,20 @@ static void ggml_backend_hexagon_buffer_get_tensor(ggml_backend_buffer_t buffer,
     HEX_VERBOSE("ggml-hex: %s get-tensor %s : data %p offset %zu size %zu usage %d flags 0x%x\n",
             sess->c_name(), tensor->name, data, offset, size, (int) buffer->usage, extra->flags);
 
+    if (ggml_backend_buffer_get_usage(buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS &&
+        ggml_hexagon_is_conv2d_hmx_weight(tensor)) {
+        GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor));
+        unpack_conv2d_f16_hmx(data, tensor, size);
+        return;
+    }
+
+    if (ggml_backend_buffer_get_usage(buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS &&
+        tensor->type == GGML_TYPE_BF16) {
+        GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor));
+        convert_f16_weights_to_bf16(data, tensor->data, size);
+        return;
+    }
+
     if ((extra->flags & GGML_HEXAGON_TENSOR_REPACK) == 0) {
         memcpy(data, (const char *) tensor->data + offset, size);
         return;
@@ -1420,6 +1631,10 @@ static void ggml_backend_hexagon_buffer_get_tensor(ggml_backend_buffer_t buffer,
             repack_tiled_mxfp4(data, tensor, offset, size);
             break;
 
+        case GGML_TYPE_F8_E4M3:
+            repack_tiled_f8_e4m3(data, tensor, offset, size);
+            break;
+
         default:
             memcpy(data, (const char *) tensor->data + offset, size);
             break;
@@ -1640,6 +1855,9 @@ static size_t ggml_backend_hexagon_buffer_type_get_alignment(ggml_backend_buffer
 }
 
 static size_t ggml_backend_hexagon_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const struct ggml_tensor * t) {
+    if (ggml_hexagon_is_conv2d_hmx_weight(t)) {
+        return (size_t) t->ne[0] * t->ne[1] * t->ne[2] * hex_round_up(t->ne[3], 32) * sizeof(ggml_fp16_t);
+    }
     if (ggml_hexagon_is_repack_type(t->type)) {
         int64_t ne0 = hex_round_up(t->ne[0], 32);
         int64_t ne1 = hex_round_up(t->ne[1], 32);
@@ -1810,7 +2028,9 @@ struct ggml_hexagon_opbatch {
         int64_t nb2 = is_repack ? nb1 * ne1      : t->nb[2];
         int64_t nb3 = is_repack ? nb2 * t->ne[2] : t->nb[3];
 
-        return (h->type == t->type) &&
+        const ggml_type storage_type = (extra->flags & GGML_HEXAGON_TENSOR_WEIGHT) != 0
+            ? ggml_hexagon_hmx_weight_storage_type(t->type) : t->type;
+        return (h->type == storage_type) &&
                (h->ne[0] == ne0) && (h->ne[1] == ne1) && (h->ne[2] == t->ne[2]) && (h->ne[3] == t->ne[3]) &&
                (h->nb[0] == t->nb[0]) && (h->nb[1] == nb1) && (h->nb[2] == nb2) && (h->nb[3] == nb3);
     }
@@ -1845,7 +2065,8 @@ struct ggml_hexagon_opbatch {
         h.bi    = add_buffer(sbuf);
         h.ti    = ti;
         h.data  = t_offset;
-        h.type  = t->type;
+        h.type  = (extra->flags & GGML_HEXAGON_TENSOR_WEIGHT) != 0
+            ? ggml_hexagon_hmx_weight_storage_type(t->type) : t->type;
 
         const bool is_repack = (extra->flags & GGML_HEXAGON_TENSOR_REPACK) != 0;
         if (is_repack) {
@@ -2272,6 +2493,167 @@ struct ggml_hexagon_opbatch {
         return true;
     }
 
+    bool try_fuse_mul_mat_scale(const htp_opnode & node) {
+        if (n_ops == 0 || (node.opcode != HTP_OP_SCALE && node.opcode != HTP_OP_MUL)) {
+            return false;
+        }
+
+        htp_opnode & last_node = ops[n_ops - 1];
+        if (last_node.opcode != HTP_OP_MUL_MAT && last_node.opcode != HTP_OP_MUL_MAT_SEGMENTED) {
+            return false;
+        }
+        if (last_node.inputs.empty() || last_node.inputs[0]->type != GGML_TYPE_F8_E4M3) {
+            return false;
+        }
+
+        const ggml_tensor * mm_out = last_node.dst();
+        if (!ggml_hexagon_tensor_is_fuseable(mm_out)) {
+            return false;
+        }
+
+        const ggml_tensor * scale = nullptr;
+        auto * kparams = (struct htp_mm_kernel_params *) last_node.kernel_params;
+        if (node.opcode == HTP_OP_SCALE) {
+            if (node.src0() != mm_out || (kparams->scale_flags & HTP_MM_SCALE_PARAM)) {
+                return false;
+            }
+            float bias;
+            memcpy(&bias, &node.node->op_params[1], sizeof(bias));
+            if (bias != 0.0f) {
+                return false;
+            }
+        } else {
+            if (kparams->scale_flags & HTP_MM_SCALE_TENSOR) {
+                return false;
+            }
+            if (node.src0() == mm_out) {
+                scale = node.src1();
+            } else if (node.src1() == mm_out) {
+                scale = node.src0();
+            } else {
+                return false;
+            }
+            if (!scale || scale->type != GGML_TYPE_F32 || ggml_nelements(scale) != 1) {
+                return false;
+            }
+        }
+
+        size_t extra_bufs = 0;
+        size_t extra_vmem = 0;
+        size_t extra_tens = 0;
+        auto fit_t = [&](const ggml_tensor * t) {
+            if (!t || t_map.count(t)) {
+                return;
+            }
+            extra_tens++;
+            auto sbuf = static_cast<ggml_hexagon_shared_buffer *>(t->buffer->context);
+            if (!b_map.count(sbuf->fd())) {
+                extra_vmem += sbuf->size();
+                extra_bufs++;
+            }
+        };
+        fit_t(scale);
+        fit_t(node.dst());
+        if (extra_bufs + n_bufs > n_bufs_max || extra_tens + n_tens > n_tens_max || extra_vmem + b_vmem > b_vmem_max) {
+            return false;
+        }
+
+        if (node.opcode == HTP_OP_SCALE) {
+            kparams->scale_flags |= HTP_MM_SCALE_PARAM;
+        } else {
+            kparams->scale_flags |= HTP_MM_SCALE_TENSOR;
+        }
+        last_node.add_fused(node.node);
+
+        htp_op_desc & o = h_ops[n_ops - 1];
+        if (node.opcode == HTP_OP_SCALE) {
+            memcpy(o.params, node.node->op_params, sizeof(o.params));
+        }
+        memcpy(o.kernel_params, last_node.kernel_params, sizeof(o.kernel_params));
+        for (uint32_t s = 0; s < HTP_OP_MAX_INPUTS; s++) {
+            o.src[s] = s < last_node.inputs.size() && last_node.inputs[s] ? add_tensor(last_node.inputs[s]) : 0xffff;
+        }
+        o.dst[0] = add_tensor(node.dst());
+        for (uint32_t d = 1; d < HTP_OP_MAX_OUTPUTS; d++) {
+            o.dst[d] = 0xffff;
+        }
+
+        HEX_VERBOSE("ggml-hex: %s fused %s (#%u)\n", sess->c_name(), last_node.name.c_str(), n_ops - 1);
+        return true;
+    }
+
+    bool try_fuse_conv2d_add(const htp_opnode & node) {
+        if (n_ops == 0 || node.opcode != HTP_OP_ADD) {
+            return false;
+        }
+
+        htp_opnode & last_node = ops[n_ops - 1];
+        if (last_node.opcode != HTP_OP_CONV_2D || last_node.inputs.size() < 2) {
+            return false;
+        }
+
+        const ggml_tensor * conv_out = last_node.dst();
+        if (!ggml_hexagon_tensor_is_fuseable(conv_out)) {
+            return false;
+        }
+        const ggml_tensor * extra = node.src0() == conv_out ? node.src1() :
+                                    node.src1() == conv_out ? node.src0() : nullptr;
+        if (!extra) {
+            return false;
+        }
+
+        auto * params = (struct htp_conv2d_kernel_params *) last_node.kernel_params;
+        uint32_t new_flags = params->flags;
+        if ((params->flags & HTP_CONV2D_BIAS) == 0) {
+            if (extra->type != GGML_TYPE_F32 || ggml_nelements(extra) != conv_out->ne[2]) {
+                return false;
+            }
+            new_flags |= HTP_CONV2D_BIAS;
+        } else {
+            if ((params->flags & HTP_CONV2D_RESIDUAL) != 0 || conv_out->type != GGML_TYPE_F32 ||
+                extra->type != conv_out->type || !ggml_are_same_shape(extra, conv_out) ||
+                !ggml_is_contiguous(extra) || last_node.inputs[1]->data == node.dst()->data) {
+                return false;
+            }
+            new_flags |= HTP_CONV2D_RESIDUAL;
+        }
+
+        size_t extra_bufs = 0;
+        size_t extra_vmem = 0;
+        size_t extra_tens = 0;
+        auto fit_t = [&](const ggml_tensor * t) {
+            if (t_map.count(t)) {
+                return;
+            }
+            extra_tens++;
+            auto sbuf = static_cast<ggml_hexagon_shared_buffer *>(t->buffer->context);
+            if (!b_map.count(sbuf->fd())) {
+                extra_vmem += sbuf->size();
+                extra_bufs++;
+            }
+        };
+        fit_t(extra);
+        fit_t(node.dst());
+        if (extra_bufs + n_bufs > n_bufs_max || extra_tens + n_tens > n_tens_max || extra_vmem + b_vmem > b_vmem_max) {
+            return false;
+        }
+
+        params->flags = new_flags;
+        last_node.add_fused(node.node);
+        htp_op_desc & o = h_ops[n_ops - 1];
+        memcpy(o.kernel_params, last_node.kernel_params, sizeof(o.kernel_params));
+        for (uint32_t s = 0; s < HTP_OP_MAX_INPUTS; s++) {
+            o.src[s] = s < last_node.inputs.size() && last_node.inputs[s] ? add_tensor(last_node.inputs[s]) : 0xffff;
+        }
+        o.dst[0] = add_tensor(node.dst());
+        for (uint32_t d = 1; d < HTP_OP_MAX_OUTPUTS; d++) {
+            o.dst[d] = 0xffff;
+        }
+
+        HEX_VERBOSE("ggml-hex: %s fused %s (#%u)\n", sess->c_name(), last_node.name.c_str(), n_ops - 1);
+        return true;
+    }
+
     bool try_fuse_mul_mat_nx(const htp_opnode & node) {
         if (n_ops == 0 || node.opcode != HTP_OP_MUL_MAT) return false;
         if (!is_mergeable_mul_mat(node.node)) return false;
@@ -2597,6 +2979,8 @@ struct ggml_hexagon_opbatch {
         if (ggml_hexagon_is_fusion_enabled(GGML_HEXAGON_FUSE_ALLREDUCE_ADD) && try_fuse_allreduce_add(node)) return true;
         if (ggml_hexagon_is_fusion_enabled(GGML_HEXAGON_FUSE_RMS_NORM_MUL)  && try_fuse_rms_norm_mul(node))  return true;
         if (ggml_hexagon_is_fusion_enabled(GGML_HEXAGON_FUSE_MUL_MAT_ADD)   && try_fuse_mul_mat_add(node))   return true;
+        if (ggml_hexagon_is_fusion_enabled(GGML_HEXAGON_FUSE_MUL_MAT_SCALE) && try_fuse_mul_mat_scale(node)) return true;
+        if (ggml_hexagon_is_fusion_enabled(GGML_HEXAGON_FUSE_CONV2D_ADD)     && try_fuse_conv2d_add(node))     return true;
         if (ggml_hexagon_is_fusion_enabled(GGML_HEXAGON_FUSE_MUL_MAT_NX)    && try_fuse_mul_mat_nx(node))    return true;
         if (ggml_hexagon_is_fusion_enabled(GGML_HEXAGON_FUSE_MUL_MAT_ID_NX) && try_fuse_mul_mat_id_nx(node)) return true;
         return false;
@@ -3838,7 +4222,7 @@ static bool ggml_hexagon_matmul_is_hmx_eligible(
     const int ne00  = src0->ne[0];
     const int ne11  = src1->ne[1];
     const int ne12  = src1->ne[2];
-    const int wtype = src0->type;
+    const int wtype = ggml_hexagon_hmx_weight_storage_type(src0->type);
 
     // HMX weight tile requires N to be 32-aligned.
     if (ne01_padded % 32 != 0) {
@@ -3868,7 +4252,7 @@ static bool ggml_hexagon_matmul_is_hmx_eligible(
     // M alignment: Use HMX when M > HTP_MM_HMX_MIN_NROWS.
     // For MUL_MAT_ID, src1 shape is [K, n_expert_used, n_tokens, 1], so n_tokens is ne12.
     const int m = is_matmul_id ? ne12 : ne11;
-    if (m <= HTP_MM_HMX_MIN_NROWS) {
+    if (m <= HTP_MM_HMX_MIN_NROWS && wtype != GGML_TYPE_F8_E4M3) {
         return false;
     }
 
@@ -4161,7 +4545,7 @@ static void ggml_hexagon_precompute_matmul_params_impl(
     const int ne12 = src1->ne[2];
     const int ne13 = src1->ne[3];
 
-    const int wtype = src0->type;
+    const int wtype = ggml_hexagon_hmx_weight_storage_type(src0->type);
     const bool is_repack = ggml_hexagon_is_repack_type((ggml_type) wtype);
     const int ne00_padded = is_repack ? hex_round_up(ne00, 32) : ne00;
     const int ne01_padded = is_repack ? hex_round_up(ne01, 32) : ne01;
@@ -4201,6 +4585,23 @@ static void ggml_hexagon_precompute_matmul_params(
     ggml_hexagon_precompute_matmul_params_impl(sess, src0, src1, dst, 0, kparams);
 }
 
+static void ggml_hexagon_precompute_segmented_matmul_params(
+    const struct ggml_hexagon_session * sess,
+    const struct ggml_tensor * src0,
+    const struct ggml_tensor * src1,
+    const struct ggml_tensor * src2,
+    const struct ggml_tensor * dst,
+    struct htp_mm_kernel_params * kparams
+) {
+    struct ggml_tensor joined = *src1;
+    joined.ne[0] = src1->ne[0] + src2->ne[0];
+    joined.nb[0] = sizeof(float);
+    joined.nb[1] = joined.ne[0] * joined.nb[0];
+    joined.nb[2] = joined.ne[1] * joined.nb[1];
+    joined.nb[3] = joined.ne[2] * joined.nb[2];
+    ggml_hexagon_precompute_matmul_params_impl(sess, src0, &joined, dst, 0, kparams);
+}
+
 static void ggml_hexagon_precompute_fused_matmul_add_params(
     const struct ggml_hexagon_session * sess,
     const struct ggml_tensor * src0,
@@ -4547,6 +4948,10 @@ static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * s
         case GGML_TYPE_Q8_0:
         case GGML_TYPE_IQ4_NL:
         case GGML_TYPE_MXFP4:
+        case GGML_TYPE_F8_E4M3:
+            if (src0->type == GGML_TYPE_F8_E4M3 && (opt_arch < 79 || src1->type != GGML_TYPE_F32)) {
+                return false;
+            }
             if (src0->ne[0] % 32) {
                 return false;
             }
@@ -4560,6 +4965,7 @@ static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * s
             }
             break;
 
+        case GGML_TYPE_BF16:
         case GGML_TYPE_F16:
             if (src0->nb[1] < src0->nb[0]) {
                 return false;
@@ -4587,6 +4993,9 @@ static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * s
 
     struct htp_mm_kernel_params kparams;
     ggml_hexagon_precompute_matmul_params(sess, src0, src1, dst, &kparams);
+    if (src0->type == GGML_TYPE_F8_E4M3 && kparams.n_hmx == 0) {
+        return false;
+    }
     if ((size_t)kparams.vtcm_size > sess->vtcm_size) {
         HEX_VERBOSE("ggml-hex: %s supported MUL_MAT VTCM size needed (%d) > budget (%zu)\n", sess->c_name(), kparams.vtcm_size, sess->vtcm_size);
         return false;
@@ -4595,6 +5004,38 @@ static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * s
     return true;
 }
 
+static bool ggml_hexagon_supported_mul_mat_segmented(const struct ggml_hexagon_session * sess, const struct ggml_tensor * dst) {
+    const struct ggml_tensor * src0 = dst->src[0];
+    const struct ggml_tensor * src1 = dst->src[1];
+    const struct ggml_tensor * src2 = dst->src[2];
+
+    if (!src0 || !src1 || !src2 || src0->type != GGML_TYPE_F8_E4M3 ||
+        src1->type != GGML_TYPE_F32 || src2->type != GGML_TYPE_F32 ||
+        dst->type != GGML_TYPE_F32 || opt_arch < 79) {
+        return false;
+    }
+    if (src0->ne[0] != src1->ne[0] + src2->ne[0] || src0->ne[1] != dst->ne[0] ||
+        src1->ne[1] != src2->ne[1] || src1->ne[1] != dst->ne[1] ||
+        src1->ne[2] != src2->ne[2] || src1->ne[2] != dst->ne[2] ||
+        src1->ne[3] != src2->ne[3] || src1->ne[3] != dst->ne[3] ||
+        src1->ne[2] != 1 || src1->ne[3] != 1) {
+        return false;
+    }
+    if (src1->ne[0] % 32 || src2->ne[0] % 32 ||
+        src1->nb[0] != sizeof(float) || src2->nb[0] != sizeof(float) ||
+        src1->nb[1] < src1->ne[0] * sizeof(float) ||
+        src2->nb[1] < src2->ne[0] * sizeof(float)) {
+        return false;
+    }
+    if (!src0->buffer) {
+        sess->needs_repack.insert(src0);
+    }
+
+    struct htp_mm_kernel_params kparams;
+    ggml_hexagon_precompute_segmented_matmul_params(sess, src0, src1, src2, dst, &kparams);
+    return kparams.n_hmx > 0 && (size_t) kparams.vtcm_size <= sess->vtcm_size;
+}
+
 static bool ggml_hexagon_supported_mul_mat_id(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
     const struct ggml_tensor * src0 = op->src[0];
     const struct ggml_tensor * src1 = op->src[1];
@@ -5129,6 +5570,188 @@ static bool ggml_hexagon_supported_im2col(const struct ggml_hexagon_session * se
     return true;
 }
 
+static bool ggml_hexagon_precompute_conv2d_params(
+        const struct ggml_hexagon_session * sess,
+        const struct ggml_tensor * op,
+        struct htp_conv2d_kernel_params * kparams) {
+    const uint32_t kh = (uint32_t) op->src[0]->ne[1];
+    const uint32_t kw = (uint32_t) op->src[0]->ne[0];
+    const uint32_t ic = (uint32_t) op->src[0]->ne[2];
+    const uint32_t oc = (uint32_t) op->src[0]->ne[3];
+    const uint32_t ow = (uint32_t) op->ne[0];
+    const uint32_t oh = (uint32_t) op->ne[1];
+    const uint32_t padded_ow =
+        (ow + HTP_CONV2D_TILE_W - 1u) & ~(HTP_CONV2D_TILE_W - 1u);
+
+    auto max_tile_rows = [&](uint32_t tile_w, struct htp_conv2d_kernel_params * result) {
+        struct htp_conv2d_kernel_params candidate;
+        if (htp_conv2d_layout_build(&candidate, kh, kw, ic, oc, tile_w, 1,
+                                    sess->n_threads) > sess->vtcm_size) {
+            return 0u;
+        }
+
+        uint32_t lo = 1;
+        uint32_t hi = oh;
+        while (lo < hi) {
+            const uint32_t mid = lo + (hi - lo + 1) / 2;
+            if (htp_conv2d_layout_build(&candidate, kh, kw, ic, oc, tile_w, mid,
+                                        sess->n_threads) <= sess->vtcm_size) {
+                lo = mid;
+            } else {
+                hi = mid - 1;
+            }
+        }
+        htp_conv2d_layout_build(result, kh, kw, ic, oc, tile_w, lo, sess->n_threads);
+        return lo;
+    };
+
+    uint32_t max_m_tiles = 0;
+    for (uint32_t tile_w = HTP_CONV2D_TILE_W; tile_w <= padded_ow; tile_w += HTP_CONV2D_TILE_W) {
+        struct htp_conv2d_kernel_params candidate;
+        const uint32_t rows = max_tile_rows(tile_w, &candidate);
+        max_m_tiles = std::max(max_m_tiles, rows * (tile_w / HTP_CONV2D_TILE_W));
+    }
+    if (max_m_tiles == 0) {
+        memset(kparams, 0, sizeof(*kparams));
+        return false;
+    }
+
+    bool found = false;
+    for (uint32_t tile_w = HTP_CONV2D_TILE_W; tile_w <= padded_ow; tile_w += HTP_CONV2D_TILE_W) {
+        struct htp_conv2d_kernel_params candidate;
+        const uint32_t rows = max_tile_rows(tile_w, &candidate);
+        const uint32_t m_tiles = rows * (tile_w / HTP_CONV2D_TILE_W);
+        const bool enough_m_tiles =
+            kh * kw == 1 ? (uint64_t) m_tiles * 5 >= (uint64_t) max_m_tiles * 4 :
+                           (uint64_t) m_tiles * 20 >= (uint64_t) max_m_tiles * 19;
+        if (enough_m_tiles) {
+            *kparams = candidate;
+            found = true;
+        }
+    }
+    if (found) {
+        if (op->op == GGML_OP_CONV_2D_UPSCALE) {
+            kparams->flags |= HTP_CONV2D_UPSCALE2;
+        }
+        if (op->op == GGML_OP_CONV_2D_BIAS || op->src[2] != nullptr) {
+            kparams->flags |= HTP_CONV2D_BIAS;
+        }
+        HEX_VERBOSE("ggml-hex: conv2d tile %ux%u (%u M tiles), VTCM %u/%zu bytes\n",
+                    kparams->tile_w, kparams->tile_h, kparams->m_tiles,
+                    kparams->vtcm_size, sess->vtcm_size);
+    }
+    return found;
+}
+
+static bool ggml_hexagon_supported_conv2d(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
+    const struct ggml_tensor * weight = op->src[0];
+    const struct ggml_tensor * src = op->src[1];
+    const struct ggml_tensor * bias = op->src[2];
+    if (sess->n_hmx == 0 || !weight || !src || weight->type != GGML_TYPE_F16 ||
+        (src->type != GGML_TYPE_F32 && src->type != GGML_TYPE_F16) ||
+        (op->type != GGML_TYPE_F32 && op->type != GGML_TYPE_F16) ||
+        (src->type == GGML_TYPE_F16 && op->type != GGML_TYPE_F16)) {
+        return false;
+    }
+    if (!ggml_is_contiguous(weight) || !ggml_is_contiguous(src) || !ggml_is_contiguous(op) ||
+        src->ne[3] != 1 || op->ne[3] != 1) {
+        return false;
+    }
+    if ((op->op == GGML_OP_CONV_2D_BIAS && !bias) ||
+        (bias && (bias->type != GGML_TYPE_F32 || !ggml_is_contiguous(bias) ||
+                  ggml_nelements(bias) != weight->ne[3]))) {
+        return false;
+    }
+
+    const int32_t s0 = ggml_get_op_params_i32(op, 0);
+    const int32_t s1 = ggml_get_op_params_i32(op, 1);
+    const int32_t p0 = ggml_get_op_params_i32(op, 2);
+    const int32_t p1 = ggml_get_op_params_i32(op, 3);
+    const int32_t d0 = ggml_get_op_params_i32(op, 4);
+    const int32_t d1 = ggml_get_op_params_i32(op, 5);
+    const int32_t upscale = op->op == GGML_OP_CONV_2D_UPSCALE
+                                ? ggml_get_op_params_i32(op, 6)
+                                : 1;
+    const int64_t kw = weight->ne[0];
+    const int64_t kh = weight->ne[1];
+    const int64_t ic = weight->ne[2];
+    const int64_t oc = weight->ne[3];
+    if (!((kw == 1 && kh == 1) || (kw == 3 && kh == 3)) ||
+        s0 != 1 || s1 != 1 || d0 != 1 || d1 != 1 ||
+        (upscale != 1 && upscale != 2) ||
+        p0 != (kw - 1) / 2 || p1 != (kh - 1) / 2 ||
+        ic % 32 != 0 ||
+        src->ne[2] != ic || op->ne[2] != oc ||
+        op->ne[0] != src->ne[0] * upscale ||
+        op->ne[1] != src->ne[1] * upscale) {
+        return false;
+    }
+
+    struct htp_conv2d_kernel_params kparams;
+    if (!ggml_hexagon_precompute_conv2d_params(sess, op, &kparams)) {
+        return false;
+    }
+
+    return true;
+}
+
+static bool ggml_hexagon_supported_group_norm(const struct ggml_tensor * op) {
+    const struct ggml_tensor * src = op->src[0];
+    if (!src || !((src->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) ||
+                  (src->type == GGML_TYPE_F16 && op->type == GGML_TYPE_F16)) ||
+        !ggml_is_contiguous(src) || !ggml_is_contiguous(op) ||
+        !ggml_are_same_shape(src, op)) {
+        return false;
+    }
+    const int32_t groups = ggml_get_op_params_i32(op, 0);
+    return groups > 0 && src->ne[2] % groups == 0 &&
+           (src->type == GGML_TYPE_F16 ||
+            (src->ne[0] * src->ne[1]) % 32 == 0);
+}
+
+static bool ggml_hexagon_supported_group_norm_affine_silu(const struct ggml_tensor * op) {
+    const struct ggml_tensor * src    = op->src[0];
+    const struct ggml_tensor * weight = op->src[1];
+    const struct ggml_tensor * bias   = op->src[2];
+    if (!src || !weight || !bias || src->type != GGML_TYPE_F16 ||
+        op->type != GGML_TYPE_F16 || weight->type != GGML_TYPE_F32 ||
+        bias->type != GGML_TYPE_F32 || !ggml_is_contiguous(src) ||
+        !ggml_is_contiguous(op) || !ggml_is_contiguous(weight) ||
+        !ggml_is_contiguous(bias) || !ggml_are_same_shape(src, op) ||
+        ggml_nelements(weight) != src->ne[2] ||
+        ggml_nelements(bias) != src->ne[2]) {
+        return false;
+    }
+
+    const int32_t groups = ggml_get_op_params_i32(op, 0);
+    return groups > 0 && src->ne[2] % groups == 0;
+}
+
+static bool ggml_hexagon_supported_qknorm_rope(const struct ggml_tensor * op) {
+    const struct ggml_tensor * src    = op->src[0];
+    const struct ggml_tensor * weight = op->src[1];
+    const struct ggml_tensor * theta  = op->src[2];
+    if (!src || !weight || !theta ||
+        src->type != GGML_TYPE_F32 || weight->type != GGML_TYPE_F32 ||
+        theta->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32) {
+        return false;
+    }
+
+    const int64_t head_dim = src->ne[0];
+    return head_dim > 0 && head_dim <= 256 && head_dim % 64 == 0 &&
+           src->ne[1] > 0 && src->ne[2] > 0 && src->ne[3] > 0 &&
+           ggml_nelements(weight) == head_dim &&
+           ggml_nelements(theta) >= src->ne[2] * 2 * head_dim &&
+           src->nb[0] == sizeof(float) &&
+           src->nb[1] % sizeof(float) == 0 &&
+           src->nb[2] % sizeof(float) == 0 &&
+           src->nb[3] % sizeof(float) == 0 &&
+           ggml_is_contiguous(weight) && ggml_is_contiguous(theta) &&
+           ggml_is_contiguous(op) &&
+           op->ne[0] == head_dim && op->ne[1] == src->ne[2] &&
+           op->ne[2] == src->ne[1] * src->ne[3] && op->ne[3] == 1;
+}
+
 static bool ggml_hexagon_supported_pad(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
     const struct ggml_tensor * src0 = op->src[0];
     const struct ggml_tensor * dst  = op;
@@ -5252,6 +5875,7 @@ static htp_op_code op_remap_to_htp(const ggml_tensor * t) {
         case GGML_OP_FLASH_ATTN_EXT:  return HTP_OP_FLASH_ATTN_EXT;
         case GGML_OP_MUL_MAT:         return HTP_OP_MUL_MAT;
         case GGML_OP_MUL_MAT_ID:      return HTP_OP_MUL_MAT_ID;
+        case GGML_OP_MUL_MAT_SEGMENTED: return HTP_OP_MUL_MAT_SEGMENTED;
         case GGML_OP_MUL:             return HTP_OP_MUL;
         case GGML_OP_ADD:             return HTP_OP_ADD;
         case GGML_OP_ADD_ID:          return HTP_OP_ADD_ID;
@@ -5285,6 +5909,12 @@ static htp_op_code op_remap_to_htp(const ggml_tensor * t) {
         case GGML_OP_TRI:             return HTP_OP_TRI;
         case GGML_OP_PAD:             return HTP_OP_PAD;
         case GGML_OP_IM2COL:          return HTP_OP_IM2COL;
+        case GGML_OP_QKNORM_ROPE:     return HTP_OP_QKNORM_ROPE;
+        case GGML_OP_CONV_2D:
+        case GGML_OP_CONV_2D_BIAS:
+        case GGML_OP_CONV_2D_UPSCALE: return HTP_OP_CONV_2D;
+        case GGML_OP_GROUP_NORM:      return HTP_OP_GROUP_NORM;
+        case GGML_OP_GROUP_NORM_AFFINE_SILU: return HTP_OP_GROUP_NORM;
 
         case GGML_OP_UNARY:
             switch (ggml_get_unary_op(t)) {
@@ -5324,13 +5954,250 @@ static inline bool op_is_compute(ggml_tensor *node)
     return !ggml_op_is_empty(node->op) && !ggml_is_empty(node) && (node->flags & GGML_TENSOR_FLAG_COMPUTE);
 }
 
+static const ggml_tensor * ggml_hexagon_unwrap_empty(const ggml_tensor * tensor) {
+    while (tensor && ggml_op_is_empty(tensor->op)) {
+        const ggml_tensor * next = tensor->src[0] ? tensor->src[0] : tensor->view_src;
+        if (!next) {
+            break;
+        }
+        tensor = next;
+    }
+    return tensor;
+}
+
+static bool ggml_hexagon_has_unwrapped_src(const ggml_tensor * node, const ggml_tensor * src) {
+    for (int i = 0; i < GGML_MAX_SRC; ++i) {
+        if (ggml_hexagon_unwrap_empty(node->src[i]) == src) {
+            return true;
+        }
+    }
+    return false;
+}
+
+static bool ggml_hexagon_match_qknorm_rope(
+        const ggml_cgraph * graph,
+        int start,
+        std::vector<int> & matched,
+        const ggml_tensor ** src,
+        const ggml_tensor ** weight,
+        const ggml_tensor ** theta) {
+    static const ggml_op expected[] = {
+        GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_CONT, GGML_OP_CONT, GGML_OP_REPEAT,
+        GGML_OP_CONT, GGML_OP_MUL, GGML_OP_REPEAT, GGML_OP_MUL, GGML_OP_ADD,
+    };
+    static const ggml_op expected_with_pe_copy[] = {
+        GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_CONT, GGML_OP_CONT, GGML_OP_REPEAT,
+        GGML_OP_CONT, GGML_OP_CONT, GGML_OP_MUL, GGML_OP_REPEAT, GGML_OP_MUL, GGML_OP_ADD,
+    };
+
+    matched.clear();
+    for (int i = start; i < graph->n_nodes && matched.size() < std::size(expected_with_pe_copy); ++i) {
+        if (op_is_compute(graph->nodes[i])) {
+            matched.push_back(i);
+        }
+    }
+    if (matched.size() < std::size(expected)) {
+        return false;
+    }
+
+    bool matches_base = true;
+    for (size_t i = 0; i < std::size(expected); ++i) {
+        if (graph->nodes[matched[i]]->op != expected[i]) {
+            matches_base = false;
+            break;
+        }
+    }
+    if (matches_base) {
+        matched.resize(std::size(expected));
+    } else {
+        if (matched.size() != std::size(expected_with_pe_copy)) {
+            return false;
+        }
+        for (size_t i = 0; i < matched.size(); ++i) {
+            if (graph->nodes[matched[i]]->op != expected_with_pe_copy[i]) {
+                return false;
+            }
+        }
+    }
+
+    const ggml_tensor * rms = graph->nodes[matched[0]];
+    const ggml_tensor * norm_mul = graph->nodes[matched[1]];
+    const ggml_tensor * cont0 = graph->nodes[matched[2]];
+    const ggml_tensor * cont1 = graph->nodes[matched[3]];
+    const ggml_tensor * repeat0 = graph->nodes[matched[4]];
+    const ggml_tensor * pe_copy = matches_base ? nullptr : graph->nodes[matched[5]];
+    const size_t pe_idx = matches_base ? 5 : 6;
+    const ggml_tensor * pe_cont = graph->nodes[matched[pe_idx]];
+    const ggml_tensor * mul0 = graph->nodes[matched[pe_idx + 1]];
+    const ggml_tensor * repeat1 = graph->nodes[matched[pe_idx + 2]];
+    const ggml_tensor * mul1 = graph->nodes[matched[pe_idx + 3]];
+    const ggml_tensor * add = graph->nodes[matched[pe_idx + 4]];
+
+    if (!ggml_hexagon_has_unwrapped_src(norm_mul, rms) ||
+        !ggml_hexagon_has_unwrapped_src(cont0, norm_mul) ||
+        !ggml_hexagon_has_unwrapped_src(cont1, cont0) ||
+        !ggml_hexagon_has_unwrapped_src(repeat0, cont1) ||
+        (pe_copy && !ggml_hexagon_has_unwrapped_src(pe_cont, pe_copy)) ||
+        !ggml_hexagon_has_unwrapped_src(mul0, repeat0) ||
+        !ggml_hexagon_has_unwrapped_src(mul0, pe_cont) ||
+        !ggml_hexagon_has_unwrapped_src(repeat1, cont1) ||
+        !ggml_hexagon_has_unwrapped_src(mul1, repeat1) ||
+        !ggml_hexagon_has_unwrapped_src(mul1, pe_cont) ||
+        !ggml_hexagon_has_unwrapped_src(add, mul0) ||
+        !ggml_hexagon_has_unwrapped_src(add, mul1)) {
+        return false;
+    }
+
+    const ggml_tensor * norm_weight = norm_mul->src[0] == rms ? norm_mul->src[1] : norm_mul->src[0];
+    const ggml_tensor * pe_base = pe_copy ? pe_copy : ggml_hexagon_unwrap_empty(pe_cont->src[0]);
+    if (pe_base && pe_base->op == GGML_OP_CONT && pe_base->ne[0] == 2 && pe_base->ne[1] == 2) {
+        pe_base = ggml_hexagon_unwrap_empty(pe_base->src[0]);
+    }
+    const ggml_tensor * pe = ggml_hexagon_unwrap_empty(pe_base);
+    const ggml_tensor * input = rms->src[0];
+    if (!input || !norm_weight || !pe) {
+        return false;
+    }
+
+    const uint64_t n_output = (uint64_t) input->ne[0] * input->ne[1] * input->ne[2] * input->ne[3];
+    if (input->type != GGML_TYPE_F32 || norm_weight->type != GGML_TYPE_F32 ||
+        pe->type != GGML_TYPE_F32 || add->type != GGML_TYPE_F32 ||
+        input->ne[0] <= 0 || input->ne[0] > 256 || input->ne[0] % 64 != 0 ||
+        ggml_nelements(norm_weight) != input->ne[0] || ggml_nelements(add) != (int64_t) n_output ||
+        ggml_nelements(pe) < input->ne[2] * 2 * input->ne[0] ||
+        input->nb[0] != sizeof(float) || !ggml_is_contiguous(add)) {
+        return false;
+    }
+
+    *src = input;
+    *weight = norm_weight;
+    *theta = pe;
+    return true;
+}
+
+static bool ggml_hexagon_fusion_region_is_closed(
+        const ggml_cgraph * graph,
+        int start,
+        int end,
+        std::initializer_list<const ggml_tensor *> outputs) {
+    auto is_output = [&](const ggml_tensor * tensor) {
+        const ggml_tensor * unwrapped = ggml_hexagon_unwrap_empty(tensor);
+        return std::any_of(outputs.begin(), outputs.end(), [&](const ggml_tensor * output) {
+            for (const ggml_tensor * alias = output; alias; alias = alias->view_src) {
+                if (alias == tensor || ggml_hexagon_unwrap_empty(alias) == unwrapped) {
+                    return true;
+                }
+            }
+            return false;
+        });
+    };
+    auto is_internal = [&](const ggml_tensor * tensor) {
+        for (int i = start; i <= end; ++i) {
+            if (graph->nodes[i] == tensor) {
+                return true;
+            }
+        }
+        return false;
+    };
+
+    for (int i = start; i <= end; ++i) {
+        if ((graph->nodes[i]->flags & GGML_TENSOR_FLAG_OUTPUT) && !is_output(graph->nodes[i])) {
+            return false;
+        }
+    }
+    for (int i = end + 1; i < graph->n_nodes; ++i) {
+        const ggml_tensor * consumer = graph->nodes[i];
+        for (int s = 0; s < GGML_MAX_SRC; ++s) {
+            const ggml_tensor * dependency = ggml_hexagon_unwrap_empty(consumer->src[s]);
+            if (dependency && is_internal(dependency) && !is_output(dependency)) {
+                return false;
+            }
+        }
+    }
+    return true;
+}
+
+static htp_opnode ggml_hexagon_make_qknorm_rope_node(
+        const ggml_cgraph * graph,
+        const std::vector<int> & matched,
+        const ggml_tensor * src,
+        const ggml_tensor * weight,
+        const ggml_tensor * theta) {
+    htp_opnode node(HTP_OP_QKNORM_ROPE, graph->nodes[matched[0]]);
+    for (size_t i = 1; i < matched.size(); ++i) {
+        node.add_fused(graph->nodes[matched[i]]);
+    }
+    node.inputs = { src, weight, theta };
+    return node;
+}
+
+static bool ggml_hexagon_fuse_qknorm_rope(
+        const ggml_cgraph * graph,
+        int * node_idx,
+        std::vector<htp_opnode> * nodes) {
+    std::vector<int> matched;
+    const ggml_tensor * src = nullptr;
+    const ggml_tensor * weight = nullptr;
+    const ggml_tensor * theta = nullptr;
+    if (!ggml_hexagon_match_qknorm_rope(graph, *node_idx, matched, &src, &weight, &theta)) {
+        return false;
+    }
+
+    const int end = matched.back();
+    const ggml_tensor * output = graph->nodes[end];
+    if (ggml_hexagon_fusion_region_is_closed(graph, *node_idx, end, { output })) {
+        nodes->push_back(ggml_hexagon_make_qknorm_rope_node(graph, matched, src, weight, theta));
+        *node_idx = end;
+        return true;
+    }
+
+    const bool has_pe_copy = matched.size() == 11;
+    const size_t pe_pos = has_pe_copy ? 6 : 5;
+    ggml_tensor * pe_copy = has_pe_copy ? graph->nodes[matched[5]] : nullptr;
+    ggml_tensor * pe_cont = graph->nodes[matched[pe_pos]];
+    if (ggml_hexagon_fusion_region_is_closed(graph, *node_idx, end, { output, pe_copy, pe_cont })) {
+        if (pe_copy) {
+            nodes->emplace_back(op_remap_to_htp(pe_copy), pe_copy);
+        }
+        nodes->emplace_back(op_remap_to_htp(pe_cont), pe_cont);
+        nodes->push_back(ggml_hexagon_make_qknorm_rope_node(graph, matched, src, weight, theta));
+        *node_idx = end;
+        return true;
+    }
+
+    std::vector<int> matched_k;
+    const ggml_tensor * src_k = nullptr;
+    const ggml_tensor * weight_k = nullptr;
+    const ggml_tensor * theta_k = nullptr;
+    int k_start = end + 1;
+    while (k_start < graph->n_nodes && !op_is_compute(graph->nodes[k_start])) {
+        ++k_start;
+    }
+    if (k_start >= graph->n_nodes ||
+        !ggml_hexagon_match_qknorm_rope(graph, k_start, matched_k, &src_k, &weight_k, &theta_k) ||
+        theta_k != theta) {
+        return false;
+    }
+
+    const int end_k = matched_k.back();
+    const ggml_tensor * output_k = graph->nodes[end_k];
+    if (!ggml_hexagon_fusion_region_is_closed(graph, *node_idx, end_k, { output, output_k })) {
+        return false;
+    }
+
+    nodes->push_back(ggml_hexagon_make_qknorm_rope_node(graph, matched, src, weight, theta));
+    nodes->push_back(ggml_hexagon_make_qknorm_rope_node(graph, matched_k, src_k, weight_k, theta_k));
+    *node_idx = end_k;
+    return true;
+}
+
 static bool mm_is_hmx_eligible(const ggml_tensor * t) {
     if (opt_nhmx == 0) { return false; }
 
     const ggml_tensor * src0 = t->src[0];
     const ggml_tensor * src1 = t->src[1];
 
-    const int wtype = src0->type;
+    const int wtype = ggml_hexagon_hmx_weight_storage_type(src0->type);
     const bool is_repack    = ggml_hexagon_is_repack_type((ggml_type) wtype);
     const bool is_matmul_id = (t->op == GGML_OP_MUL_MAT_ID);
     const bool is_batched   = (src0->ne[2] * src0->ne[3] > 1 || src1->ne[2] * src1->ne[3] > 1);
@@ -5341,6 +6208,9 @@ static bool mm_is_hmx_eligible(const ggml_tensor * t) {
 }
 
 static bool is_supported_mul_mat_nx_kernel(const ggml_tensor * src0, const struct htp_mm_kernel_params * kparams) {
+    if (src0->type == GGML_TYPE_F8_E4M3) {
+        return false;
+    }
     if (kparams->n_hmx) {
         return kparams->kernel_type == HTP_MM_KERNEL_HMX_2D;
     }
@@ -5373,7 +6243,7 @@ static bool is_mergeable_mul_mat(const ggml_tensor * t) {
     if (src0->ne[2] != 1 || src0->ne[3] != 1) return false;
 
     if (mm_is_hmx_eligible(t)) {
-        return ggml_hexagon_is_hmx_weight_type(src0->type);
+        return ggml_hexagon_is_hmx_weight_type(ggml_hexagon_hmx_weight_storage_type(src0->type));
     }
 
     return ggml_hexagon_is_repack_type(src0->type);
@@ -5456,11 +6326,20 @@ static ggml_status ggml_backend_hexagon_graph_compute(ggml_backend_t backend, gg
 
             if (graph->nodes[i]->op == GGML_OP_RMS_NORM && ggml_can_fuse(graph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL })) {
                 extra->flags |= GGML_HEXAGON_TENSOR_FUSEABLE;
-            } else if (graph->nodes[i]->op == GGML_OP_MUL_MAT || graph->nodes[i]->op == GGML_OP_MUL_MAT_ID) {
+            } else if (graph->nodes[i]->op == GGML_OP_MUL_MAT || graph->nodes[i]->op == GGML_OP_MUL_MAT_ID ||
+                       graph->nodes[i]->op == GGML_OP_MUL_MAT_SEGMENTED) {
                 if ((i + 1 < graph->n_nodes && graph->nodes[i + 1]->op == GGML_OP_ADD && ggml_can_fuse(graph, i, { graph->nodes[i]->op, GGML_OP_ADD })) ||
                     ggml_node_has_n_uses(graph, i, 1)) {
                     extra->flags |= GGML_HEXAGON_TENSOR_FUSEABLE;
                 }
+            } else if (graph->nodes[i]->op == GGML_OP_SCALE && ggml_node_has_n_uses(graph, i, 1)) {
+                extra->flags |= GGML_HEXAGON_TENSOR_FUSEABLE;
+            } else if ((graph->nodes[i]->op == GGML_OP_CONV_2D ||
+                        graph->nodes[i]->op == GGML_OP_CONV_2D_BIAS ||
+                        graph->nodes[i]->op == GGML_OP_CONV_2D_UPSCALE ||
+                        graph->nodes[i]->op == GGML_OP_ADD) &&
+                       ggml_node_has_n_uses(graph, i, 1)) {
+                extra->flags |= GGML_HEXAGON_TENSOR_FUSEABLE;
             }
         }
 
@@ -5472,6 +6351,10 @@ static ggml_status ggml_backend_hexagon_graph_compute(ggml_backend_t backend, gg
                 continue;
             }
 
+            if (n->op == GGML_OP_RMS_NORM && ggml_hexagon_fuse_qknorm_rope(graph, &i, &computed_nodes)) {
+                continue;
+            }
+
             htp_opnode node(HTP_OP_INVALID, n);
             node.opcode = op_remap_to_htp(n);
             if (node.opcode == HTP_OP_MUL_MAT || node.opcode == HTP_OP_MUL_MAT_ID) {
@@ -5479,11 +6362,21 @@ static ggml_status ggml_backend_hexagon_graph_compute(ggml_backend_t backend, gg
                     node.node->src[0], node.node->src[1], node.node,
                     (struct htp_mm_kernel_params *)node.kernel_params
                 );
+            } else if (node.opcode == HTP_OP_MUL_MAT_SEGMENTED) {
+                ggml_hexagon_precompute_segmented_matmul_params(sess,
+                    node.node->src[0], node.node->src[1], node.node->src[2], node.node,
+                    (struct htp_mm_kernel_params *) node.kernel_params
+                );
             } else if (node.opcode == HTP_OP_FLASH_ATTN_EXT) {
                 ggml_hexagon_precompute_flash_attn_params(sess,
                     node.node,
                     (struct htp_fa_kernel_params *)node.kernel_params
                 );
+            } else if (node.opcode == HTP_OP_CONV_2D) {
+                ggml_hexagon_precompute_conv2d_params(sess, node.node, (struct htp_conv2d_kernel_params *) node.kernel_params);
+            } else if (node.node->op == GGML_OP_GROUP_NORM_AFFINE_SILU) {
+                auto * params = (struct htp_group_norm_kernel_params *) node.kernel_params;
+                params->flags = HTP_GROUP_NORM_AFFINE | HTP_GROUP_NORM_SILU;
             } else if (htp_op_is_unary(node.opcode)) {
                 auto inputs = node.get_inputs();
                 const struct ggml_tensor * src0 = inputs[0];
@@ -6226,6 +7119,14 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
         return false;
     }
 
+    if (op->view_src) {
+        const ggml_tensor * producer = ggml_hexagon_unwrap_empty(op->view_src);
+        if (producer && producer != op && producer->op != GGML_OP_NONE &&
+            !ggml_backend_hexagon_device_supports_op(dev, producer)) {
+            return false;
+        }
+    }
+
     bool supp = false;
     switch (op->op) {
         case GGML_OP_NONE:
@@ -6247,6 +7148,14 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
             supp = ggml_hexagon_supported_mul_mat(sess, op);
             break;
 
+        case GGML_OP_MUL_MAT_SEGMENTED:
+            supp = ggml_hexagon_supported_mul_mat_segmented(sess, op);
+            break;
+
+        case GGML_OP_QKNORM_ROPE:
+            supp = ggml_hexagon_supported_qknorm_rope(op);
+            break;
+
         case GGML_OP_MUL_MAT_ID:
             supp = ggml_hexagon_supported_mul_mat_id(sess, op);
             break;
@@ -6352,6 +7261,20 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
             supp = ggml_hexagon_supported_im2col(sess, op);
             break;
 
+        case GGML_OP_CONV_2D:
+        case GGML_OP_CONV_2D_BIAS:
+        case GGML_OP_CONV_2D_UPSCALE:
+            supp = ggml_hexagon_supported_conv2d(sess, op);
+            break;
+
+        case GGML_OP_GROUP_NORM:
+            supp = ggml_hexagon_supported_group_norm(op);
+            break;
+
+        case GGML_OP_GROUP_NORM_AFFINE_SILU:
+            supp = ggml_hexagon_supported_group_norm_affine_silu(op);
+            break;
+
         case GGML_OP_GATED_DELTA_NET:
             supp = ggml_hexagon_supported_gated_delta_net(sess, op);
             break;
@@ -6765,6 +7688,8 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) {
                   "please update hexagon_type to match ggml_type");
     static_assert((unsigned int) HTP_TYPE_IQ4_NL == (unsigned int) GGML_TYPE_IQ4_NL,
                   "please update hexagon_type to match ggml_type");
+    static_assert((unsigned int) HTP_TYPE_F8_E4M3 == (unsigned int) GGML_TYPE_F8_E4M3,
+                  "please update hexagon_type to match ggml_type");
 
     const char * str_verbose  = getenv("GGML_HEXAGON_VERBOSE");
     const char * str_opbatch  = getenv("GGML_HEXAGON_OPBATCH");
diff --git a/ggml/src/ggml-hexagon/htp/CMakeLists.txt b/ggml/src/ggml-hexagon/htp/CMakeLists.txt
index 77f3ee39dd3c..c1f22682b0fc 100644
--- a/ggml/src/ggml-hexagon/htp/CMakeLists.txt
+++ b/ggml/src/ggml-hexagon/htp/CMakeLists.txt
@@ -44,6 +44,9 @@ add_library(${HTP_LIB} SHARED
     argsort-ops.c
     im2col-ops.c
     allreduce-ops.c
+    conv2d-ops.c
+    groupnorm-ops.c
+    qknorm-rope-ops.c
 )
 
 target_compile_definitions(${HTP_LIB} PRIVATE
diff --git a/ggml/src/ggml-hexagon/htp/conv2d-ops.c b/ggml/src/ggml-hexagon/htp/conv2d-ops.c
new file mode 100644
index 000000000000..3ee5eab4fc30
--- /dev/null
+++ b/ggml/src/ggml-hexagon/htp/conv2d-ops.c
@@ -0,0 +1,1246 @@
+#pragma clang diagnostic ignored "-Wunused-function"
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <HAP_farf.h>
+#include <hexagon_protos.h>
+#include <hexagon_types.h>
+#include <string.h>
+
+#define GGML_COMMON_DECL_C
+#include "ggml-common.h"
+#include "conv2d-ops.h"
+#include "hex-dma.h"
+#include "hex-profile.h"
+#include "matmul-ops.h"
+#include "hmx-mm-kernels-tiled.h"
+#include "hmx-utils.h"
+#include "htp-ctx.h"
+#include "htp-ops.h"
+#include "hvx-utils.h"
+#include "hvx-sigmoid.h"
+
+static const uint32_t conv2d_gather_rows_32x32[32] __attribute__((aligned(VLEN))) = {
+    0 * 128,  1 * 128,  2 * 128,  3 * 128,  4 * 128,  5 * 128,  6 * 128,  7 * 128,
+    8 * 128,  9 * 128, 10 * 128, 11 * 128, 12 * 128, 13 * 128, 14 * 128, 15 * 128,
+   16 * 128, 17 * 128, 18 * 128, 19 * 128, 20 * 128, 21 * 128, 22 * 128, 23 * 128,
+   24 * 128, 25 * 128, 26 * 128, 27 * 128, 28 * 128, 29 * 128, 30 * 128, 31 * 128,
+};
+
+static const uint16_t conv2d_scatter_cols_32x32[64] __attribute__((aligned(VLEN))) = {
+     0,    2,  128,  130,  256,  258,  384,  386,
+   512,  514,  640,  642,  768,  770,  896,  898,
+  1024, 1026, 1152, 1154, 1280, 1282, 1408, 1410,
+  1536, 1538, 1664, 1666, 1792, 1794, 1920, 1922,
+     0,    0,    0,    0,    0,    0,    0,    0,
+     0,    0,    0,    0,    0,    0,    0,    0,
+     0,    0,    0,    0,    0,    0,    0,    0,
+     0,    0,    0,    0,    0,    0,    0,    0,
+};
+
+#define CONV2D_PAIR_DECLS(prefix) \
+    HVX_VectorPair prefix##01, prefix##23, prefix##45, prefix##67; \
+    HVX_VectorPair prefix##89, prefix##1011, prefix##1213, prefix##1415; \
+    HVX_VectorPair prefix##1617, prefix##1819, prefix##2021, prefix##2223; \
+    HVX_VectorPair prefix##2425, prefix##2627, prefix##2829, prefix##3031
+
+#define CONV2D_TRANSPOSE_FINISH(STORE_PAIR) do { \
+    CONV2D_PAIR_DECLS(s); \
+    s01   = Q6_W_vshuff_VVR(Q6_V_lo_W(d23),   Q6_V_lo_W(d01),   -8); \
+    s23   = Q6_W_vshuff_VVR(Q6_V_hi_W(d23),   Q6_V_hi_W(d01),   -8); \
+    s45   = Q6_W_vshuff_VVR(Q6_V_lo_W(d67),   Q6_V_lo_W(d45),   -8); \
+    s67   = Q6_W_vshuff_VVR(Q6_V_hi_W(d67),   Q6_V_hi_W(d45),   -8); \
+    s89   = Q6_W_vshuff_VVR(Q6_V_lo_W(d1011), Q6_V_lo_W(d89),   -8); \
+    s1011 = Q6_W_vshuff_VVR(Q6_V_hi_W(d1011), Q6_V_hi_W(d89),   -8); \
+    s1213 = Q6_W_vshuff_VVR(Q6_V_lo_W(d1415), Q6_V_lo_W(d1213), -8); \
+    s1415 = Q6_W_vshuff_VVR(Q6_V_hi_W(d1415), Q6_V_hi_W(d1213), -8); \
+    s1617 = Q6_W_vshuff_VVR(Q6_V_lo_W(d1819), Q6_V_lo_W(d1617), -8); \
+    s1819 = Q6_W_vshuff_VVR(Q6_V_hi_W(d1819), Q6_V_hi_W(d1617), -8); \
+    s2021 = Q6_W_vshuff_VVR(Q6_V_lo_W(d2223), Q6_V_lo_W(d2021), -8); \
+    s2223 = Q6_W_vshuff_VVR(Q6_V_hi_W(d2223), Q6_V_hi_W(d2021), -8); \
+    s2425 = Q6_W_vshuff_VVR(Q6_V_lo_W(d2627), Q6_V_lo_W(d2425), -8); \
+    s2627 = Q6_W_vshuff_VVR(Q6_V_hi_W(d2627), Q6_V_hi_W(d2425), -8); \
+    s2829 = Q6_W_vshuff_VVR(Q6_V_lo_W(d3031), Q6_V_lo_W(d2829), -8); \
+    s3031 = Q6_W_vshuff_VVR(Q6_V_hi_W(d3031), Q6_V_hi_W(d2829), -8); \
+    d01   = Q6_W_vshuff_VVR(Q6_V_lo_W(s45),   Q6_V_lo_W(s01),   -16); \
+    d23   = Q6_W_vshuff_VVR(Q6_V_hi_W(s45),   Q6_V_hi_W(s01),   -16); \
+    d45   = Q6_W_vshuff_VVR(Q6_V_lo_W(s67),   Q6_V_lo_W(s23),   -16); \
+    d67   = Q6_W_vshuff_VVR(Q6_V_hi_W(s67),   Q6_V_hi_W(s23),   -16); \
+    d89   = Q6_W_vshuff_VVR(Q6_V_lo_W(s1213), Q6_V_lo_W(s89),   -16); \
+    d1011 = Q6_W_vshuff_VVR(Q6_V_hi_W(s1213), Q6_V_hi_W(s89),   -16); \
+    d1213 = Q6_W_vshuff_VVR(Q6_V_lo_W(s1415), Q6_V_lo_W(s1011), -16); \
+    d1415 = Q6_W_vshuff_VVR(Q6_V_hi_W(s1415), Q6_V_hi_W(s1011), -16); \
+    d1617 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2021), Q6_V_lo_W(s1617), -16); \
+    d1819 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2021), Q6_V_hi_W(s1617), -16); \
+    d2021 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2223), Q6_V_lo_W(s1819), -16); \
+    d2223 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2223), Q6_V_hi_W(s1819), -16); \
+    d2425 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2829), Q6_V_lo_W(s2425), -16); \
+    d2627 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2829), Q6_V_hi_W(s2425), -16); \
+    d2829 = Q6_W_vshuff_VVR(Q6_V_lo_W(s3031), Q6_V_lo_W(s2627), -16); \
+    d3031 = Q6_W_vshuff_VVR(Q6_V_hi_W(s3031), Q6_V_hi_W(s2627), -16); \
+    s01   = Q6_W_vshuff_VVR(Q6_V_lo_W(d89),   Q6_V_lo_W(d01),   -32); \
+    s23   = Q6_W_vshuff_VVR(Q6_V_hi_W(d89),   Q6_V_hi_W(d01),   -32); \
+    s45   = Q6_W_vshuff_VVR(Q6_V_lo_W(d1011), Q6_V_lo_W(d23),   -32); \
+    s67   = Q6_W_vshuff_VVR(Q6_V_hi_W(d1011), Q6_V_hi_W(d23),   -32); \
+    s89   = Q6_W_vshuff_VVR(Q6_V_lo_W(d1213), Q6_V_lo_W(d45),   -32); \
+    s1011 = Q6_W_vshuff_VVR(Q6_V_hi_W(d1213), Q6_V_hi_W(d45),   -32); \
+    s1213 = Q6_W_vshuff_VVR(Q6_V_lo_W(d1415), Q6_V_lo_W(d67),   -32); \
+    s1415 = Q6_W_vshuff_VVR(Q6_V_hi_W(d1415), Q6_V_hi_W(d67),   -32); \
+    s1617 = Q6_W_vshuff_VVR(Q6_V_lo_W(d2425), Q6_V_lo_W(d1617), -32); \
+    s1819 = Q6_W_vshuff_VVR(Q6_V_hi_W(d2425), Q6_V_hi_W(d1617), -32); \
+    s2021 = Q6_W_vshuff_VVR(Q6_V_lo_W(d2627), Q6_V_lo_W(d1819), -32); \
+    s2223 = Q6_W_vshuff_VVR(Q6_V_hi_W(d2627), Q6_V_hi_W(d1819), -32); \
+    s2425 = Q6_W_vshuff_VVR(Q6_V_lo_W(d2829), Q6_V_lo_W(d2021), -32); \
+    s2627 = Q6_W_vshuff_VVR(Q6_V_hi_W(d2829), Q6_V_hi_W(d2021), -32); \
+    s2829 = Q6_W_vshuff_VVR(Q6_V_lo_W(d3031), Q6_V_lo_W(d2223), -32); \
+    s3031 = Q6_W_vshuff_VVR(Q6_V_hi_W(d3031), Q6_V_hi_W(d2223), -32); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s1617), Q6_V_lo_W(s01), -64); STORE_PAIR( 0, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s1617), Q6_V_hi_W(s01), -64); STORE_PAIR( 1, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s1819), Q6_V_lo_W(s23), -64); STORE_PAIR( 2, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s1819), Q6_V_hi_W(s23), -64); STORE_PAIR( 3, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2021), Q6_V_lo_W(s45), -64); STORE_PAIR( 4, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2021), Q6_V_hi_W(s45), -64); STORE_PAIR( 5, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2223), Q6_V_lo_W(s67), -64); STORE_PAIR( 6, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2223), Q6_V_hi_W(s67), -64); STORE_PAIR( 7, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2425), Q6_V_lo_W(s89), -64); STORE_PAIR( 8, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2425), Q6_V_hi_W(s89), -64); STORE_PAIR( 9, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2627), Q6_V_lo_W(s1011), -64); STORE_PAIR(10, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2627), Q6_V_hi_W(s1011), -64); STORE_PAIR(11, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s2829), Q6_V_lo_W(s1213), -64); STORE_PAIR(12, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s2829), Q6_V_hi_W(s1213), -64); STORE_PAIR(13, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_lo_W(s3031), Q6_V_lo_W(s1415), -64); STORE_PAIR(14, d01); \
+    d01 = Q6_W_vshuff_VVR(Q6_V_hi_W(s3031), Q6_V_hi_W(s1415), -64); STORE_PAIR(15, d01); \
+} while (0)
+
+static __attribute__((noinline)) void conv2d_pack_tile_f32_to_f16(
+        __fp16 * dst,
+        const float * src,
+        size_t channel_stride) {
+    CONV2D_PAIR_DECLS(d);
+#define LOAD_CHANNEL(c) hvx_vmemu(src + (size_t) (c) * channel_stride)
+    d01   = Q6_W_vshuff_VVR(LOAD_CHANNEL( 1), LOAD_CHANNEL( 0), -4);
+    d23   = Q6_W_vshuff_VVR(LOAD_CHANNEL( 3), LOAD_CHANNEL( 2), -4);
+    d45   = Q6_W_vshuff_VVR(LOAD_CHANNEL( 5), LOAD_CHANNEL( 4), -4);
+    d67   = Q6_W_vshuff_VVR(LOAD_CHANNEL( 7), LOAD_CHANNEL( 6), -4);
+    d89   = Q6_W_vshuff_VVR(LOAD_CHANNEL( 9), LOAD_CHANNEL( 8), -4);
+    d1011 = Q6_W_vshuff_VVR(LOAD_CHANNEL(11), LOAD_CHANNEL(10), -4);
+    d1213 = Q6_W_vshuff_VVR(LOAD_CHANNEL(13), LOAD_CHANNEL(12), -4);
+    d1415 = Q6_W_vshuff_VVR(LOAD_CHANNEL(15), LOAD_CHANNEL(14), -4);
+    d1617 = Q6_W_vshuff_VVR(LOAD_CHANNEL(17), LOAD_CHANNEL(16), -4);
+    d1819 = Q6_W_vshuff_VVR(LOAD_CHANNEL(19), LOAD_CHANNEL(18), -4);
+    d2021 = Q6_W_vshuff_VVR(LOAD_CHANNEL(21), LOAD_CHANNEL(20), -4);
+    d2223 = Q6_W_vshuff_VVR(LOAD_CHANNEL(23), LOAD_CHANNEL(22), -4);
+    d2425 = Q6_W_vshuff_VVR(LOAD_CHANNEL(25), LOAD_CHANNEL(24), -4);
+    d2627 = Q6_W_vshuff_VVR(LOAD_CHANNEL(27), LOAD_CHANNEL(26), -4);
+    d2829 = Q6_W_vshuff_VVR(LOAD_CHANNEL(29), LOAD_CHANNEL(28), -4);
+    d3031 = Q6_W_vshuff_VVR(LOAD_CHANNEL(31), LOAD_CHANNEL(30), -4);
+#undef LOAD_CHANNEL
+#define PACK_STORE_PAIR(i, p) \
+    ((HVX_Vector *) dst)[i] = hvx_vec_f32_to_f16_shuff(Q6_V_lo_W(p), Q6_V_hi_W(p))
+    CONV2D_TRANSPOSE_FINISH(PACK_STORE_PAIR);
+#undef PACK_STORE_PAIR
+}
+
+static __attribute__((noinline)) void conv2d_store_tile_f16_to_f32(
+        float * dst,
+        const __fp16 * src,
+        size_t channel_stride) {
+    CONV2D_PAIR_DECLS(d);
+    const HVX_Vector * tiles = (const HVX_Vector *) src;
+    HVX_VectorPair rows;
+#define LOAD_ROW_PAIR(i, dpair) do { \
+    rows = hvx_vec_f16_to_f32_shuff(tiles[i]); \
+    dpair = Q6_W_vshuff_VVR(Q6_V_hi_W(rows), Q6_V_lo_W(rows), -4); \
+} while (0)
+    LOAD_ROW_PAIR( 0, d01);
+    LOAD_ROW_PAIR( 1, d23);
+    LOAD_ROW_PAIR( 2, d45);
+    LOAD_ROW_PAIR( 3, d67);
+    LOAD_ROW_PAIR( 4, d89);
+    LOAD_ROW_PAIR( 5, d1011);
+    LOAD_ROW_PAIR( 6, d1213);
+    LOAD_ROW_PAIR( 7, d1415);
+    LOAD_ROW_PAIR( 8, d1617);
+    LOAD_ROW_PAIR( 9, d1819);
+    LOAD_ROW_PAIR(10, d2021);
+    LOAD_ROW_PAIR(11, d2223);
+    LOAD_ROW_PAIR(12, d2425);
+    LOAD_ROW_PAIR(13, d2627);
+    LOAD_ROW_PAIR(14, d2829);
+    LOAD_ROW_PAIR(15, d3031);
+#undef LOAD_ROW_PAIR
+#define OUTPUT_STORE_PAIR(i, p) do { \
+    hvx_vmem(dst + (size_t) (2 * (i) + 0) * channel_stride) = Q6_V_lo_W(p); \
+    hvx_vmem(dst + (size_t) (2 * (i) + 1) * channel_stride) = Q6_V_hi_W(p); \
+} while (0)
+    CONV2D_TRANSPOSE_FINISH(OUTPUT_STORE_PAIR);
+#undef OUTPUT_STORE_PAIR
+}
+
+struct conv2d_prepare_state {
+    struct htp_context * ctx;
+    float * x;
+    const float * src;
+    __fp16 * a;
+    uint32_t iw;
+    uint32_t ih;
+    uint32_t src_w;
+    uint32_t src_h;
+    uint32_t halo_w;
+    uint32_t halo_h;
+    uint32_t sx_begin;
+    uint32_t sy_begin;
+    uint32_t copy_w;
+    uint32_t copy_h;
+    uint32_t dx;
+    uint32_t dy;
+    int32_t sx0;
+    int32_t sy0;
+    uint32_t kh;
+    uint32_t kw;
+    uint32_t ic_blocks;
+    uint32_t n_groups;
+    uint32_t group_channels;
+    uint32_t k_tiles;
+    uint32_t activation_k_tiles;
+    uint32_t tile_w;
+    uint32_t tile_h;
+    uint32_t x_tiles;
+    uint32_t m_tiles;
+    bool clear;
+    bool wait_weight_dma;
+    bool upscale2;
+    bool direct_upscale2;
+    bool input_silu;
+    bool input_f16;
+};
+
+static inline HVX_Vector conv2d_silu(HVX_Vector x) {
+    const HVX_Vector sigmoid = hvx_vec_fast_sigmoid_f32_guard(
+        x, hvx_vec_splat_f32(1.0f), hvx_vec_splat_f32(87.0f),
+        hvx_vec_splat_f32(-87.0f));
+    return Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(x, sigmoid));
+}
+
+static void conv2d_silu_inplace(float * data, uint32_t n) {
+    HVX_Vector * vectors = (HVX_Vector *) data;
+    const uint32_t nvec = n / 32u;
+    uint32_t i = 0;
+    for (; i + 3u < nvec; i += 4u) {
+        const HVX_Vector x0 = vectors[i + 0u];
+        const HVX_Vector x1 = vectors[i + 1u];
+        const HVX_Vector x2 = vectors[i + 2u];
+        const HVX_Vector x3 = vectors[i + 3u];
+        vectors[i + 0u] = conv2d_silu(x0);
+        vectors[i + 1u] = conv2d_silu(x1);
+        vectors[i + 2u] = conv2d_silu(x2);
+        vectors[i + 3u] = conv2d_silu(x3);
+    }
+    for (; i < nvec; ++i) {
+        vectors[i] = conv2d_silu(vectors[i]);
+    }
+    const uint32_t tail = n % 32u;
+    if (tail) {
+        hvx_vec_store_a(vectors + nvec, tail * sizeof(float),
+                        conv2d_silu(vectors[nvec]));
+    }
+}
+
+static void conv2d_silu_f16_inplace(__fp16 * data, uint32_t n) {
+    HVX_Vector * vectors = (HVX_Vector *) data;
+    const uint32_t nvec = n / 64u;
+    for (uint32_t i = 0; i < nvec; ++i) {
+        const HVX_Vector x = vectors[i];
+        vectors[i] = hvx_vec_mul_f16_f16(x, hvx_vec_fast_sigmoid_f16(x));
+    }
+    const uint32_t tail = n % 64u;
+    if (tail) {
+        const HVX_Vector x = vectors[nvec];
+        hvx_vec_store_a(vectors + nvec, tail * sizeof(__fp16),
+                        hvx_vec_mul_f16_f16(x, hvx_vec_fast_sigmoid_f16(x)));
+    }
+}
+
+static inline HVX_Vector conv2d_join_f16_halves(HVX_Vector lo, HVX_Vector hi) {
+    const HVX_VectorPred lower_half = Q6_Q_vsetq_R(64);
+    return Q6_V_vmux_QVV(lower_half, lo, Q6_V_vror_VR(hi, 64));
+}
+
+static inline HVX_Vector conv2d_pack_f16_pair(HVX_Vector row0, HVX_Vector row1) {
+    return Q6_Vh_vshuff_Vh(conv2d_join_f16_halves(row0, row1));
+}
+
+static inline void conv2d_upscale2_row(float * dst, const float * src,
+                                       uint32_t logical_x, uint32_t count) {
+    while (count > 0) {
+        const uint32_t take = count < 32u ? count : 32u;
+        const HVX_Vector input = hvx_vmemu(src + logical_x / 2u);
+        const HVX_VectorPair dup = Q6_W_vshuff_VVR(input, input, -4);
+        HVX_Vector output = Q6_V_lo_W(dup);
+        if (logical_x & 1u) {
+            output = Q6_V_valign_VVR(Q6_V_hi_W(dup), output, sizeof(float));
+        }
+        hvx_vec_store_u(dst, take * sizeof(float), output);
+        logical_x += take;
+        dst += take;
+        count -= take;
+    }
+}
+
+static inline void conv2d_upscale2_row_f16(__fp16 * dst, const __fp16 * src,
+                                           uint32_t logical_x, uint32_t count) {
+    while (count > 0) {
+        const uint32_t take = count < 64u ? count : 64u;
+        const HVX_Vector input = hvx_vmemu(src + logical_x / 2u);
+        const HVX_VectorPair dup = Q6_W_vshuff_VVR(input, input, -2);
+        HVX_Vector output = Q6_V_lo_W(dup);
+        if (logical_x & 1u) {
+            output = Q6_V_valign_VVR(Q6_V_hi_W(dup), output, sizeof(__fp16));
+        }
+        hvx_vec_store_u(dst, take * sizeof(__fp16), output);
+        logical_x += take;
+        dst += take;
+        count -= take;
+    }
+}
+
+static void conv2d_push_halo_group(
+        const struct conv2d_prepare_state * st,
+        dma_queue * q,
+        uint32_t group,
+        void * x_ptr) {
+    const size_t channel_elms = (size_t) st->halo_w * st->halo_h;
+    if (st->input_f16) {
+        __fp16 * x = (__fp16 *) x_ptr;
+        if (st->clear) {
+            memset(x, 0, st->group_channels * channel_elms * sizeof(__fp16));
+        }
+        const uint32_t channel_first = group * st->group_channels;
+        if (st->upscale2) {
+            for (uint32_t c = 0; c < st->group_channels; ++c) {
+                __fp16 * d = x + ((size_t) c * st->halo_h + st->dy) * st->halo_w + st->dx;
+                const __fp16 * channel = (const __fp16 *) st->src +
+                    (size_t) (channel_first + c) * st->src_h * st->src_w;
+                for (uint32_t y = 0; y < st->copy_h; ++y) {
+                    const uint32_t logical_y = st->sy_begin + y;
+                    const __fp16 * s = channel + (size_t) (logical_y / 2u) * st->src_w;
+                    conv2d_upscale2_row_f16(d + (size_t) y * st->halo_w, s,
+                                            st->sx_begin, st->copy_w);
+                }
+            }
+            return;
+        }
+        for (uint32_t c = 0; c < st->group_channels; ++c) {
+            __fp16 * d = x + ((size_t) c * st->halo_h + st->dy) * st->halo_w + st->dx;
+            const __fp16 * s = (const __fp16 *) st->src +
+                ((size_t) (channel_first + c) * st->ih + st->sy_begin) * st->iw +
+                st->sx_begin;
+            while (!dma_queue_push(q, dma_make_ptr(d, s),
+                                   st->halo_w * sizeof(__fp16), st->iw * sizeof(__fp16),
+                                   st->copy_w * sizeof(__fp16), st->copy_h)) {
+                dma_queue_flush(q);
+            }
+        }
+        return;
+    }
+    float * x = (float *) x_ptr;
+    if (st->clear) {
+        memset(x, 0, st->group_channels * channel_elms * sizeof(float));
+    }
+    const uint32_t channel_first = group * st->group_channels;
+    if (st->upscale2) {
+        for (uint32_t c = 0; c < st->group_channels; ++c) {
+            float * d = x + ((size_t) c * st->halo_h + st->dy) * st->halo_w + st->dx;
+            const float * channel = st->src +
+                (size_t) (channel_first + c) * st->src_h * st->src_w;
+            for (uint32_t y = 0; y < st->copy_h; ++y) {
+                const uint32_t logical_y = st->sy_begin + y;
+                const float * s = channel + (size_t) (logical_y / 2u) * st->src_w;
+                conv2d_upscale2_row(d + (size_t) y * st->halo_w, s,
+                                    st->sx_begin, st->copy_w);
+            }
+        }
+        return;
+    }
+    for (uint32_t c = 0; c < st->group_channels; ++c) {
+        float * d = x + ((size_t) c * st->halo_h + st->dy) * st->halo_w + st->dx;
+        const float * s = st->src + ((size_t) (channel_first + c) * st->ih + st->sy_begin) * st->iw + st->sx_begin;
+        while (!dma_queue_push(q, dma_make_ptr(d, s),
+                               st->halo_w * sizeof(float), st->iw * sizeof(float),
+                               st->copy_w * sizeof(float), st->copy_h)) {
+            dma_queue_flush(q);
+        }
+    }
+}
+
+static inline void conv2d_store_activation_triplet(
+        const struct conv2d_prepare_state * st,
+        uint32_t source_mt,
+        uint32_t icb,
+        uint32_t cp,
+        HVX_Vector p0,
+        HVX_Vector p1,
+    HVX_Vector p2) {
+    HVX_Vector * tile0 = (HVX_Vector *)
+        (st->a + ((size_t) source_mt * st->activation_k_tiles +
+         0 * st->ic_blocks + icb) * HTP_MM_HMX_TILE_N_ELMS);
+    HVX_Vector * tile1 = (HVX_Vector *)
+        (st->a + ((size_t) source_mt * st->activation_k_tiles +
+         1 * st->ic_blocks + icb) * HTP_MM_HMX_TILE_N_ELMS);
+    HVX_Vector * tile2 = (HVX_Vector *)
+        (st->a + ((size_t) source_mt * st->activation_k_tiles +
+         2 * st->ic_blocks + icb) * HTP_MM_HMX_TILE_N_ELMS);
+    tile0[cp] = p0;
+    tile1[cp] = p1;
+    tile2[cp] = p2;
+}
+
+static void conv2d_pack_activation_group_f16(
+        const struct conv2d_prepare_state * st,
+        uint32_t group,
+        __fp16 * x) {
+    const size_t channel_elms = (size_t) st->halo_w * st->halo_h;
+    const uint32_t channel_first = group * st->group_channels;
+    const uint32_t icb = channel_first / 32u;
+    const uint32_t cp_first = (channel_first % 32u) / 2u;
+    const uint32_t n_pairs = st->group_channels / 2u;
+    if (st->input_silu) {
+        conv2d_silu_f16_inplace(x, st->group_channels * channel_elms);
+    }
+    if (st->kw == 3) {
+        for (uint32_t cp = 0; cp < n_pairs; ++cp) {
+            for (uint32_t sy = 0; sy < st->halo_h; ++sy) {
+                const __fp16 * row = x + (size_t) sy * st->halo_w;
+                const __fp16 * row0 = row + (2u * cp + 0u) * channel_elms;
+                const __fp16 * row1 = row + (2u * cp + 1u) * channel_elms;
+                for (uint32_t tx = 0; tx < st->x_tiles; ++tx) {
+                    const uint32_t x0 = tx * HTP_CONV2D_TILE_W;
+                    conv2d_store_activation_triplet(
+                        st, sy * st->x_tiles + tx, icb, cp_first + cp,
+                        conv2d_pack_f16_pair(hvx_vmemu(row0 + x0 + 0u),
+                                             hvx_vmemu(row1 + x0 + 0u)),
+                        conv2d_pack_f16_pair(hvx_vmemu(row0 + x0 + 1u),
+                                             hvx_vmemu(row1 + x0 + 1u)),
+                        conv2d_pack_f16_pair(hvx_vmemu(row0 + x0 + 2u),
+                                             hvx_vmemu(row1 + x0 + 2u)));
+                }
+            }
+        }
+        return;
+    }
+    for (uint32_t oy = 0; oy < st->tile_h; ++oy) {
+        for (uint32_t tx = 0; tx < st->x_tiles; ++tx) {
+            const uint32_t mt = oy * st->x_tiles + tx;
+            HVX_Vector * tile = (HVX_Vector *)
+                (st->a + ((size_t) mt * st->k_tiles + icb) *
+                 HTP_MM_HMX_TILE_N_ELMS);
+            const __fp16 * src = x + (size_t) oy * st->halo_w +
+                                  tx * HTP_CONV2D_TILE_W;
+            for (uint32_t cp = 0; cp < n_pairs; ++cp) {
+                tile[cp_first + cp] = conv2d_pack_f16_pair(
+                    hvx_vmemu(src + (2u * cp + 0u) * channel_elms),
+                    hvx_vmemu(src + (2u * cp + 1u) * channel_elms));
+            }
+        }
+    }
+}
+
+static inline HVX_Vector conv2d_upscale2_shift(
+        HVX_Vector lo,
+        HVX_Vector hi,
+        uint32_t shift_words) {
+    return shift_words == 0 ? lo
+                            : Q6_V_valign_VVR(hi, lo, shift_words * sizeof(float));
+}
+
+static void conv2d_prepare_upscale_group_direct(
+        const struct conv2d_prepare_state * st,
+        uint32_t group) {
+    const uint32_t channel_first = group * st->group_channels;
+    const uint32_t icb = channel_first / 32;
+    const uint32_t cp_first = (channel_first % 32) / 2;
+    const uint32_t n_pairs = st->group_channels / 2;
+    const uint32_t x_parity = (uint32_t) st->sx0 & 1u;
+
+    for (uint32_t cp = 0; cp < n_pairs; ++cp) {
+        const uint32_t c0 = channel_first + 2 * cp;
+        const uint32_t c1 = c0 + 1;
+        for (uint32_t sy = 0; sy < st->halo_h; ++sy) {
+            const uint32_t iy = (uint32_t) (st->sy0 + (int32_t) sy) / 2u;
+            const float * row0 = st->src + ((size_t) c0 * st->src_h + iy) * st->src_w;
+            const float * row1 = st->src + ((size_t) c1 * st->src_h + iy) * st->src_w;
+            for (uint32_t tx = 0; tx < st->x_tiles; ++tx) {
+                const uint32_t logical_x = (uint32_t) st->sx0 + tx * HTP_CONV2D_TILE_W;
+                const uint32_t ix = logical_x / 2u;
+                const HVX_Vector in0 = hvx_vmemu(row0 + ix);
+                const HVX_Vector in1 = hvx_vmemu(row1 + ix);
+                const HVX_VectorPair dup0 = Q6_W_vshuff_VVR(in0, in0, -4);
+                const HVX_VectorPair dup1 = Q6_W_vshuff_VVR(in1, in1, -4);
+                HVX_Vector x00 = conv2d_upscale2_shift(
+                    Q6_V_lo_W(dup0), Q6_V_hi_W(dup0), x_parity + 0u);
+                HVX_Vector x01 = conv2d_upscale2_shift(
+                    Q6_V_lo_W(dup1), Q6_V_hi_W(dup1), x_parity + 0u);
+                HVX_Vector x10 = conv2d_upscale2_shift(
+                    Q6_V_lo_W(dup0), Q6_V_hi_W(dup0), x_parity + 1u);
+                HVX_Vector x11 = conv2d_upscale2_shift(
+                    Q6_V_lo_W(dup1), Q6_V_hi_W(dup1), x_parity + 1u);
+                HVX_Vector x20 = conv2d_upscale2_shift(
+                    Q6_V_lo_W(dup0), Q6_V_hi_W(dup0), x_parity + 2u);
+                HVX_Vector x21 = conv2d_upscale2_shift(
+                    Q6_V_lo_W(dup1), Q6_V_hi_W(dup1), x_parity + 2u);
+                if (st->input_silu) {
+                    x00 = conv2d_silu(x00);
+                    x01 = conv2d_silu(x01);
+                    x10 = conv2d_silu(x10);
+                    x11 = conv2d_silu(x11);
+                    x20 = conv2d_silu(x20);
+                    x21 = conv2d_silu(x21);
+                }
+                conv2d_store_activation_triplet(
+                    st, sy * st->x_tiles + tx, icb, cp_first + cp,
+                    hvx_vec_f32_to_f16_shuff(x00, x01),
+                    hvx_vec_f32_to_f16_shuff(x10, x11),
+                    hvx_vec_f32_to_f16_shuff(x20, x21));
+            }
+        }
+    }
+}
+
+static inline HVX_Vector conv2d_upscale2_shift_f16(
+        HVX_Vector lo,
+        HVX_Vector hi,
+        uint32_t shift) {
+    return shift == 0 ? lo : Q6_V_valign_VVR(hi, lo, shift * sizeof(__fp16));
+}
+
+static void conv2d_prepare_upscale_group_direct_f16(
+        const struct conv2d_prepare_state * st,
+        uint32_t group) {
+    const uint32_t channel_first = group * st->group_channels;
+    const uint32_t icb = channel_first / 32u;
+    const uint32_t cp_first = (channel_first % 32u) / 2u;
+    const uint32_t n_pairs = st->group_channels / 2u;
+    const uint32_t x_parity = (uint32_t) st->sx0 & 1u;
+    const __fp16 * src = (const __fp16 *) st->src;
+
+    for (uint32_t cp = 0; cp < n_pairs; ++cp) {
+        const uint32_t c0 = channel_first + 2u * cp;
+        const uint32_t c1 = c0 + 1u;
+        for (uint32_t sy = 0; sy < st->halo_h; ++sy) {
+            const uint32_t iy = (uint32_t) (st->sy0 + (int32_t) sy) / 2u;
+            const __fp16 * row0 = src + ((size_t) c0 * st->src_h + iy) * st->src_w;
+            const __fp16 * row1 = src + ((size_t) c1 * st->src_h + iy) * st->src_w;
+            for (uint32_t tx = 0; tx < st->x_tiles; ++tx) {
+                const uint32_t logical_x = (uint32_t) st->sx0 + tx * HTP_CONV2D_TILE_W;
+                const uint32_t ix = logical_x / 2u;
+                const HVX_Vector in0 = hvx_vmemu(row0 + ix);
+                const HVX_Vector in1 = hvx_vmemu(row1 + ix);
+                const HVX_VectorPair dup0 = Q6_W_vshuff_VVR(in0, in0, -2);
+                const HVX_VectorPair dup1 = Q6_W_vshuff_VVR(in1, in1, -2);
+                HVX_Vector x00 = conv2d_upscale2_shift_f16(
+                    Q6_V_lo_W(dup0), Q6_V_hi_W(dup0), x_parity + 0u);
+                HVX_Vector x01 = conv2d_upscale2_shift_f16(
+                    Q6_V_lo_W(dup1), Q6_V_hi_W(dup1), x_parity + 0u);
+                HVX_Vector x10 = conv2d_upscale2_shift_f16(
+                    Q6_V_lo_W(dup0), Q6_V_hi_W(dup0), x_parity + 1u);
+                HVX_Vector x11 = conv2d_upscale2_shift_f16(
+                    Q6_V_lo_W(dup1), Q6_V_hi_W(dup1), x_parity + 1u);
+                HVX_Vector x20 = conv2d_upscale2_shift_f16(
+                    Q6_V_lo_W(dup0), Q6_V_hi_W(dup0), x_parity + 2u);
+                HVX_Vector x21 = conv2d_upscale2_shift_f16(
+                    Q6_V_lo_W(dup1), Q6_V_hi_W(dup1), x_parity + 2u);
+                if (st->input_silu) {
+                    x00 = hvx_vec_mul_f16_f16(x00, hvx_vec_fast_sigmoid_f16(x00));
+                    x01 = hvx_vec_mul_f16_f16(x01, hvx_vec_fast_sigmoid_f16(x01));
+                    x10 = hvx_vec_mul_f16_f16(x10, hvx_vec_fast_sigmoid_f16(x10));
+                    x11 = hvx_vec_mul_f16_f16(x11, hvx_vec_fast_sigmoid_f16(x11));
+                    x20 = hvx_vec_mul_f16_f16(x20, hvx_vec_fast_sigmoid_f16(x20));
+                    x21 = hvx_vec_mul_f16_f16(x21, hvx_vec_fast_sigmoid_f16(x21));
+                }
+                conv2d_store_activation_triplet(
+                    st, sy * st->x_tiles + tx, icb, cp_first + cp,
+                    conv2d_pack_f16_pair(x00, x01),
+                    conv2d_pack_f16_pair(x10, x11),
+                    conv2d_pack_f16_pair(x20, x21));
+            }
+        }
+    }
+}
+
+static void conv2d_prepare_activation_worker(unsigned int nth, unsigned int ith, void * data) {
+    struct conv2d_prepare_state * st = (struct conv2d_prepare_state *) data;
+    const uint32_t group_first = fastdiv(st->n_groups * ith, &st->ctx->n_threads_div);
+    const uint32_t group_last = fastdiv(st->n_groups * (ith + 1), &st->ctx->n_threads_div);
+    const size_t channel_elms = (size_t) st->halo_w * st->halo_h;
+    const uint32_t element_size = st->input_f16 ? sizeof(__fp16) : sizeof(float);
+    const uint32_t x_slot_first = 2 * ith;
+    uint8_t * x_slots = (uint8_t *) st->x +
+        (size_t) x_slot_first * st->group_channels * channel_elms * element_size;
+    dma_queue * q = st->ctx->dma[ith];
+    if (ith == 0 && st->wait_weight_dma) {
+        dma_queue_pop(q);
+    }
+    if (group_first == group_last) {
+        return;
+    }
+    if (st->direct_upscale2) {
+        for (uint32_t group = group_first; group < group_last; ++group) {
+            if (st->input_f16) {
+                conv2d_prepare_upscale_group_direct_f16(st, group);
+            } else {
+                conv2d_prepare_upscale_group_direct(st, group);
+            }
+        }
+        return;
+    }
+
+    if (st->copy_w != 0 && st->copy_h != 0) {
+        conv2d_push_halo_group(st, q, group_first, x_slots);
+        if (group_first + 1 < group_last) {
+            conv2d_push_halo_group(st, q, group_first + 1,
+                                   x_slots + st->group_channels * channel_elms * element_size);
+        }
+    } else if (st->clear) {
+        memset(x_slots, 0, (group_last - group_first < 2 ? group_last - group_first : 2) *
+                            st->group_channels * channel_elms * element_size);
+    }
+
+    for (uint32_t group = group_first; group < group_last; ++group) {
+        if (!st->upscale2 && st->copy_w != 0 && st->copy_h != 0) {
+            for (uint32_t c = 0; c < st->group_channels; ++c) {
+                dma_queue_pop(q);
+            }
+        }
+
+        uint8_t * x_ptr = x_slots + (size_t) ((group - group_first) & 1u) *
+                                   st->group_channels * channel_elms * element_size;
+        if (st->input_f16) {
+            conv2d_pack_activation_group_f16(st, group, (__fp16 *) x_ptr);
+        } else {
+            const uint32_t channel_first = group * st->group_channels;
+            const uint32_t icb = channel_first / 32;
+            const uint32_t cp_first = (channel_first % 32) / 2;
+            const uint32_t n_pairs = st->group_channels / 2;
+            float * x = (float *) x_ptr;
+            if (st->input_silu) {
+                conv2d_silu_inplace(x, st->group_channels * channel_elms);
+            }
+            if (st->kw == 3) {
+            for (uint32_t cp = 0; cp < n_pairs; ++cp) {
+                for (uint32_t sy = 0; sy < st->halo_h; ++sy) {
+                    const float * row = x + (size_t) sy * st->halo_w;
+                        const float * row0 = row + (2 * cp + 0) * channel_elms;
+                        const float * row1 = row + (2 * cp + 1) * channel_elms;
+                        HVX_Vector lo0 = hvx_vmemu(row0);
+                        HVX_Vector lo1 = hvx_vmemu(row1);
+                        for (uint32_t tx = 0; tx < st->x_tiles; ++tx) {
+                            HVX_Vector hi0 = hvx_vmemu(
+                                row0 + (tx + 1) * HTP_CONV2D_TILE_W);
+                            HVX_Vector hi1 = hvx_vmemu(
+                                row1 + (tx + 1) * HTP_CONV2D_TILE_W);
+                            const HVX_Vector p0 = hvx_vec_f32_to_f16_shuff(lo0, lo1);
+                            const HVX_Vector p1 = hvx_vec_f32_to_f16_shuff(
+                                Q6_V_valign_VVR(hi0, lo0, sizeof(float)),
+                                Q6_V_valign_VVR(hi1, lo1, sizeof(float)));
+                            const HVX_Vector p2 = hvx_vec_f32_to_f16_shuff(
+                                Q6_V_valign_VVR(hi0, lo0, 2 * sizeof(float)),
+                                Q6_V_valign_VVR(hi1, lo1, 2 * sizeof(float)));
+                            conv2d_store_activation_triplet(
+                                st, sy * st->x_tiles + tx, icb,
+                                cp_first + cp, p0, p1, p2);
+                            lo0 = hi0;
+                            lo1 = hi1;
+                        }
+                    }
+            }
+            } else {
+            for (uint32_t oy = 0; oy < st->tile_h; ++oy) {
+                for (uint32_t tx = 0; tx < st->x_tiles; ++tx) {
+                    const uint32_t mt = oy * st->x_tiles + tx;
+                    HVX_Vector * tile = (HVX_Vector *)
+                        (st->a + ((size_t) mt * st->k_tiles + icb) *
+                         HTP_MM_HMX_TILE_N_ELMS);
+                    const float * src = x + (size_t) oy * st->halo_w +
+                                        tx * HTP_CONV2D_TILE_W;
+                    for (uint32_t cp = 0; cp < n_pairs; ++cp) {
+                        HVX_Vector row0 =
+                            hvx_vmemu(src + (2 * cp + 0) * channel_elms);
+                        HVX_Vector row1 =
+                            hvx_vmemu(src + (2 * cp + 1) * channel_elms);
+                        tile[cp_first + cp] =
+                            hvx_vec_f32_to_f16_shuff(row0, row1);
+                    }
+                }
+            }
+            }
+        }
+        if (st->copy_w != 0 && st->copy_h != 0 && group + 2 < group_last) {
+            conv2d_push_halo_group(st, q, group + 2, x_ptr);
+        }
+    }
+}
+
+struct conv2d_store_state {
+    struct htp_context * ctx;
+    const __fp16 * c;
+    float * dst;
+    const float * bias;
+    const float * residual;
+    uint32_t ow;
+    uint32_t oh;
+    uint32_t oc;
+    uint32_t n_tiles;
+    uint32_t x0;
+    uint32_t y0;
+    uint32_t tile_w;
+    uint32_t valid_w;
+    uint32_t tile_h;
+    uint32_t x_tiles;
+    uint32_t m_tiles;
+    bool final_block;
+    atomic_uint * next_pair;
+    uint32_t flags;
+    bool output_f16;
+};
+
+static void conv2d_store_output_f16_worker(unsigned int nth, unsigned int ith, void * data) {
+    struct conv2d_store_state * st = (struct conv2d_store_state *) data;
+    const uint32_t n_pairs = (st->oc + 1) / 2;
+    const HVX_Vector zero = Q6_V_vzero();
+    const HVX_VectorPred all = Q6_Q_vcmp_eq_VbVb(zero, zero);
+    const uint32_t pair_chunk = 1;
+    for (;;) {
+        const uint32_t pair_first = atomic_fetch_add_explicit(
+            st->next_pair, pair_chunk, memory_order_relaxed);
+        if (pair_first >= n_pairs) {
+            break;
+        }
+        const uint32_t pair_last = hex_smin(pair_first + pair_chunk, n_pairs);
+        for (uint32_t pair = pair_first; pair < pair_last; ++pair) {
+            const uint32_t channel = pair * 2u;
+            const uint32_t nt = channel / 32u;
+            const uint32_t rp = (channel % 32u) / 2u;
+            const HVX_Vector bias0 = hvx_vec_splat_f16(
+                (st->flags & HTP_CONV2D_BIAS) ? st->bias[channel] : 0.0f);
+            const HVX_Vector bias1 = hvx_vec_splat_f16(
+                (st->flags & HTP_CONV2D_BIAS) && channel + 1u < st->oc
+                    ? st->bias[channel + 1u] : 0.0f);
+            const HVX_Vector packed_bias = Q6_Vh_vshuff_Vh(
+                conv2d_join_f16_halves(bias0, bias1));
+            for (uint32_t oy = 0; oy < st->tile_h; ++oy) {
+                for (uint32_t tx = 0; tx < st->x_tiles; tx += 2u) {
+                    HVX_Vector planar[2];
+                    const uint32_t count = hex_smin(st->x_tiles - tx, 2u);
+                    const uint32_t valid_count = hex_smin(
+                        st->valid_w - tx * HTP_CONV2D_TILE_W,
+                        count * HTP_CONV2D_TILE_W);
+                    for (uint32_t j = 0; j < count; ++j) {
+                        const uint32_t mt = oy * st->x_tiles + tx + j;
+                        const HVX_Vector * tile = (const HVX_Vector *)
+                            (st->c + ((size_t) nt * st->m_tiles + mt) *
+                             HTP_MM_HMX_TILE_N_ELMS);
+                        HVX_Vector output = hvx_vec_add_f16_f16(tile[rp], packed_bias);
+                        __fp16 * dst0 = (__fp16 *) st->dst +
+                            ((size_t) channel * st->oh + st->y0 + oy) * st->ow +
+                            st->x0 + (tx + j) * HTP_CONV2D_TILE_W;
+                        if (st->flags & HTP_CONV2D_RESIDUAL) {
+                            const __fp16 * residual0 = (const __fp16 *) st->residual +
+                                (dst0 - (__fp16 *) st->dst);
+                            const __fp16 * residual1 = residual0 + (size_t) st->oh * st->ow;
+                            const HVX_Vector packed_residual = Q6_Vh_vshuff_Vh(
+                                conv2d_join_f16_halves(
+                                    hvx_vmemu(residual0), hvx_vmemu(residual1)));
+                            output = hvx_vec_add_f16_f16(output, packed_residual);
+                        }
+                        planar[j] = Q6_Vh_vdeal_Vh(output);
+                    }
+                    __fp16 * dst0 = (__fp16 *) st->dst +
+                        ((size_t) channel * st->oh + st->y0 + oy) * st->ow +
+                        st->x0 + tx * HTP_CONV2D_TILE_W;
+                    const HVX_Vector row0 = count == 2u
+                        ? conv2d_join_f16_halves(planar[0], planar[1]) : planar[0];
+                    if (valid_count == 2u * HTP_CONV2D_TILE_W &&
+                        ((uintptr_t) dst0 & 127u) == 0) {
+                        Q6_vmem_QRIV_nt(all, (HVX_Vector *) dst0, row0);
+                    } else {
+                        hvx_vec_store_u(dst0, valid_count * sizeof(__fp16), row0);
+                    }
+                    if (channel + 1u < st->oc) {
+                        __fp16 * dst1 = dst0 + (size_t) st->oh * st->ow;
+                        const HVX_Vector p0 = Q6_V_vror_VR(planar[0], 64);
+                        const HVX_Vector p1 = count == 2u
+                            ? Q6_V_vror_VR(planar[1], 64) : p0;
+                        const HVX_Vector row1 = count == 2u
+                            ? conv2d_join_f16_halves(p0, p1) : p0;
+                        if (valid_count == 2u * HTP_CONV2D_TILE_W &&
+                            ((uintptr_t) dst1 & 127u) == 0) {
+                            Q6_vmem_QRIV_nt(all, (HVX_Vector *) dst1, row1);
+                        } else {
+                            hvx_vec_store_u(dst1, valid_count * sizeof(__fp16), row1);
+                        }
+                    }
+                }
+            }
+        }
+    }
+    if (st->final_block) {
+        asm volatile("syncht" ::: "memory");
+    }
+}
+
+static void conv2d_store_output_worker(unsigned int nth, unsigned int ith, void * data) {
+    struct conv2d_store_state * st = (struct conv2d_store_state *) data;
+    if (st->output_f16) {
+        conv2d_store_output_f16_worker(nth, ith, data);
+        return;
+    }
+    const uint32_t n_pairs = (st->oc + 1) / 2;
+    const HVX_Vector zero = Q6_V_vzero();
+    const HVX_VectorPred all = Q6_Q_vcmp_eq_VbVb(zero, zero);
+
+    const uint32_t pair_chunk = 4;
+    for (;;) {
+        const uint32_t pair_first = atomic_fetch_add_explicit(
+            st->next_pair, pair_chunk, memory_order_relaxed);
+        if (pair_first >= n_pairs) {
+            break;
+        }
+        const uint32_t pair_last = hex_smin(pair_first + pair_chunk, n_pairs);
+        for (uint32_t pair = pair_first; pair < pair_last; ++pair) {
+            const uint32_t channel = pair * 2;
+            const uint32_t nt = channel / 32;
+            const uint32_t rp = (channel % 32) / 2;
+            const HVX_Vector bias0 = hvx_vec_splat_f32(
+                (st->flags & HTP_CONV2D_BIAS) ? st->bias[channel] : 0.0f);
+            const HVX_Vector bias1 = hvx_vec_splat_f32(
+                (st->flags & HTP_CONV2D_BIAS) && channel + 1 < st->oc
+                    ? st->bias[channel + 1] : 0.0f);
+            const size_t plane = (size_t) st->oh * st->ow;
+            const float * residual0 = (st->flags & HTP_CONV2D_RESIDUAL)
+                ? st->residual + (size_t) channel * plane +
+                    (size_t) st->y0 * st->ow + st->x0
+                : NULL;
+            const float * residual1 = residual0 && channel + 1 < st->oc
+                ? residual0 + plane : NULL;
+            if (residual0) {
+                const uint32_t rows = hex_smin(st->tile_h, 4u);
+                hex_l2fetch(residual0, st->valid_w * sizeof(float),
+                            st->ow * sizeof(float), rows);
+                if (residual1) {
+                    hex_l2fetch(residual1, st->valid_w * sizeof(float),
+                                st->ow * sizeof(float), rows);
+                }
+            }
+            for (uint32_t oy = 0; oy < st->tile_h; ++oy) {
+                if (residual0 && (oy & 3u) == 0 && oy + 4u < st->tile_h) {
+                    const uint32_t rows = hex_smin(st->tile_h - oy - 4u, 4u);
+                    hex_l2fetch(residual0 + (size_t) (oy + 4u) * st->ow,
+                                st->valid_w * sizeof(float), st->ow * sizeof(float), rows);
+                    if (residual1) {
+                        hex_l2fetch(residual1 + (size_t) (oy + 4u) * st->ow,
+                                    st->valid_w * sizeof(float), st->ow * sizeof(float), rows);
+                    }
+                }
+                for (uint32_t tx = 0; tx < st->x_tiles; ++tx) {
+                    const uint32_t valid_count = hex_smin(
+                        st->valid_w - tx * HTP_CONV2D_TILE_W,
+                        HTP_CONV2D_TILE_W);
+                    const uint32_t mt = oy * st->x_tiles + tx;
+                    const HVX_Vector * tile = (const HVX_Vector *)
+                        (st->c + ((size_t) nt * st->m_tiles + mt) *
+                         HTP_MM_HMX_TILE_N_ELMS);
+                    const HVX_VectorPair rows = hvx_vec_f16_to_f32_shuff(tile[rp]);
+                    float * dst = st->dst + ((size_t) channel * st->oh + st->y0 + oy) *
+                                  st->ow + st->x0 + tx * HTP_CONV2D_TILE_W;
+                    HVX_Vector row0 = hvx_vec_add_f32_f32(Q6_V_lo_W(rows), bias0);
+                    if (st->flags & HTP_CONV2D_RESIDUAL) {
+                        row0 = hvx_vec_add_f32_f32(
+                            row0, hvx_vmemu(st->residual + (dst - st->dst)));
+                    }
+                    if (valid_count == HTP_CONV2D_TILE_W) {
+                        Q6_vmem_QRIV_nt(all, (HVX_Vector *) dst, row0);
+                    } else {
+                        hvx_vec_store_u(dst, valid_count * sizeof(float), row0);
+                    }
+                    if (channel + 1 < st->oc) {
+                        float * dst1 = dst + (size_t) st->oh * st->ow;
+                        HVX_Vector row1 = hvx_vec_add_f32_f32(Q6_V_hi_W(rows), bias1);
+                        if (st->flags & HTP_CONV2D_RESIDUAL) {
+                            row1 = hvx_vec_add_f32_f32(
+                                row1, hvx_vmemu(st->residual + (dst1 - st->dst)));
+                        }
+                        if (valid_count == HTP_CONV2D_TILE_W) {
+                            Q6_vmem_QRIV_nt(all, (HVX_Vector *) dst1, row1);
+                        } else {
+                            hvx_vec_store_u(dst1, valid_count * sizeof(float), row1);
+                        }
+                    }
+                }
+            }
+        }
+    }
+    if (st->final_block) {
+        asm volatile("syncht" ::: "memory");
+    }
+}
+
+struct conv2d_hmx_job {
+    __fp16 * c;
+    const __fp16 * a;
+    const __fp16 * b;
+    const __fp16 * scales;
+    uint32_t mt;
+    uint32_t nt;
+    uint32_t kt;
+    uint32_t kh;
+    uint32_t x_tiles;
+    uint32_t tile_h;
+    uint32_t ic_blocks;
+};
+
+static inline void conv2d_hmx_mpy_tiles(
+        const __fp16 ** row,
+        const __fp16 ** col,
+        uint32_t n_tiles) {
+    while (n_tiles >= 32) {
+        asm volatile(HMX_LOAD_MPY_DEEP_F16("%1", "%2", "%0") :
+                     : "r"(65535), "r"(*row), "r"(*col));
+        *row += 32 * HTP_MM_HMX_TILE_N_ELMS;
+        *col += 32 * HTP_MM_HMX_TILE_N_ELMS;
+        n_tiles -= 32;
+    }
+    if (n_tiles > 0) {
+        const uint32_t range = 2048u * n_tiles - 1u;
+        asm volatile(HMX_LOAD_MPY_DEEP_F16("%1", "%2", "%0") :
+                     : "r"(range), "r"(*row), "r"(*col));
+    }
+}
+
+static void conv2d_hmx_3x3(const struct conv2d_hmx_job * job) {
+    asm volatile(HMX_SET_BIAS("%0") :: "r"((unsigned int) job->scales));
+
+    const uint32_t activation_k_tiles = 3 * job->ic_blocks;
+    const size_t activation_stride =
+        (size_t) activation_k_tiles * HTP_MM_HMX_TILE_N_ELMS;
+    const size_t weight_row_stride =
+        (size_t) 3 * activation_stride;
+    const uint32_t m_tiles = job->x_tiles * job->tile_h;
+
+    for (uint32_t r = 0; r < job->mt; ++r) {
+        const __fp16 * weight_row = job->a + r * weight_row_stride;
+        for (uint32_t oy = 0; oy < job->tile_h; ++oy) {
+            for (uint32_t tx = 0; tx < job->x_tiles; ++tx) {
+                asm volatile(HMX_CLRACC_F16());
+                for (uint32_t ky = 0; ky < 3; ++ky) {
+                    const __fp16 * row = weight_row + ky * activation_stride;
+                    const __fp16 * col = job->b +
+                        ((size_t) (oy + ky) * job->x_tiles + tx) * activation_stride;
+                    conv2d_hmx_mpy_tiles(&row, &col, activation_k_tiles);
+                }
+                __fp16 * out = job->c +
+                    ((size_t) r * m_tiles + oy * job->x_tiles + tx) *
+                    HTP_MM_HMX_TILE_N_ELMS;
+                asm volatile(HMX_STORE_AFTER_F16("%0", "%1") :
+                             : "r"(out), "r"(0) : "memory");
+            }
+        }
+    }
+}
+
+static void conv2d_hmx_worker(void * data) {
+    struct conv2d_hmx_job * job = (struct conv2d_hmx_job *) data;
+    if (job->kh == 3) {
+        conv2d_hmx_3x3(job);
+    } else {
+        core_dot_chunk_fp16(job->c, job->a, job->b, job->scales,
+                            job->mt, job->nt, job->kt);
+    }
+}
+
+static void conv2d_init_prepare_state(
+        struct htp_ops_context * octx,
+        float * x,
+        __fp16 * a,
+        uint32_t x0,
+        uint32_t y0,
+        uint32_t tile_w,
+        uint32_t tile_h,
+        bool wait_weight_dma,
+        const struct htp_conv2d_kernel_params * kp,
+        struct conv2d_prepare_state * state) {
+    const struct htp_tensor * src = octx->src[1];
+    const int32_t p0 = octx->op_params[2];
+    const int32_t p1 = octx->op_params[3];
+    const bool upscale2 = (kp->flags & HTP_CONV2D_UPSCALE2) != 0;
+    const uint32_t iw = upscale2 ? octx->dst->ne[0] : src->ne[0];
+    const uint32_t ih = upscale2 ? octx->dst->ne[1] : src->ne[1];
+    const uint32_t halo_w = tile_w + octx->src[0]->ne[0] - 1;
+    const uint32_t halo_h = tile_h + octx->src[0]->ne[1] - 1;
+    const int32_t sx0 = (int32_t) x0 - p0;
+    const int32_t sy0 = (int32_t) y0 - p1;
+    const int32_t sx_begin = sx0 < 0 ? 0 : sx0;
+    const int32_t sy_begin = sy0 < 0 ? 0 : sy0;
+    const int32_t sx_end = hex_smin((int32_t) iw, sx0 + (int32_t) halo_w);
+    const int32_t sy_end = hex_smin((int32_t) ih, sy0 + (int32_t) halo_h);
+
+    *state = (struct conv2d_prepare_state) {
+        .ctx = octx->ctx,
+        .x = x,
+        .src = (const float *) src->data,
+        .a = a,
+        .iw = iw,
+        .ih = ih,
+        .src_w = src->ne[0],
+        .src_h = src->ne[1],
+        .halo_w = halo_w,
+        .halo_h = halo_h,
+        .sx_begin = sx_begin,
+        .sy_begin = sy_begin,
+        .copy_w = sx_end > sx_begin ? (uint32_t) (sx_end - sx_begin) : 0,
+        .copy_h = sy_end > sy_begin ? (uint32_t) (sy_end - sy_begin) : 0,
+        .dx = (uint32_t) (sx_begin - sx0),
+        .dy = (uint32_t) (sy_begin - sy0),
+        .sx0 = sx0,
+        .sy0 = sy0,
+        .kh = octx->src[0]->ne[1],
+        .kw = octx->src[0]->ne[0],
+        .ic_blocks = src->ne[2] / 32,
+        .n_groups = src->ne[2] / kp->activation_group_channels,
+        .group_channels = kp->activation_group_channels,
+        .k_tiles = kp->k_tiles,
+        .activation_k_tiles = kp->activation_k_tiles,
+        .tile_w = tile_w,
+        .tile_h = tile_h,
+        .x_tiles = tile_w / HTP_CONV2D_TILE_W,
+        .m_tiles = (tile_w / HTP_CONV2D_TILE_W) * tile_h,
+        .clear = sx_end - sx_begin != (int32_t) halo_w ||
+                 sy_end - sy_begin != (int32_t) halo_h,
+        .wait_weight_dma = wait_weight_dma,
+        .upscale2 = upscale2,
+        .direct_upscale2 = upscale2 && sx0 >= 0 && sy0 >= 0 &&
+                           sx0 + (int32_t) halo_w <= (int32_t) iw &&
+                           sy0 + (int32_t) halo_h <= (int32_t) ih &&
+                           (uint32_t) (sx0 / 2) +
+                               (tile_w / HTP_CONV2D_TILE_W - 1u) *
+                               (HTP_CONV2D_TILE_W / 2u) +
+                               (src->type == HTP_TYPE_F16 ? 64u : 32u) <= src->ne[0],
+        .input_silu = (kp->flags & HTP_CONV2D_INPUT_SILU) != 0,
+        .input_f16 = src->type == HTP_TYPE_F16,
+    };
+}
+
+static void conv2d_prepare_activation(
+        struct htp_ops_context * octx,
+        float * x,
+        __fp16 * a,
+        uint32_t x0,
+        uint32_t y0,
+        uint32_t tile_w,
+        uint32_t tile_h,
+        bool wait_weight_dma,
+        const struct htp_conv2d_kernel_params * kp) {
+    struct conv2d_prepare_state state;
+    conv2d_init_prepare_state(octx, x, a, x0, y0, tile_w, tile_h,
+                              wait_weight_dma, kp, &state);
+    work_queue_run(octx->ctx->work_queue, conv2d_prepare_activation_worker, &state, octx->n_threads);
+}
+
+struct conv2d_pipeline_state {
+    struct conv2d_prepare_state prepare;
+    struct conv2d_store_state store;
+    struct conv2d_hmx_job * next_job;
+    atomic_uint prepare_barrier;
+    atomic_uint next_store_pair;
+    atomic_bool hmx_ready;
+    uint32_t block;
+    bool has_next;
+};
+
+static void conv2d_pipeline_worker(unsigned int nth, unsigned int ith, void * data) {
+    struct conv2d_pipeline_state * st = (struct conv2d_pipeline_state *) data;
+    if (st->has_next) {
+        conv2d_prepare_activation_worker(nth, ith, &st->prepare);
+    }
+
+    atomic_fetch_sub_explicit(&st->prepare_barrier, 1, memory_order_release);
+    if (ith == 0) {
+        hmx_queue_pop(st->store.ctx->hmx_queue);
+        htp_trace_event_start(&st->store.ctx->trace[0], HTP_TRACE_EVT_HVX_O_PROC, st->block);
+        atomic_store_explicit(&st->hmx_ready, true, memory_order_release);
+
+        if (st->has_next) {
+            while (atomic_load_explicit(&st->prepare_barrier, memory_order_acquire) != 0) {
+                hex_pause();
+            }
+            htp_trace_event_stop(&st->store.ctx->trace[0], HTP_TRACE_EVT_HVX_A_PREP, st->block + 1);
+            hmx_queue_push(st->store.ctx->hmx_queue,
+                           hmx_queue_make_desc(conv2d_hmx_worker, st->next_job));
+        }
+    } else {
+        while (!atomic_load_explicit(&st->hmx_ready, memory_order_acquire)) {
+            hex_pause();
+        }
+    }
+
+    conv2d_store_output_worker(nth, ith, &st->store);
+}
+
+int op_conv2d(struct htp_ops_context * octx) {
+    const struct htp_tensor * weight = octx->src[0];
+    const struct htp_tensor * src = octx->src[1];
+    const struct htp_tensor * dst = octx->dst;
+    const struct htp_conv2d_kernel_params * kp =
+        (const struct htp_conv2d_kernel_params *) octx->kernel_params;
+
+    if (weight->type != HTP_TYPE_F16 ||
+        (src->type != HTP_TYPE_F32 && src->type != HTP_TYPE_F16) ||
+        (dst->type != HTP_TYPE_F32 && dst->type != HTP_TYPE_F16) ||
+        (src->type == HTP_TYPE_F16 && dst->type != HTP_TYPE_F16) ||
+        src->ne[3] != 1 || dst->ne[3] != 1 || kp->vtcm_size > octx->ctx->vtcm_size) {
+        return HTP_STATUS_NO_SUPPORT;
+    }
+    if ((kp->flags & HTP_CONV2D_BIAS) &&
+        (!octx->src[2] || octx->src[2]->type != HTP_TYPE_F32 ||
+         (uint64_t) octx->src[2]->ne[0] * octx->src[2]->ne[1] *
+         octx->src[2]->ne[2] * octx->src[2]->ne[3] != weight->ne[3])) {
+        return HTP_STATUS_NO_SUPPORT;
+    }
+    if ((kp->flags & HTP_CONV2D_RESIDUAL) &&
+        (!octx->src[3] ||
+         octx->src[3]->ne[0] != dst->ne[0] || octx->src[3]->ne[1] != dst->ne[1] ||
+         octx->src[3]->ne[2] != dst->ne[2] || octx->src[3]->ne[3] != dst->ne[3] ||
+         octx->src[3]->type != dst->type)) {
+        return HTP_STATUS_NO_SUPPORT;
+    }
+    if ((kp->flags & HTP_CONV2D_UPSCALE2) &&
+        (dst->ne[0] != 2u * src->ne[0] || dst->ne[1] != 2u * src->ne[1] ||
+         dst->ne[2] != weight->ne[3] || src->ne[2] != weight->ne[2])) {
+        return HTP_STATUS_NO_SUPPORT;
+    }
+    if (octx->flags & HTP_OPFLAGS_SKIP_COMPUTE) {
+        return HTP_STATUS_OK;
+    }
+
+    const uint32_t kh = weight->ne[1];
+    const uint32_t kw = weight->ne[0];
+    const uint32_t ic = weight->ne[2];
+    const uint32_t oc = weight->ne[3];
+    const uint32_t ow = dst->ne[0];
+    const uint32_t oh = dst->ne[1];
+    uint8_t * base = (uint8_t *) octx->ctx->vtcm_base;
+    __fp16 * b = VTCM_LAYOUT_PTR(__fp16, base, kp->off_weight);
+    float * xbuf[2] = {
+        VTCM_LAYOUT_PTR(float, base, kp->off_x[0]),
+        VTCM_LAYOUT_PTR(float, base, kp->off_x[1]),
+    };
+    __fp16 * abuf[2] = {
+        VTCM_LAYOUT_PTR(__fp16, base, kp->off_a[0]),
+        VTCM_LAYOUT_PTR(__fp16, base, kp->off_a[1]),
+    };
+    __fp16 * cbuf[2] = {
+        VTCM_LAYOUT_PTR(__fp16, base, kp->off_c[0]),
+        VTCM_LAYOUT_PTR(__fp16, base, kp->off_c[1]),
+    };
+    __fp16 * scales = VTCM_LAYOUT_PTR(__fp16, base, kp->off_scales);
+
+    dma_queue * q = octx->ctx->dma[0];
+    if (!dma_queue_push(q, dma_make_ptr(b, (const void *) (uintptr_t) weight->data), HTP_CONV2D_TILE_BYTES,
+                        HTP_CONV2D_TILE_BYTES, HTP_CONV2D_TILE_BYTES,
+                        kp->k_tiles * kp->n_tiles)) {
+        return HTP_STATUS_INTERNAL_ERR;
+    }
+    hmx_init_column_scales(scales, Q6_V_vsplat_R(0x3c00));
+
+    const uint32_t nx = (ow + kp->tile_w - 1) / kp->tile_w;
+    const uint32_t ny = (oh + kp->tile_h - 1) / kp->tile_h;
+    const uint32_t nblocks = nx * ny;
+    float * dst_data = (float *) dst->data;
+    const float * bias = (kp->flags & HTP_CONV2D_BIAS)
+                       ? (const float *) octx->src[2]->data : NULL;
+    const float * residual = (kp->flags & HTP_CONV2D_RESIDUAL)
+                           ? (const float *) octx->src[3]->data : NULL;
+    struct conv2d_hmx_job jobs[2];
+
+    const uint32_t first_valid_w = hex_smin(kp->tile_w, ow);
+    const uint32_t first_tile_w =
+        htp_conv2d_align_up(first_valid_w, HTP_CONV2D_TILE_W);
+    const uint32_t first_tile_h = hex_smin(kp->tile_h, oh);
+    const uint32_t first_m_tiles = (first_tile_w / HTP_CONV2D_TILE_W) * first_tile_h;
+    htp_trace_event_start(&octx->ctx->trace[0], HTP_TRACE_EVT_HVX_A_PREP, 0);
+    conv2d_prepare_activation(octx, xbuf[0], abuf[0], 0, 0,
+                              first_tile_w, first_tile_h, true, kp);
+    htp_trace_event_stop(&octx->ctx->trace[0], HTP_TRACE_EVT_HVX_A_PREP, 0);
+
+    jobs[0] = (struct conv2d_hmx_job) {
+        .c = cbuf[0], .a = b, .b = abuf[0], .scales = scales,
+        .mt = kp->n_tiles, .nt = first_m_tiles, .kt = kp->k_tiles,
+        .kh = kh, .x_tiles = first_tile_w / HTP_CONV2D_TILE_W,
+        .tile_h = first_tile_h, .ic_blocks = ic / 32,
+    };
+    hmx_queue_push(octx->ctx->hmx_queue, hmx_queue_make_desc(conv2d_hmx_worker, &jobs[0]));
+
+    for (uint32_t block = 0; block < nblocks; ++block) {
+        const uint32_t slot = block & 1u;
+        const uint32_t bx = block % nx;
+        const uint32_t by = block / nx;
+        const uint32_t x0 = bx * kp->tile_w;
+        const uint32_t y0 = by * kp->tile_h;
+        const uint32_t valid_w = hex_smin(kp->tile_w, ow - x0);
+        const uint32_t tile_w =
+            htp_conv2d_align_up(valid_w, HTP_CONV2D_TILE_W);
+        const uint32_t tile_h = hex_smin(kp->tile_h, oh - y0);
+        const uint32_t x_tiles = tile_w / HTP_CONV2D_TILE_W;
+        const uint32_t m_tiles = x_tiles * tile_h;
+
+        struct conv2d_pipeline_state pipeline = {
+            .store = {
+                .ctx = octx->ctx, .c = cbuf[slot], .dst = dst_data,
+                .bias = bias, .residual = residual,
+                .ow = ow, .oh = oh, .oc = oc, .n_tiles = kp->n_tiles,
+                .x0 = x0, .y0 = y0, .tile_w = tile_w, .valid_w = valid_w,
+                .tile_h = tile_h,
+                .x_tiles = x_tiles, .m_tiles = m_tiles,
+                .final_block = block + 1 == nblocks,
+                .next_pair = NULL,
+                .flags = kp->flags,
+                .output_f16 = dst->type == HTP_TYPE_F16,
+            },
+            .next_job = NULL,
+            .block = block,
+            .has_next = block + 1 < nblocks,
+        };
+
+        if (pipeline.has_next) {
+            const uint32_t next = block + 1;
+            const uint32_t next_slot = next & 1u;
+            const uint32_t next_x0 = (next % nx) * kp->tile_w;
+            const uint32_t next_y0 = (next / nx) * kp->tile_h;
+            const uint32_t next_valid_w = hex_smin(kp->tile_w, ow - next_x0);
+            const uint32_t next_tile_w =
+                htp_conv2d_align_up(next_valid_w, HTP_CONV2D_TILE_W);
+            const uint32_t next_tile_h = hex_smin(kp->tile_h, oh - next_y0);
+            const uint32_t next_m_tiles =
+                (next_tile_w / HTP_CONV2D_TILE_W) * next_tile_h;
+
+            htp_trace_event_start(&octx->ctx->trace[0], HTP_TRACE_EVT_HVX_A_PREP, next);
+            conv2d_init_prepare_state(octx, xbuf[next_slot], abuf[next_slot],
+                                      next_x0, next_y0, next_tile_w, next_tile_h,
+                                      false, kp, &pipeline.prepare);
+            jobs[next_slot] = (struct conv2d_hmx_job) {
+                .c = cbuf[next_slot], .a = b, .b = abuf[next_slot], .scales = scales,
+                .mt = kp->n_tiles, .nt = next_m_tiles, .kt = kp->k_tiles,
+                .kh = kh, .x_tiles = next_tile_w / HTP_CONV2D_TILE_W,
+                .tile_h = next_tile_h, .ic_blocks = ic / 32,
+            };
+            pipeline.next_job = &jobs[next_slot];
+        }
+
+        atomic_init(&pipeline.prepare_barrier, octx->n_threads);
+        atomic_init(&pipeline.next_store_pair, 0);
+        atomic_init(&pipeline.hmx_ready, false);
+        pipeline.store.next_pair = &pipeline.next_store_pair;
+        work_queue_run(octx->ctx->work_queue, conv2d_pipeline_worker, &pipeline, octx->n_threads);
+        htp_trace_event_stop(&octx->ctx->trace[0], HTP_TRACE_EVT_HVX_O_PROC, block);
+    }
+
+    return HTP_STATUS_OK;
+}
diff --git a/ggml/src/ggml-hexagon/htp/conv2d-ops.h b/ggml/src/ggml-hexagon/htp/conv2d-ops.h
new file mode 100644
index 000000000000..40f4753956dc
--- /dev/null
+++ b/ggml/src/ggml-hexagon/htp/conv2d-ops.h
@@ -0,0 +1,121 @@
+#ifndef HTP_CONV2D_OPS_H
+#define HTP_CONV2D_OPS_H
+
+#include <stdint.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#define HTP_CONV2D_TILE_W 32u
+#define HTP_CONV2D_TILE_BYTES 2048u
+
+enum htp_conv2d_flags {
+    HTP_CONV2D_BIAS     = 1u << 0,
+    HTP_CONV2D_RESIDUAL = 1u << 1,
+    HTP_CONV2D_UPSCALE2 = 1u << 2,
+    HTP_CONV2D_INPUT_SILU = 1u << 3,
+};
+
+struct htp_conv2d_kernel_params {
+    uint32_t tile_w;
+    uint32_t tile_h;
+    uint32_t m_tiles;
+    uint32_t k_tiles;
+    uint32_t activation_k_tiles;
+    uint32_t n_tiles;
+    uint32_t halo_w;
+    uint32_t halo_h;
+    uint32_t activation_group_channels;
+
+    uint32_t off_weight;
+    uint32_t weight_bytes;
+    uint32_t off_x[2];
+    uint32_t x_slot_bytes;
+    uint32_t off_a[2];
+    uint32_t a_slot_bytes;
+    uint32_t off_c[2];
+    uint32_t c_slot_bytes;
+    uint32_t off_scales;
+    uint32_t vtcm_size;
+    uint32_t flags;
+};
+
+static inline uint32_t htp_conv2d_align_up(uint32_t value, uint32_t alignment) {
+    return (value + alignment - 1u) & ~(alignment - 1u);
+}
+
+static inline uint32_t htp_conv2d_layout_build(
+        struct htp_conv2d_kernel_params * p,
+        uint32_t kh,
+        uint32_t kw,
+        uint32_t ic,
+        uint32_t oc,
+        uint32_t tile_w,
+        uint32_t tile_h,
+        uint32_t n_threads) {
+    const uint32_t k = kh * kw * ic;
+    uint32_t off = 0;
+
+    p->tile_w = tile_w;
+    p->flags = 0;
+    p->tile_h = tile_h;
+    p->m_tiles = (tile_w / HTP_CONV2D_TILE_W) * tile_h;
+    p->k_tiles = k / 32u;
+    p->activation_k_tiles = kw * ic / 32u;
+    p->n_tiles = (oc + 31u) / 32u;
+    p->halo_w = tile_w + kw - 1u;
+    p->halo_h = tile_h + kh - 1u;
+
+    p->off_weight = off;
+    p->weight_bytes = k * p->n_tiles * 32u * sizeof(uint16_t);
+    off = htp_conv2d_align_up(off + p->weight_bytes, HTP_CONV2D_TILE_BYTES);
+
+    const uint32_t ic_blocks = ic / 32u;
+    uint32_t groups_per_ic_block = 1;
+    while (groups_per_ic_block < 16u &&
+           ic_blocks * groups_per_ic_block < 2u * n_threads) {
+        groups_per_ic_block *= 2u;
+    }
+    p->activation_group_channels = 32u / groups_per_ic_block;
+    const uint32_t n_groups = ic / p->activation_group_channels;
+    uint32_t x_slots = 0;
+    for (uint32_t i = 0; i < n_threads; ++i) {
+        const uint32_t first = (uint64_t) n_groups * i / n_threads;
+        const uint32_t last = (uint64_t) n_groups * (i + 1u) / n_threads;
+        const uint32_t count = last - first;
+        x_slots += count < 2u ? count : 2u;
+    }
+    const uint32_t x_f32_bytes = x_slots * p->activation_group_channels *
+                                 p->halo_w * p->halo_h * sizeof(float);
+    p->off_x[0] = off;
+    p->x_slot_bytes = x_f32_bytes;
+    off = htp_conv2d_align_up(off + p->x_slot_bytes, HTP_CONV2D_TILE_BYTES);
+    p->off_x[1] = p->off_x[0];
+
+    p->a_slot_bytes = HTP_CONV2D_TILE_W * (tile_w / HTP_CONV2D_TILE_W) *
+                      p->halo_h * (kw * ic) * sizeof(uint16_t);
+    for (uint32_t i = 0; i < 2; ++i) {
+        p->off_a[i] = off;
+        off = htp_conv2d_align_up(off + p->a_slot_bytes, HTP_CONV2D_TILE_BYTES);
+    }
+
+    p->c_slot_bytes = HTP_CONV2D_TILE_W * p->m_tiles *
+                      p->n_tiles * 32u * sizeof(uint16_t);
+    for (uint32_t i = 0; i < 2; ++i) {
+        p->off_c[i] = off;
+        off = htp_conv2d_align_up(off + p->c_slot_bytes, HTP_CONV2D_TILE_BYTES);
+    }
+
+    p->off_scales = off;
+    off = htp_conv2d_align_up(off + 256u, HTP_CONV2D_TILE_BYTES);
+
+    p->vtcm_size = off;
+    return off;
+}
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
diff --git a/ggml/src/ggml-hexagon/htp/cpy-ops.c b/ggml/src/ggml-hexagon/htp/cpy-ops.c
index 7f01a8c1e043..2dd5412e8a89 100644
--- a/ggml/src/ggml-hexagon/htp/cpy-ops.c
+++ b/ggml/src/ggml-hexagon/htp/cpy-ops.c
@@ -49,6 +49,100 @@ struct htp_copy_context {
     struct fastdiv_values div_ne02_ne01_ne00;
 };
 
+static inline void hvx_transpose_32x32_f32(HVX_Vector matrix[32]) {
+    HVX_Vector temp[32];
+
+    for (int i = 0; i < 16; ++i) {
+        const HVX_VectorPair pair = Q6_W_vshuff_VVR(matrix[2 * i + 1], matrix[2 * i], -4);
+        temp[2 * i]     = Q6_V_lo_W(pair);
+        temp[2 * i + 1] = Q6_V_hi_W(pair);
+    }
+    for (int base = 0; base < 32; base += 4) {
+        const HVX_VectorPair pair0 = Q6_W_vshuff_VVR(temp[base + 2], temp[base], -8);
+        const HVX_VectorPair pair1 = Q6_W_vshuff_VVR(temp[base + 3], temp[base + 1], -8);
+        matrix[base]     = Q6_V_lo_W(pair0);
+        matrix[base + 1] = Q6_V_hi_W(pair0);
+        matrix[base + 2] = Q6_V_lo_W(pair1);
+        matrix[base + 3] = Q6_V_hi_W(pair1);
+    }
+    for (int base = 0; base < 32; base += 8) {
+        for (int i = 0; i < 4; ++i) {
+            const HVX_VectorPair pair = Q6_W_vshuff_VVR(matrix[base + i + 4], matrix[base + i], -16);
+            temp[base + 2 * i]     = Q6_V_lo_W(pair);
+            temp[base + 2 * i + 1] = Q6_V_hi_W(pair);
+        }
+    }
+    for (int base = 0; base < 32; base += 16) {
+        for (int i = 0; i < 8; ++i) {
+            const HVX_VectorPair pair = Q6_W_vshuff_VVR(temp[base + i + 8], temp[base + i], -32);
+            matrix[base + 2 * i]     = Q6_V_lo_W(pair);
+            matrix[base + 2 * i + 1] = Q6_V_hi_W(pair);
+        }
+    }
+    for (int i = 0; i < 16; ++i) {
+        const HVX_VectorPair pair = Q6_W_vshuff_VVR(matrix[i + 16], matrix[i], -64);
+        temp[2 * i]     = Q6_V_lo_W(pair);
+        temp[2 * i + 1] = Q6_V_hi_W(pair);
+    }
+    for (int i = 0; i < 32; ++i) {
+        matrix[i] = temp[i];
+    }
+}
+
+struct htp_transpose_f32_context {
+    const float * src;
+    float * dst;
+    uint32_t rows;
+    uint32_t cols;
+};
+
+static void cpy_thread_transpose_f32(unsigned int nth, unsigned int ith, void * data) {
+    const struct htp_transpose_f32_context * ctx = (const struct htp_transpose_f32_context *) data;
+    const uint32_t row_tiles = ctx->rows / 32u;
+    const uint32_t col_tiles = ctx->cols / 32u;
+    const uint32_t tiles = row_tiles * col_tiles;
+    HVX_Vector matrix[32] __attribute__((aligned(128)));
+
+    for (uint32_t tile = ith; tile < tiles; tile += nth) {
+        const uint32_t row_tile = tile / col_tiles;
+        const uint32_t col_tile = tile - row_tile * col_tiles;
+        const uint32_t row0 = row_tile * 32u;
+        const uint32_t col0 = col_tile * 32u;
+        hex_l2fetch(ctx->src + (size_t) row0 * ctx->cols + col0,
+                    32u * sizeof(float), ctx->cols * sizeof(float), 32u);
+        for (uint32_t row = 0; row < 32u; ++row) {
+            matrix[row] = hvx_vmemu(ctx->src + (size_t) (row0 + row) * ctx->cols + col0);
+        }
+        hvx_transpose_32x32_f32(matrix);
+        for (uint32_t row = 0; row < 32u; ++row) {
+            hvx_vmemu(ctx->dst + (size_t) (col0 + row) * ctx->rows + row0) = matrix[row];
+        }
+    }
+}
+
+static bool cpy_is_contiguous_transpose_f32(
+        const struct htp_tensor * src,
+        const struct htp_tensor * dst,
+        uint32_t * rows,
+        uint32_t * cols) {
+    const uint64_t rest = (uint64_t) src->ne[1] * src->ne[2] * src->ne[3];
+    if (src->type != HTP_TYPE_F32 || dst->type != HTP_TYPE_F32 ||
+        src->ne[3] != 1 || src->ne[0] % 32u != 0 || rest == 0 ||
+        rest > UINT32_MAX || rest % 32u != 0 ||
+        src->nb[1] != sizeof(float) ||
+        (src->ne[2] > 1 && src->nb[2] != (uint64_t) src->ne[1] * sizeof(float)) ||
+        src->nb[0] != rest * sizeof(float) ||
+        dst->nb[0] != sizeof(float) ||
+        dst->nb[1] != (uint64_t) dst->ne[0] * sizeof(float) ||
+        dst->nb[2] != (uint64_t) dst->ne[1] * dst->nb[1] ||
+        dst->nb[3] != (uint64_t) dst->ne[2] * dst->nb[2]) {
+        return false;
+    }
+    *rows = src->ne[0];
+    *cols = (uint32_t) rest;
+    return true;
+}
+
 #define cpy_preamble                              \
     const struct htp_tensor *src0 = octx->src[0]; \
     const struct htp_tensor *dst  = octx->dst;    \
@@ -298,6 +392,23 @@ static int exec_cpy(struct htp_ops_context * octx, bool * use_dma) {
     cpy_preamble;
     *use_dma = false;
 
+    uint32_t transpose_rows;
+    uint32_t transpose_cols;
+    if (octx->ctx->mdev.count <= 1 &&
+        cpy_is_contiguous_transpose_f32(src0, dst, &transpose_rows, &transpose_cols)) {
+        if (!(octx->flags & HTP_OPFLAGS_SKIP_COMPUTE)) {
+            const struct htp_transpose_f32_context transpose = {
+                .src = (const float *) src0->data,
+                .dst = (float *) dst->data,
+                .rows = transpose_rows,
+                .cols = transpose_cols,
+            };
+            work_queue_run(octx->ctx->work_queue, cpy_thread_transpose_f32,
+                           (void *) &transpose, octx->n_threads);
+        }
+        return HTP_STATUS_OK;
+    }
+
     struct htp_copy_context ct;
     ct.octx = octx;
 
diff --git a/ggml/src/ggml-hexagon/htp/groupnorm-ops.c b/ggml/src/ggml-hexagon/htp/groupnorm-ops.c
new file mode 100644
index 000000000000..5d88566a1a80
--- /dev/null
+++ b/ggml/src/ggml-hexagon/htp/groupnorm-ops.c
@@ -0,0 +1,396 @@
+#include <hexagon_protos.h>
+#include <math.h>
+#include <stdatomic.h>
+#include <stdint.h>
+
+#include "groupnorm-ops.h"
+#include "htp-ctx.h"
+#include "hvx-reduce.h"
+#include "hvx-sigmoid.h"
+#include "hvx-utils.h"
+
+struct group_norm_state {
+    struct htp_context * ctx;
+    const void * src;
+    const float * weight;
+    const float * bias;
+    void * dst;
+    uint32_t spatial;
+    uint32_t width;
+    uint32_t channels;
+    uint32_t groups;
+    uint32_t channels_per_group;
+    uint32_t batches;
+    uint32_t element_size;
+    float epsilon;
+    uint32_t flags;
+    uint8_t * vtcm;
+    uint32_t vtcm_size;
+    atomic_uint next_job;
+};
+
+struct group_norm_accum {
+    HVX_Vector sum[4];
+    HVX_Vector square[4];
+};
+
+typedef void (*group_norm_accumulate_fn)(struct group_norm_accum *, const void *, uint32_t);
+typedef void (*group_norm_transform_fn)(const struct group_norm_state *, uint32_t, uint32_t, const void *, void *, uint32_t, HVX_Vector, HVX_Vector);
+
+static inline void group_norm_accum_init(struct group_norm_accum * acc) {
+    const HVX_Vector zero = Q6_V_vzero();
+    for (int i = 0; i < 4; ++i) {
+        acc->sum[i] = zero;
+        acc->square[i] = zero;
+    }
+}
+
+static inline void group_norm_accumulate_vector(struct group_norm_accum * acc, int lane, HVX_Vector x) {
+#if __HVX_ARCH__ >= 79
+    acc->sum[lane] = Q6_Vsf_vadd_VsfVsf(acc->sum[lane], x);
+    acc->square[lane] = Q6_Vsf_vadd_VsfVsf(acc->square[lane], Q6_Vsf_vmpy_VsfVsf(x, x));
+#else
+    const HVX_Vector zero = Q6_V_vzero();
+    acc->sum[lane] = Q6_Vqf32_vadd_Vqf32Vqf32(acc->sum[lane], Q6_Vqf32_vadd_VsfVsf(x, zero));
+    acc->square[lane] = Q6_Vqf32_vadd_Vqf32Vqf32(acc->square[lane], Q6_Vqf32_vmpy_VsfVsf(x, x));
+#endif
+}
+
+static inline __attribute__((always_inline)) void group_norm_accumulate_f32(struct group_norm_accum * acc, const void * data, uint32_t n) {
+    const HVX_Vector * src = (const HVX_Vector *) data;
+    const uint32_t nvec = n / 32u;
+    uint32_t i = 0;
+
+    for (; i + 3u < nvec; i += 4u) {
+#pragma unroll(4)
+        for (int lane = 0; lane < 4; ++lane) {
+            group_norm_accumulate_vector(acc, lane, src[i + lane]);
+        }
+    }
+    for (; i < nvec; ++i) {
+        group_norm_accumulate_vector(acc, 0, src[i]);
+    }
+}
+
+static inline __attribute__((always_inline)) void group_norm_accumulate_f16(struct group_norm_accum * acc, const void * data, uint32_t n) {
+    const __fp16 * src = (const __fp16 *) data;
+    const uint32_t nvec = n / 64u;
+    uint32_t i = 0;
+
+    for (; i + 1u < nvec; i += 2u) {
+        const HVX_VectorPair p0 = hvx_vec_f16_to_f32(hvx_vmemu(src + (size_t) i * 64u));
+        const HVX_VectorPair p1 = hvx_vec_f16_to_f32(hvx_vmemu(src + (size_t) (i + 1u) * 64u));
+        const HVX_Vector x0 = Q6_V_lo_W(p0);
+        const HVX_Vector x1 = Q6_V_hi_W(p0);
+        const HVX_Vector x2 = Q6_V_lo_W(p1);
+        const HVX_Vector x3 = Q6_V_hi_W(p1);
+        acc->sum[0] = hvx_vec_add_f32_f32(acc->sum[0], x0);
+        acc->sum[1] = hvx_vec_add_f32_f32(acc->sum[1], x1);
+        acc->sum[2] = hvx_vec_add_f32_f32(acc->sum[2], x2);
+        acc->sum[3] = hvx_vec_add_f32_f32(acc->sum[3], x3);
+        acc->square[0] = hvx_vec_add_f32_f32(acc->square[0], hvx_vec_mul_f32_f32(x0, x0));
+        acc->square[1] = hvx_vec_add_f32_f32(acc->square[1], hvx_vec_mul_f32_f32(x1, x1));
+        acc->square[2] = hvx_vec_add_f32_f32(acc->square[2], hvx_vec_mul_f32_f32(x2, x2));
+        acc->square[3] = hvx_vec_add_f32_f32(acc->square[3], hvx_vec_mul_f32_f32(x3, x3));
+    }
+    if (i < nvec) {
+        const HVX_VectorPair p = hvx_vec_f16_to_f32(hvx_vmemu(src + (size_t) i * 64u));
+        const HVX_Vector x0 = Q6_V_lo_W(p);
+        const HVX_Vector x1 = Q6_V_hi_W(p);
+        acc->sum[0] = hvx_vec_add_f32_f32(acc->sum[0], x0);
+        acc->sum[1] = hvx_vec_add_f32_f32(acc->sum[1], x1);
+        acc->square[0] = hvx_vec_add_f32_f32(acc->square[0], hvx_vec_mul_f32_f32(x0, x0));
+        acc->square[1] = hvx_vec_add_f32_f32(acc->square[1], hvx_vec_mul_f32_f32(x1, x1));
+    }
+
+    const uint32_t tail = n % 64u;
+    if (tail) {
+        HVX_Vector x = hvx_vmemu(src + (size_t) nvec * 64u);
+        x = Q6_V_vmux_QVV(Q6_Q_vsetq2_R(tail * sizeof(__fp16)), x, Q6_V_vzero());
+        const HVX_VectorPair p = hvx_vec_f16_to_f32(x);
+        const HVX_Vector x0 = Q6_V_lo_W(p);
+        const HVX_Vector x1 = Q6_V_hi_W(p);
+        acc->sum[0] = hvx_vec_add_f32_f32(acc->sum[0], x0);
+        acc->sum[1] = hvx_vec_add_f32_f32(acc->sum[1], x1);
+        acc->square[0] = hvx_vec_add_f32_f32(acc->square[0], hvx_vec_mul_f32_f32(x0, x0));
+        acc->square[1] = hvx_vec_add_f32_f32(acc->square[1], hvx_vec_mul_f32_f32(x1, x1));
+    }
+}
+
+static inline void group_norm_accum_finish(const struct group_norm_accum * acc, uint32_t n, float * mean, float * mean_square) {
+#if __HVX_ARCH__ >= 79
+    const HVX_Vector sum = Q6_Vsf_vadd_VsfVsf(Q6_Vsf_vadd_VsfVsf(acc->sum[0], acc->sum[1]), Q6_Vsf_vadd_VsfVsf(acc->sum[2], acc->sum[3]));
+    const HVX_Vector square = Q6_Vsf_vadd_VsfVsf(Q6_Vsf_vadd_VsfVsf(acc->square[0], acc->square[1]), Q6_Vsf_vadd_VsfVsf(acc->square[2], acc->square[3]));
+    *mean = hvx_vec_get_f32(hvx_vec_reduce_sum_f32(sum)) / (float) n;
+    *mean_square = hvx_vec_get_f32(hvx_vec_reduce_sum_f32(square)) / (float) n;
+#else
+    const HVX_Vector sum = Q6_Vqf32_vadd_Vqf32Vqf32(Q6_Vqf32_vadd_Vqf32Vqf32(acc->sum[0], acc->sum[1]), Q6_Vqf32_vadd_Vqf32Vqf32(acc->sum[2], acc->sum[3]));
+    const HVX_Vector square = Q6_Vqf32_vadd_Vqf32Vqf32(Q6_Vqf32_vadd_Vqf32Vqf32(acc->square[0], acc->square[1]), Q6_Vqf32_vadd_Vqf32Vqf32(acc->square[2], acc->square[3]));
+    *mean = hvx_vec_get_f32(hvx_vec_reduce_sum_f32(Q6_Vsf_equals_Vqf32(sum))) / (float) n;
+    *mean_square = hvx_vec_get_f32(hvx_vec_reduce_sum_f32(Q6_Vsf_equals_Vqf32(square))) / (float) n;
+#endif
+}
+
+static inline void group_norm_dma_push(dma_queue * queue, void * dst, const void * src, uint32_t bytes) {
+    while (!dma_queue_push(queue, dma_make_ptr(dst, src), bytes, bytes, bytes, 1)) {
+        dma_queue_flush(queue);
+    }
+}
+
+static inline __attribute__((always_inline)) void group_norm_stats_dma(const struct group_norm_state * st, dma_queue * queue, const void * src, uint32_t n, void * slots[2], uint32_t chunk, group_norm_accumulate_fn accumulate, float * mean, float * mean_square) {
+    const uint32_t chunks = (n + chunk - 1u) / chunk;
+    for (uint32_t i = 0; i < chunks && i < 2u; ++i) {
+        const uint32_t count = hex_smin(n - i * chunk, chunk);
+        group_norm_dma_push(queue, slots[i], (const uint8_t *) src + (size_t) i * chunk * st->element_size, count * st->element_size);
+    }
+
+    struct group_norm_accum acc;
+    group_norm_accum_init(&acc);
+    for (uint32_t i = 0; i < chunks; ++i) {
+        const uint32_t count = hex_smin(n - i * chunk, chunk);
+        const dma_ptr ready = dma_queue_pop(queue);
+        accumulate(&acc, (const void *) ready.dst, count);
+        if (i + 2u < chunks) {
+            const uint32_t next = i + 2u;
+            const uint32_t next_count = hex_smin(n - next * chunk, chunk);
+            group_norm_dma_push(queue, slots[i & 1u], (const uint8_t *) src + (size_t) next * chunk * st->element_size, next_count * st->element_size);
+        }
+    }
+    dma_queue_flush(queue);
+    group_norm_accum_finish(&acc, n, mean, mean_square);
+}
+
+static inline __attribute__((always_inline)) void group_norm_transform_f32(const struct group_norm_state * st, uint32_t channel0, uint32_t group_offset, const void * src_data, void * dst_data, uint32_t n, HVX_Vector mean, HVX_Vector inv_std) {
+    const float * src = (const float *) src_data;
+    float * dst = (float *) dst_data;
+    const bool affine = (st->flags & HTP_GROUP_NORM_AFFINE) != 0;
+    const bool silu = (st->flags & HTP_GROUP_NORM_SILU) != 0;
+    uint32_t done = 0;
+
+    while (done < n) {
+        const uint32_t offset = group_offset + done;
+        const uint32_t channel = channel0 + offset / st->spatial;
+        const uint32_t in_channel = offset % st->spatial;
+        const uint32_t count = hex_smin(st->spatial - in_channel, n - done);
+        const HVX_Vector weight = hvx_vec_splat_f32(affine ? st->weight[channel] : 1.0f);
+        const HVX_Vector bias = hvx_vec_splat_f32(affine ? st->bias[channel] : 0.0f);
+        const HVX_Vector scale = hvx_vec_mul_f32_f32(inv_std, weight);
+        const HVX_Vector offset_vec = hvx_vec_sub_f32_f32(bias, hvx_vec_mul_f32_f32(mean, scale));
+        const HVX_Vector * input = (const HVX_Vector *) (src + done);
+        HVX_Vector * output = (HVX_Vector *) (dst + done);
+        const uint32_t nvec = count / 32u;
+
+#pragma unroll(4)
+        for (uint32_t i = 0; i < nvec; ++i) {
+            output[i] = hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(input[i], scale), offset_vec);
+        }
+        done += count;
+    }
+
+    if (silu) {
+        for (uint32_t i = 0; i < n; i += st->width) {
+            hvx_sigmoid_f32_aa((uint8_t *) (dst + i), (const uint8_t *) (dst + i), st->width);
+            hvx_mul_f32_aaa((uint8_t *) (dst + i), (const uint8_t *) (src + i), (const uint8_t *) (dst + i), st->width);
+        }
+    }
+}
+
+static const uint16_t group_norm_silu_lut_a[64] __attribute__((aligned(128))) = {
+    0x997f, 0x9c3f, 0x9e88, 0xa0ff, 0xa399, 0xa5bb, 0xa849, 0xaa54,
+    0xac9b, 0xae91, 0xb08e, 0xb212, 0xb3a1, 0xb461, 0xb44e, 0xb20a,
+    0x0000, 0x34fb, 0x39d9, 0x3ce8, 0x3f0c, 0x409f, 0x41b7, 0x42cb,
+    0x43db, 0x4473, 0x44f7, 0x457a, 0x45fc, 0x467e, 0x46fe, 0x477f,
+};
+
+static const uint16_t group_norm_silu_lut_b[64] __attribute__((aligned(128))) = {
+    0x94b9, 0x973b, 0x9981, 0x9c28, 0x9e3a, 0xa09c, 0xa2ba, 0xa4d1,
+    0xa6b7, 0xa881, 0xa9b1, 0xaa81, 0xaa14, 0xa60e, 0x283e, 0x301d,
+    0x3404, 0x35f7, 0x3779, 0x382f, 0x3860, 0x3867, 0x385a, 0x3848,
+    0x3836, 0x3827, 0x381b, 0x3812, 0x380c, 0x3808, 0x3806, 0x3804,
+};
+
+static const uint16_t group_norm_silu_lut_c[64] __attribute__((aligned(128))) = {
+    0x8d14, 0x8fa2, 0x91ae, 0x942c, 0x9607, 0x9841, 0x99d0, 0x9b8f,
+    0x9c85, 0x9c99, 0x99ed, 0x182f, 0x2245, 0x2741, 0x2a05, 0x2bb7,
+    0x2bb7, 0x2a05, 0x2741, 0x2245, 0x182f, 0x99ed, 0x9c99, 0x9c85,
+    0x9b8f, 0x99d0, 0x9841, 0x9607, 0x942c, 0x91ae, 0x8fa2, 0x8d14,
+};
+
+static inline HVX_Vector group_norm_lut32_f16(HVX_Vector index, HVX_Vector table) {
+    HVX_VectorPair values = Q6_Wh_vlut16_VbVhR(index, table, 0);
+    values = Q6_Wh_vlut16or_WhVbVhR(values, index, table, 1);
+    return Q6_V_lo_W(values);
+}
+
+static inline HVX_Vector group_norm_silu_f16(HVX_Vector x, HVX_Vector lut_a, HVX_Vector lut_b, HVX_Vector lut_c) {
+    const HVX_Vector scaled = hvx_vec_mul_f16_f16(hvx_vec_add_f16_f16(x, Q6_Vh_vsplat_R(0x4800)), Q6_Vh_vsplat_R(0x4000));
+    const HVX_Vector index_h = Q6_Vh_equals_Vhf(scaled);
+    const HVX_Vector fraction = hvx_vec_sub_f16_f16(scaled, Q6_Vhf_equals_Vh(index_h));
+    const HVX_Vector index = Q6_Vb_vshuffe_VbVb(index_h, index_h);
+    const HVX_Vector a = group_norm_lut32_f16(index, lut_a);
+    const HVX_Vector b = group_norm_lut32_f16(index, lut_b);
+    const HVX_Vector c = group_norm_lut32_f16(index, lut_c);
+    HVX_Vector y = hvx_vec_add_f16_f16(a, hvx_vec_mul_f16_f16(fraction, hvx_vec_add_f16_f16(b, hvx_vec_mul_f16_f16(fraction, c))));
+    y = Q6_V_vmux_QVV(Q6_Q_vcmp_gt_VhfVhf(Q6_Vh_vsplat_R(0x4800), x), y, x);
+    return Q6_V_vmux_QVV(Q6_Q_vcmp_gt_VhfVhf(Q6_Vh_vsplat_R(0xc800), x), Q6_V_vzero(), y);
+}
+
+static inline __attribute__((always_inline)) void group_norm_transform_f16(const struct group_norm_state * st, uint32_t channel0, uint32_t group_offset, const void * src_data, void * dst_data, uint32_t n, HVX_Vector mean, HVX_Vector inv_std) {
+    const __fp16 * src = (const __fp16 *) src_data;
+    __fp16 * dst = (__fp16 *) dst_data;
+    const bool affine = (st->flags & HTP_GROUP_NORM_AFFINE) != 0;
+    const bool silu = (st->flags & HTP_GROUP_NORM_SILU) != 0;
+    const HVX_Vector lut_a = Q6_Vh_vshuff_Vh(hvx_vmem(group_norm_silu_lut_a));
+    const HVX_Vector lut_b = Q6_Vh_vshuff_Vh(hvx_vmem(group_norm_silu_lut_b));
+    const HVX_Vector lut_c = Q6_Vh_vshuff_Vh(hvx_vmem(group_norm_silu_lut_c));
+    const float mean_f32 = hvx_vec_get_f32(mean);
+    const float inv_std_f32 = hvx_vec_get_f32(inv_std);
+    uint32_t done = 0;
+
+    while (done < n) {
+        const uint32_t offset = group_offset + done;
+        const uint32_t channel = channel0 + offset / st->spatial;
+        const uint32_t in_channel = offset % st->spatial;
+        const uint32_t count = hex_smin(st->spatial - in_channel, n - done);
+        const float weight = affine ? st->weight[channel] : 1.0f;
+        const HVX_Vector scale = hvx_vec_splat_f32(weight * inv_std_f32);
+        const HVX_Vector offset_vec = hvx_vec_splat_f32((affine ? st->bias[channel] : 0.0f) - mean_f32 * weight * inv_std_f32);
+        const uint32_t nvec = count / 64u;
+
+#pragma unroll(4)
+        for (uint32_t i = 0; i < nvec; ++i) {
+            const HVX_VectorPair p = hvx_vec_f16_to_f32(hvx_vmemu(src + done + (size_t) i * 64u));
+            const HVX_Vector y0 = hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(Q6_V_lo_W(p), scale), offset_vec);
+            const HVX_Vector y1 = hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(Q6_V_hi_W(p), scale), offset_vec);
+            HVX_Vector y = hvx_vec_f32_to_f16(y0, y1);
+            if (silu) {
+                y = group_norm_silu_f16(y, lut_a, lut_b, lut_c);
+            }
+            hvx_vmemu(dst + done + (size_t) i * 64u) = y;
+        }
+
+        const uint32_t tail = count % 64u;
+        if (tail) {
+            const HVX_VectorPair p = hvx_vec_f16_to_f32(hvx_vmemu(src + done + (size_t) nvec * 64u));
+            const HVX_Vector y0 = hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(Q6_V_lo_W(p), scale), offset_vec);
+            const HVX_Vector y1 = hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(Q6_V_hi_W(p), scale), offset_vec);
+            HVX_Vector y = hvx_vec_f32_to_f16(y0, y1);
+            if (silu) {
+                y = group_norm_silu_f16(y, lut_a, lut_b, lut_c);
+            }
+            hvx_vec_store_u(dst + done + (size_t) nvec * 64u, tail * sizeof(__fp16), y);
+        }
+        done += count;
+    }
+}
+
+static inline __attribute__((always_inline)) void group_norm_output_dma(const struct group_norm_state * st, dma_queue * queue, const void * src, void * dst, uint32_t channel0, uint32_t n, void * slots[2], uint32_t chunk, group_norm_transform_fn transform, HVX_Vector mean, HVX_Vector inv_std) {
+    const uint32_t chunks = (n + chunk - 1u) / chunk;
+    for (uint32_t i = 0; i < chunks && i < 2u; ++i) {
+        const uint32_t count = hex_smin(n - i * chunk, chunk);
+        group_norm_dma_push(queue, slots[i], (const uint8_t *) src + (size_t) i * chunk * st->element_size, count * st->element_size);
+    }
+
+    for (uint32_t i = 0; i < chunks; ++i) {
+        if (i >= 2u) {
+            (void) dma_queue_pop(queue);
+        }
+        const dma_ptr ready = dma_queue_pop(queue);
+        const uint32_t count = hex_smin(n - i * chunk, chunk);
+        transform(st, channel0, i * chunk, (const void *) ready.dst, (void *) ready.dst, count, mean, inv_std);
+        group_norm_dma_push(queue, (uint8_t *) dst + (size_t) i * chunk * st->element_size, (const void *) ready.dst, count * st->element_size);
+        if (i + 2u < chunks) {
+            const uint32_t next = i + 2u;
+            const uint32_t next_count = hex_smin(n - next * chunk, chunk);
+            group_norm_dma_push(queue, slots[i & 1u], (const uint8_t *) src + (size_t) next * chunk * st->element_size, next_count * st->element_size);
+        }
+    }
+    dma_queue_flush(queue);
+}
+
+#define GROUP_NORM_WORKER(name, type, chunk_expr, accumulate, transform) \
+static void name(unsigned int nth, unsigned int ith, void * data) { \
+    struct group_norm_state * st = (struct group_norm_state *) data; \
+    const uint32_t per_thread = (st->vtcm_size / nth) & ~2047u; \
+    const uint32_t slot_bytes = (per_thread / 2u) & ~2047u; \
+    const uint32_t chunk = (chunk_expr); \
+    uint8_t * thread_vtcm = st->vtcm + (size_t) ith * per_thread; \
+    void * slots[2] = { thread_vtcm, thread_vtcm + slot_bytes }; \
+    dma_queue * queue = st->ctx->dma[ith]; \
+    const uint32_t jobs = st->groups * st->batches; \
+    for (;;) { \
+        const uint32_t job = atomic_fetch_add_explicit(&st->next_job, 1u, memory_order_relaxed); \
+        if (job >= jobs) { \
+            break; \
+        } \
+        const uint32_t batch = job / st->groups; \
+        const uint32_t group = job - batch * st->groups; \
+        const uint32_t channel0 = group * st->channels_per_group; \
+        const uint32_t group_elems = st->channels_per_group * st->spatial; \
+        const size_t element_offset = (size_t) batch * st->channels * st->spatial + (size_t) channel0 * st->spatial; \
+        const type * group_src = (const type *) st->src + element_offset; \
+        type * group_dst = (type *) st->dst + element_offset; \
+        float mean; \
+        float mean_square; \
+        group_norm_stats_dma(st, queue, group_src, group_elems, slots, chunk, accumulate, &mean, &mean_square); \
+        const float variance = fmaxf(mean_square - mean * mean, 0.0f); \
+        const float inv_std = 1.0f / sqrtf(variance + st->epsilon); \
+        group_norm_output_dma(st, queue, group_src, group_dst, channel0, group_elems, slots, chunk, transform, hvx_vec_splat_f32(mean), hvx_vec_splat_f32(inv_std)); \
+    } \
+}
+
+GROUP_NORM_WORKER(group_norm_worker_f32, float, ((slot_bytes / sizeof(float)) / st->width) * st->width, group_norm_accumulate_f32, group_norm_transform_f32)
+GROUP_NORM_WORKER(group_norm_worker_f16, __fp16, (slot_bytes / sizeof(__fp16)) & ~63u, group_norm_accumulate_f16, group_norm_transform_f16)
+
+int op_group_norm(struct htp_ops_context * octx) {
+    const struct htp_tensor * src = octx->src[0];
+    const struct htp_tensor * dst = octx->dst;
+    const struct htp_group_norm_kernel_params * params = (const struct htp_group_norm_kernel_params *) octx->kernel_params;
+    const bool affine = (params->flags & HTP_GROUP_NORM_AFFINE) != 0;
+    const struct htp_tensor * weight = affine ? octx->src[1] : NULL;
+    const struct htp_tensor * bias = affine ? octx->src[2] : NULL;
+    const uint32_t groups = (uint32_t) octx->op_params[0];
+    const bool f16 = src && src->type == HTP_TYPE_F16;
+    float epsilon;
+    __builtin_memcpy(&epsilon, &octx->op_params[1], sizeof(epsilon));
+
+    if (!src || !dst || (!f16 && (params->flags & HTP_GROUP_NORM_SILU)) ||
+        !((src->type == HTP_TYPE_F32 && dst->type == HTP_TYPE_F32) || (f16 && dst->type == HTP_TYPE_F16)) ||
+        groups == 0 || src->ne[2] % groups != 0 || (!f16 && (src->ne[0] * src->ne[1]) % 32u != 0) ||
+        src->ne[0] != dst->ne[0] || src->ne[1] != dst->ne[1] ||
+        src->ne[2] != dst->ne[2] || src->ne[3] != dst->ne[3]) {
+        return HTP_STATUS_NO_SUPPORT;
+    }
+
+    const uint64_t weight_elems = weight ? (uint64_t) weight->ne[0] * weight->ne[1] * weight->ne[2] * weight->ne[3] : 0;
+    const uint64_t bias_elems = bias ? (uint64_t) bias->ne[0] * bias->ne[1] * bias->ne[2] * bias->ne[3] : 0;
+    if (affine && (!weight || !bias || weight->type != HTP_TYPE_F32 || bias->type != HTP_TYPE_F32 || weight_elems != src->ne[2] || bias_elems != src->ne[2])) {
+        return HTP_STATUS_NO_SUPPORT;
+    }
+    if (octx->flags & HTP_OPFLAGS_SKIP_COMPUTE) {
+        return HTP_STATUS_OK;
+    }
+
+    struct group_norm_state state = {
+        .ctx = octx->ctx,
+        .src = (const void *) src->data,
+        .weight = affine ? (const float *) weight->data : NULL,
+        .bias = affine ? (const float *) bias->data : NULL,
+        .dst = (void *) dst->data,
+        .spatial = src->ne[0] * src->ne[1],
+        .width = src->ne[0],
+        .channels = src->ne[2],
+        .groups = groups,
+        .channels_per_group = src->ne[2] / groups,
+        .batches = src->ne[3],
+        .element_size = f16 ? sizeof(__fp16) : sizeof(float),
+        .epsilon = epsilon,
+        .flags = params->flags,
+        .vtcm = (uint8_t *) octx->ctx->vtcm_base,
+        .vtcm_size = octx->ctx->vtcm_size,
+        .next_job = 0,
+    };
+    work_queue_run(octx->ctx->worker_pool, f16 ? group_norm_worker_f16 : group_norm_worker_f32, &state, octx->n_threads);
+    asm volatile("syncht" ::: "memory");
+    return HTP_STATUS_OK;
+}
diff --git a/ggml/src/ggml-hexagon/htp/groupnorm-ops.h b/ggml/src/ggml-hexagon/htp/groupnorm-ops.h
new file mode 100644
index 000000000000..55df834a495c
--- /dev/null
+++ b/ggml/src/ggml-hexagon/htp/groupnorm-ops.h
@@ -0,0 +1,16 @@
+#ifndef HTP_GROUPNORM_OPS_H
+#define HTP_GROUPNORM_OPS_H
+
+#include <stdint.h>
+
+enum htp_group_norm_flags {
+    HTP_GROUP_NORM_AFFINE = 1u << 0,
+    HTP_GROUP_NORM_SILU   = 1u << 1,
+};
+
+struct htp_group_norm_kernel_params {
+    uint32_t flags;
+};
+
+
+#endif
diff --git a/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h b/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
index 0011abba5a8a..2c225a6b06cd 100644
--- a/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
+++ b/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
@@ -676,6 +676,63 @@ static void core_dot_chunk_fp16(__fp16 *restrict output, const __fp16 *restrict
     }
 }
 
+static void core_dot_chunk_fp16_f8(__fp16 *restrict output, const __fp16 *restrict activation,
+                                   const uint8_t *restrict weight, const __fp16 *restrict scales,
+                                   uint32_t n_row_tiles, uint32_t n_col_tiles, uint32_t n_dot_tiles) {
+#if defined(__HEXAGON_ARCH__) && (__HEXAGON_ARCH__ >= 79)
+    __builtin_assume(n_row_tiles > 0);
+    __builtin_assume(n_col_tiles > 0);
+    __builtin_assume(n_dot_tiles > 0);
+
+    asm volatile(HMX_SET_BIAS("%0") :: "r"((unsigned int) scales));
+
+    const size_t act_dot_stride = (size_t) n_dot_tiles * HTP_MM_HMX_TILE_N_ELMS;
+    const size_t weight_dot_stride = (size_t) n_dot_tiles * HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3;
+
+    for (uint32_t r = 0; r < n_row_tiles; ++r) {
+        const __fp16 * row_base = activation + r * act_dot_stride;
+        const uint8_t * col_base = weight;
+        __fp16 * out_tile = output + (size_t) r * n_col_tiles * HTP_MM_HMX_TILE_N_ELMS;
+
+        for (uint32_t c = 0; c < n_col_tiles; ++c) {
+            const __fp16 * row_tiles = row_base;
+            const uint8_t * col_tiles = col_base;
+
+            asm volatile(HMX_CLRACC_F16());
+
+            for (uint32_t kt = 0; kt < n_dot_tiles; kt += 32) {
+                const uint32_t count = hex_smin(32, n_dot_tiles - kt);
+                const uint32_t act_range = count * HTP_MM_HMX_TILE_SIZE - 1;
+                const uint32_t weight_range = count * HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3 - 1;
+                asm volatile(
+                    "{\n"
+                    "    activation.hf = mxmem(%0, %1):deep\n"
+                    "    weight.f8 = mxmem(%2, %3)\n"
+                    "}\n"
+                    :: "r"(row_tiles), "r"(act_range), "r"(col_tiles), "r"(weight_range));
+                row_tiles += (size_t) count * HTP_MM_HMX_TILE_N_ELMS;
+                col_tiles += (size_t) count * HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3;
+            }
+
+            asm volatile(
+                "cvt.hf = acc(%0)\n"
+                "mxmem(%1, %2) = cvt\n"
+                :: "r"(2), "r"(out_tile), "r"(0) : "memory");
+            col_base += weight_dot_stride;
+            out_tile += HTP_MM_HMX_TILE_N_ELMS;
+        }
+    }
+#else
+    (void) output;
+    (void) activation;
+    (void) weight;
+    (void) scales;
+    (void) n_row_tiles;
+    (void) n_col_tiles;
+    (void) n_dot_tiles;
+#endif
+}
+
 static void core_mma_chunk_fp16_short(__fp16 *restrict c, const __fp16 *restrict a, const __fp16 *restrict b,
                                 const __fp16 *restrict col_scales, const __fp16 *restrict eye_tile,
                                 uint32_t n_row_tiles, uint32_t n_col_tiles, uint32_t n_dot_tiles, bool zero_init) {
@@ -777,13 +834,16 @@ static void transfer_output_chunk_fp16_to_fp32_col_chunk(
     uint32_t total_n_cols,
     uint32_t dst_stride,
     uint32_t src2_stride,
-    uint32_t dst_cols
+    uint32_t dst_cols,
+    float output_scale
 ) {
     assert(c_len % HTP_MM_HMX_TILE_N_COLS == 0);
     assert(total_n_cols % HTP_MM_HMX_TILE_N_COLS == 0);
     const size_t tile_row_stride = (total_n_cols / HTP_MM_HMX_TILE_N_COLS) * HTP_MM_HMX_TILE_N_ELMS;
 
     const HVX_Vector one = hvx_vec_splat_f16(1.0);
+    const HVX_Vector scale = hvx_vec_splat_f32(output_scale);
+    const bool apply_scale = output_scale != 1.0f;
 
     const size_t limit_c         = hex_smin(c_len, dst_cols);
     const size_t limit_c_aligned = (limit_c & ~31);
@@ -807,6 +867,9 @@ static void transfer_output_chunk_fp16_to_fp32_col_chunk(
             HVX_Vector *pv_out1 = (HVX_Vector *) (output_row_base + c + dst_stride);
 
             HVX_Vector v_out0 = Q6_Vsf_equals_Vqf32(Q6_V_lo_W(vp));
+            if (apply_scale) {
+                v_out0 = hvx_vec_mul_f32_f32(v_out0, scale);
+            }
             if (src2_row_base) {
                 HVX_Vector v_src2_0 = hvx_vmemu(src2_row_base + c + 0);
                 v_out0 = hvx_vec_add_f32_f32(v_out0, v_src2_0);
@@ -815,6 +878,9 @@ static void transfer_output_chunk_fp16_to_fp32_col_chunk(
 
             if (r + 1 < n_rows) {
                 HVX_Vector v_out1 = Q6_Vsf_equals_Vqf32(Q6_V_hi_W(vp));
+                if (apply_scale) {
+                    v_out1 = hvx_vec_mul_f32_f32(v_out1, scale);
+                }
                 if (src2_row_base) {
                     HVX_Vector v_src2_1 = hvx_vmemu(src2_row_base + c + src2_stride);
                     v_out1 = hvx_vec_add_f32_f32(v_out1, v_src2_1);
@@ -832,6 +898,9 @@ static void transfer_output_chunk_fp16_to_fp32_col_chunk(
             HVX_VectorPair vp = Q6_Wqf32_vmpy_VhfVhf(v, one);
 
             HVX_Vector v_out0 = Q6_Vsf_equals_Vqf32(Q6_V_lo_W(vp));
+            if (apply_scale) {
+                v_out0 = hvx_vec_mul_f32_f32(v_out0, scale);
+            }
             if (src2_row_base) {
                 HVX_Vector v_src2_0 = hvx_vmemu(src2_row_base + c + 0);
                 v_out0 = hvx_vec_add_f32_f32(v_out0, v_src2_0);
@@ -840,6 +909,9 @@ static void transfer_output_chunk_fp16_to_fp32_col_chunk(
 
             if (r + 1 < n_rows) {
                 HVX_Vector v_out1 = Q6_Vsf_equals_Vqf32(Q6_V_hi_W(vp));
+                if (apply_scale) {
+                    v_out1 = hvx_vec_mul_f32_f32(v_out1, scale);
+                }
                 if (src2_row_base) {
                     HVX_Vector v_src2_1 = hvx_vmemu(src2_row_base + c + src2_stride);
                     v_out1 = hvx_vec_add_f32_f32(v_out1, v_src2_1);
@@ -862,7 +934,7 @@ static inline void transfer_output_chunk_fp16_to_fp32(
     uint32_t dst_cols
 ) {
     transfer_output_chunk_fp16_to_fp32_col_chunk(
-        dst, src2, vtcm_src, start_row, n_rows, n_cols, n_cols, dst_stride, src2_stride, dst_cols
+        dst, src2, vtcm_src, start_row, n_rows, n_cols, n_cols, dst_stride, src2_stride, dst_cols, 1.0f
     );
 }
 
@@ -877,6 +949,7 @@ typedef struct {
     uint32_t       dst_stride;  // DDR row stride
     uint32_t       src2_stride; // DDR row stride for residual
     uint32_t       dst_cols;    // Actual output columns
+    float          output_scale;
     struct htp_thread_trace * traces;
 } output_transfer_task_state_t;
 
diff --git a/ggml/src/ggml-hexagon/htp/htp-ctx.h b/ggml/src/ggml-hexagon/htp/htp-ctx.h
index 3b60c8bdb08c..d564ccd89656 100644
--- a/ggml/src/ggml-hexagon/htp/htp-ctx.h
+++ b/ggml/src/ggml-hexagon/htp/htp-ctx.h
@@ -149,6 +149,7 @@ static inline void htp_ops_context_set_status(struct htp_ops_context * octx, int
 }
 
 int op_matmul(struct htp_ops_context * octx);
+int op_matmul_segmented(struct htp_ops_context * octx);
 int op_matmul_id(struct htp_ops_context * octx);
 int op_matmul_nx(struct htp_ops_context * octx);
 int op_matmul_id_nx(struct htp_ops_context * octx);
@@ -175,5 +176,8 @@ int op_gated_delta_net(struct htp_ops_context * octx);
 int op_pad(struct htp_ops_context * octx);
 int op_im2col(struct htp_ops_context * octx);
 int op_allreduce(struct htp_ops_context * octx);
+int op_qknorm_rope(struct htp_ops_context * octx);
+int op_conv2d(struct htp_ops_context * octx);
+int op_group_norm(struct htp_ops_context * octx);
 
 #endif /* HTP_CTX_H */
diff --git a/ggml/src/ggml-hexagon/htp/htp-ops.h b/ggml/src/ggml-hexagon/htp/htp-ops.h
index 869b19b8c2de..3f719f3e3eb9 100644
--- a/ggml/src/ggml-hexagon/htp/htp-ops.h
+++ b/ggml/src/ggml-hexagon/htp/htp-ops.h
@@ -26,12 +26,14 @@ enum htp_data_type {
     HTP_TYPE_I32    = 26,
     HTP_TYPE_I64    = 27,
     HTP_TYPE_MXFP4  = 39,
+    HTP_TYPE_F8_E4M3 = 43,
 
     // types used internally for repack, dyn.quant, etc
     HTP_TYPE_Q4_0_TILED = 200,
     HTP_TYPE_Q4_1_TILED,
     HTP_TYPE_Q8_0_TILED,
     HTP_TYPE_MXFP4_TILED,
+    HTP_TYPE_F8_E4M3_TILED,
 
     HTP_TYPE_INVALID
 };
@@ -40,6 +42,7 @@ enum htp_data_type {
 #define QK_Q4_0_TILED  256  // 32x32 Q4_0 tiled layout
 #define QK_Q8_0_TILED  128  // 32x32 Q8_0 tiled layout
 #define QK_MXFP4_TILED 256  // 32x32 MXFP4 tiled layout
+#define QK_F8_E4M3_TILED 1024 // 32x32 E4M3 tiled layout
 
 
 
@@ -102,6 +105,10 @@ enum htp_op_code {
     HTP_OP_ALLREDUCE_ADD,
     HTP_OP_GLU_SWIGLU_CLAMP,
     HTP_OP_MDEV_GROUP,
+    HTP_OP_QKNORM_ROPE,
+    HTP_OP_CONV_2D,
+    HTP_OP_GROUP_NORM,
+    HTP_OP_MUL_MAT_SEGMENTED,
 
     HTP_OP_INVALID
 };
diff --git a/ggml/src/ggml-hexagon/htp/hvx-sigmoid.h b/ggml/src/ggml-hexagon/htp/hvx-sigmoid.h
index 552017309d19..71c1902833e7 100644
--- a/ggml/src/ggml-hexagon/htp/hvx-sigmoid.h
+++ b/ggml/src/ggml-hexagon/htp/hvx-sigmoid.h
@@ -10,7 +10,44 @@
 #define FAST_SIGMOID_C2    (0x3e8d74bd)  // 0.276281267
 #define FAST_SIGMOID_C3    (0x3f000000)  // 0.5
 
+#if __HVX_ARCH__ >= 79
+static inline HVX_Vector hvx_vec_inverse_f32_v79(HVX_Vector v) {
+    const HVX_Vector two = hvx_vec_splat_f32(2.0f);
+    HVX_Vector r = Q6_Vw_vsub_VwVw(Q6_V_vsplat_R(0x7EEEEBB3), v);
+    r = Q6_Vsf_vmpy_VsfVsf(r, Q6_Vsf_vsub_VsfVsf(two, Q6_Vsf_vmpy_VsfVsf(r, v)));
+    return Q6_Vsf_vmpy_VsfVsf(r, Q6_Vsf_vsub_VsfVsf(two, Q6_Vsf_vmpy_VsfVsf(r, v)));
+}
+
+static inline HVX_Vector hvx_vec_fast_sigmoid_f32_v79(HVX_Vector v) {
+    v = Q6_Vsf_vmpy_VsfVsf(v, Q6_V_vsplat_R(FAST_SIGMOID_LOG2F));
+    v = Q6_Vsf_vmpy_VsfVsf(v, Q6_V_vsplat_R(FAST_SIGMOID_C3));
+
+    const HVX_Vector in_int = hvx_vec_truncate_f32(v);
+    const HVX_Vector x = Q6_Vsf_vsub_VsfVsf(v, Q6_Vsf_equals_Vw(in_int));
+    const HVX_Vector xx = Q6_Vsf_vmpy_VsfVsf(x, x);
+
+    HVX_Vector v1 = Q6_Vsf_vmpy_VsfVsf(xx, Q6_V_vsplat_R(FAST_SIGMOID_C2));
+    v1 = Q6_Vsf_vadd_VsfVsf(v1, Q6_V_vsplat_R(FAST_SIGMOID_LOG2F));
+
+    HVX_Vector v2 = Q6_Vsf_vmpy_VsfVsf(x, Q6_V_vsplat_R(FAST_SIGMOID_C1));
+    v2 = Q6_Vsf_vadd_VsfVsf(Q6_Vsf_vmpy_VsfVsf(v2, xx), x);
+
+    HVX_Vector v3 = Q6_Vsf_vadd_VsfVsf(v2, v1);
+    HVX_Vector v3_exponent = Q6_Vw_vasl_VwR(v3, 1);
+    v3_exponent = Q6_Vuw_vlsr_VuwR(v3_exponent, 24);
+    v3_exponent = Q6_Vw_vadd_VwVw(in_int, v3_exponent);
+    v3 = Q6_Vw_vaslacc_VwVwR(v3, in_int, 24);
+
+    const HVX_Vector v4 = Q6_Vsf_vsub_VsfVsf(v2, v1);
+    const HVX_Vector v5 = Q6_Vsf_vsub_VsfVsf(v3, v4);
+    return Q6_Vsf_vmpy_VsfVsf(v3, hvx_vec_inverse_f32_v79(v5));
+}
+#endif
+
 static inline HVX_Vector hvx_vec_fast_sigmoid_f32(HVX_Vector v) {
+#if __HVX_ARCH__ >= 79
+    return hvx_vec_fast_sigmoid_f32_v79(v);
+#else
     v = Q6_Vqf32_vmpy_VsfVsf(v, Q6_V_vsplat_R(FAST_SIGMOID_LOG2F));
     v = Q6_Vqf32_vmpy_VsfVsf(Q6_Vsf_equals_Vqf32(v), Q6_V_vsplat_R(FAST_SIGMOID_C3));
 
@@ -38,6 +75,7 @@ static inline HVX_Vector hvx_vec_fast_sigmoid_f32(HVX_Vector v) {
     res            = Q6_Vqf32_vmpy_VsfVsf(v3, res);
 
     return Q6_Vsf_equals_Vqf32(res);
+#endif
 }
 
 static inline HVX_Vector hvx_vec_fast_sigmoid_f32_guard(HVX_Vector v,
diff --git a/ggml/src/ggml-hexagon/htp/main.c b/ggml/src/ggml-hexagon/htp/main.c
index 1d291e16b463..a60a8daf2a34 100644
--- a/ggml/src/ggml-hexagon/htp/main.c
+++ b/ggml/src/ggml-hexagon/htp/main.c
@@ -796,6 +796,9 @@ static int execute_op(struct htp_ops_context * octx) {
         case HTP_OP_MUL_MAT_NX:
             return op_matmul_nx(octx);
 
+        case HTP_OP_MUL_MAT_SEGMENTED:
+            return op_matmul_segmented(octx);
+
         case HTP_OP_MUL:
         case HTP_OP_ADD:
         case HTP_OP_SUB:
@@ -885,6 +888,15 @@ static int execute_op(struct htp_ops_context * octx) {
         case HTP_OP_GATED_DELTA_NET:
             return op_gated_delta_net(octx);
 
+        case HTP_OP_QKNORM_ROPE:
+            return op_qknorm_rope(octx);
+
+        case HTP_OP_CONV_2D:
+            return op_conv2d(octx);
+
+        case HTP_OP_GROUP_NORM:
+            return op_group_norm(octx);
+
         case HTP_OP_TRI:
             return op_unary(octx);
 
diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.c b/ggml/src/ggml-hexagon/htp/matmul-ops.c
index 1b597dcd9f20..7efd195436c5 100644
--- a/ggml/src/ggml-hexagon/htp/matmul-ops.c
+++ b/ggml/src/ggml-hexagon/htp/matmul-ops.c
@@ -1784,7 +1784,7 @@ static void transfer_output_chunk_worker_fn(unsigned int n, unsigned int i, void
 
         float        *dst      = st->dst      + chunk_idx * st->dst_stride;
         const float  *src2     = st->src2     ? (st->src2     + chunk_idx * st->src2_stride) : NULL;
-        transfer_output_chunk_fp16_to_fp32(dst, src2, st->vtcm_src, chunk_idx, chunk_size, st->n_cols, st->dst_stride, st->src2_stride, st->dst_cols);
+        transfer_output_chunk_fp16_to_fp32_col_chunk(dst, src2, st->vtcm_src, chunk_idx, chunk_size, st->n_cols, st->n_cols, st->dst_stride, st->src2_stride, st->dst_cols, st->output_scale);
     }
 
     htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_O_PROC, start_chunk_idx);
@@ -2209,6 +2209,7 @@ typedef struct {
     uint32_t                  dst_stride;
     uint32_t                  src2_stride;
     uint32_t                  dst_cols;
+    float                     output_scale;
     struct fastdiv_values     n_threads_div;
 } output_transfer_col_chunk_state_t;
 
@@ -2236,15 +2237,15 @@ static void transfer_output_chunk_col_chunk_worker_fn(unsigned int n, unsigned i
     if (chunk_dst_cols > 0) {
         transfer_output_chunk_fp16_to_fp32_col_chunk(
             dst, src2, vtcm_src, 0, st->n_rows, c_len, st->n_cols,
-            st->dst_stride, st->src2_stride, (uint32_t)chunk_dst_cols
+            st->dst_stride, st->src2_stride, (uint32_t)chunk_dst_cols, st->output_scale
         );
     }
 
     htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_O_PROC, c_first);
 }
 
-static void transfer_output_chunk_threaded(struct htp_context *ctx, float *dst, const float *src2, const __fp16 *vtcm_src,
-                                              int n_rows, int n_cols, int dst_stride, uint32_t src2_stride, int dst_cols, int n_threads) {
+static void transfer_output_chunk_threaded_scaled(struct htp_context *ctx, float *dst, const float *src2, const __fp16 *vtcm_src,
+                                                   int n_rows, int n_cols, int dst_stride, uint32_t src2_stride, int dst_cols, int n_threads, float output_scale) {
     assert(n_cols % HTP_MM_HMX_TILE_N_COLS == 0);
 
     if (n_rows <= 0) return;
@@ -2261,6 +2262,7 @@ static void transfer_output_chunk_threaded(struct htp_context *ctx, float *dst,
         col_state.dst_stride = (uint32_t)dst_stride;
         col_state.src2_stride = src2_stride;
         col_state.dst_cols = (uint32_t)dst_cols;
+        col_state.output_scale = output_scale;
         col_state.n_threads_div = n_threads_div;
         col_state.traces = ctx->trace;
         col_state.ctx = ctx;
@@ -2286,6 +2288,7 @@ static void transfer_output_chunk_threaded(struct htp_context *ctx, float *dst,
     state.dst_stride        = dst_stride;
     state.src2_stride       = src2_stride;
     state.dst_cols          = dst_cols;
+    state.output_scale      = output_scale;
     state.traces            = ctx->trace;
 
     if (actual_threads <= 1) {
@@ -2295,6 +2298,11 @@ static void transfer_output_chunk_threaded(struct htp_context *ctx, float *dst,
     }
 }
 
+static void transfer_output_chunk_threaded(struct htp_context *ctx, float *dst, const float *src2, const __fp16 *vtcm_src,
+                                           int n_rows, int n_cols, int dst_stride, uint32_t src2_stride, int dst_cols, int n_threads) {
+    transfer_output_chunk_threaded_scaled(ctx, dst, src2, vtcm_src, n_rows, n_cols, dst_stride, src2_stride, dst_cols, n_threads, 1.0f);
+}
+
 struct activation_transfer_params {
     struct htp_context *          ctx;
     __fp16 *                      dst;
@@ -2403,32 +2411,121 @@ static void transfer_activation_chunk_threaded(const struct activation_transfer_
         worker_pool_run_func(ctx->worker_pool, transfer_activation_chunk_worker_fn, &state, active_threads);
     }
 }
+
+struct activation_segmented_transfer_state {
+    __fp16 * dst;
+    const float * src0;
+    const float * src1;
+    uint32_t n_rows;
+    uint32_t k;
+    uint32_t k0;
+    uint32_t stride0;
+    uint32_t stride1;
+    struct htp_thread_trace * traces;
+};
+
+static void transfer_activation_chunk_segmented_worker_fn(unsigned int n, unsigned int i, void * data) {
+    struct activation_segmented_transfer_state * st = data;
+    struct htp_thread_trace * tr = &st->traces[i];
+    const uint32_t n_rows_padded = hex_align_up(st->n_rows, HTP_MM_HMX_TILE_N_ROWS);
+    const uint32_t k1 = st->k - st->k0;
+
+    htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_A_PREP, i);
+    for (uint32_t r = 2 * i; r < n_rows_padded; r += 2 * n) {
+        const bool row0_valid = r < st->n_rows;
+        const bool row1_valid = r + 1 < st->n_rows;
+        const float * src00 = row0_valid ? st->src0 + r * st->stride0 : NULL;
+        const float * src01 = row1_valid ? st->src0 + (r + 1) * st->stride0 : NULL;
+        const float * src10 = row0_valid ? st->src1 + r * st->stride1 : NULL;
+        const float * src11 = row1_valid ? st->src1 + (r + 1) * st->stride1 : NULL;
+
+        transfer_activation_row_pair_fp32_to_fp16_col_chunk(
+            st->dst, src00, src01, r, st->k, 0, st->k0, st->k0, row0_valid, row1_valid);
+        transfer_activation_row_pair_fp32_to_fp16_col_chunk(
+            st->dst, src10, src11, r, st->k, st->k0, k1, k1, row0_valid, row1_valid);
+    }
+    htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_A_PREP, i);
+}
+
+static void transfer_activation_chunk_segmented_threaded(
+        struct htp_context * ctx,
+        __fp16 * dst,
+        const float * src0,
+        const float * src1,
+        uint32_t n_rows,
+        uint32_t k,
+        uint32_t k0,
+        uint32_t stride0,
+        uint32_t stride1,
+        uint32_t n_threads) {
+    assert(k0 < k);
+    assert(k % HTP_MM_HMX_TILE_N_COLS == 0);
+    assert(k0 % HTP_MM_HMX_TILE_N_COLS == 0);
+
+    struct activation_segmented_transfer_state state = {
+        .dst = dst,
+        .src0 = src0,
+        .src1 = src1,
+        .n_rows = n_rows,
+        .k = k,
+        .k0 = k0,
+        .stride0 = stride0,
+        .stride1 = stride1,
+        .traces = ctx->trace,
+    };
+    const uint32_t active_threads = hex_smin(n_threads, hex_align_up(n_rows, 2) / 2);
+    if (active_threads <= 1) {
+        transfer_activation_chunk_segmented_worker_fn(1, 0, &state);
+    } else {
+        worker_pool_run_func(ctx->worker_pool, transfer_activation_chunk_segmented_worker_fn, &state, active_threads);
+    }
+}
+
+static void transfer_matmul_activation_chunk(
+        const struct activation_transfer_params * params,
+        const float * src1,
+        int k0,
+        int stride1) {
+    if (src1) {
+        transfer_activation_chunk_segmented_threaded(
+            params->ctx, params->dst, params->src, src1, params->n_rows, params->k_block,
+            k0, params->k_stride, stride1, params->n_threads);
+    } else {
+        transfer_activation_chunk_threaded(params);
+    }
+}
 // --- Async HMX matmul job (for pipeline overlap) ---
 
 typedef struct {
     __fp16 *       output;
     const __fp16 * activation;
-    const __fp16 * weight;
+    const void *   weight;
     const __fp16 * scales;
     uint32_t       n_row_tiles;
     uint32_t       n_col_tiles;
     uint32_t       n_dot_tiles;
+    uint32_t       weight_type;
 } hmx_matmul_job_t;
 
 static void hmx_matmul_worker_fn(void * data) {
     hmx_matmul_job_t * job = (hmx_matmul_job_t *) data;
     FARF(HIGH, "hmx-mm-job: n_row_tiles %u n_col_tiles %u n_dot_tiles %u", job->n_row_tiles, job->n_col_tiles, job->n_dot_tiles);
-    core_dot_chunk_fp16(job->output, job->activation, job->weight, job->scales, job->n_row_tiles, job->n_col_tiles, job->n_dot_tiles);
+    if (job->weight_type == HTP_TYPE_F8_E4M3) {
+        core_dot_chunk_fp16_f8(job->output, job->activation, (const uint8_t *) job->weight, job->scales, job->n_row_tiles, job->n_col_tiles, job->n_dot_tiles);
+    } else {
+        core_dot_chunk_fp16(job->output, job->activation, job->weight, job->scales, job->n_row_tiles, job->n_col_tiles, job->n_dot_tiles);
+    }
 }
 
 static inline void hmx_matmul_job_init(hmx_matmul_job_t * job,
                                        __fp16 *           output,
                                        const __fp16 *     activation,
-                                       const __fp16 *     weight,
+                                       const void *       weight,
                                        const __fp16 *     scales,
                                        uint32_t           n_row_tiles,
                                        uint32_t           n_col_tiles,
-                                       uint32_t           n_dot_tiles) {
+                                       uint32_t           n_dot_tiles,
+                                       uint32_t           weight_type) {
     job->output      = output;
     job->activation  = activation;
     job->weight      = weight;
@@ -2436,6 +2533,7 @@ static inline void hmx_matmul_job_init(hmx_matmul_job_t * job,
     job->n_row_tiles = n_row_tiles;
     job->n_col_tiles = n_col_tiles;
     job->n_dot_tiles = n_dot_tiles;
+    job->weight_type = weight_type;
 }
 
 static int hmx_mm_2d_f32(struct htp_context *ctx,
@@ -2460,18 +2558,25 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
                                   const struct fastdiv_values * k_div,
                                   int tile_size,
                                   int aligned_tile_size,
-                                  int vtcm_size) {
+                                  int vtcm_size,
+                                  float output_scale,
+                                  const float * activation1,
+                                  int act_stride1,
+                                  int k0) {
     struct htp_thread_trace * tr = &ctx->trace[0];
     htp_trace_event_start(tr, HTP_TRACE_EVT_INIT, 0);
 
     if (k % 32 != 0 || n % 32 != 0) { return -1; }
-    if (!hex_is_aligned(dst, VLEN) || !hex_is_aligned(activation, VLEN)) { return -1; }
+    if (!hex_is_aligned(dst, VLEN) || !hex_is_aligned(activation, VLEN) ||
+        (activation1 && !hex_is_aligned(activation1, VLEN))) { return -1; }
+    if (activation1 && (k0 <= 0 || k0 >= k || k0 % HTP_MM_HMX_TILE_N_COLS != 0)) { return -1; }
 
     size_t row_stride = htp_mm_get_tiled_row_stride(weight_type, k);
     if (row_stride == 0) {
         return -1;
     }
 
+    const bool direct_f8 = weight_type == HTP_TYPE_F8_E4M3;
     worker_callback_t dequant_worker_fn = NULL;
     switch (weight_type) {
         case HTP_TYPE_Q4_0:   dequant_worker_fn = dequantize_tiled_worker_loop_q4_0; break;
@@ -2481,6 +2586,7 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
         case HTP_TYPE_Q8_0:   dequant_worker_fn = dequantize_tiled_worker_loop_q8_0; break;
         case HTP_TYPE_F16:    dequant_worker_fn = convert_f16_worker_loop; break;
         case HTP_TYPE_F32:    dequant_worker_fn = quantize_f32_worker_loop; break;
+        case HTP_TYPE_F8_E4M3: break;
         default:
             return -1;
     }
@@ -2559,7 +2665,9 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
                 .vtcm_f32_act = vtcm_f32_act,
                 .vtcm_f32_act_bytes = L.act_f32_bytes,
             };
-            transfer_activation_chunk_threaded(&act_params);
+            transfer_matmul_activation_chunk(&act_params,
+                                             activation1 ? activation1 + mr * act_stride1 : NULL,
+                                             k0, act_stride1);
 
             // Prologue: push A0 and optionally A1 (if n_chunk_cnt > 1)
             const size_t   n_cols_A0 = hex_smin(n - 0 * n_chunk_n_cols, n_chunk_n_cols);
@@ -2586,13 +2694,12 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
                 void * curr_raw = dma_queue_pop(ctx->dma[0]).dst;
 
                 // 2. dequantize A_i
-                dequantize_tiled_weight_chunk_to_fp16_tiles(
-                    ctx, vtcm_weight_bufs[i % 2], curr_raw,
-                    n_cols, k, row_stride, weight_type,
-                    n_k_tiles, n_k_tiles_div, dequant_worker_fn, n_threads);
+                if (!direct_f8) {
+                    dequantize_tiled_weight_chunk_to_fp16_tiles(ctx, vtcm_weight_bufs[i % 2], curr_raw, n_cols, k, row_stride, weight_type, n_k_tiles, n_k_tiles_div, dequant_worker_fn, n_threads);
+                }
 
                 // 3. push A_{i+2} (if i+2 < n_chunk_cnt)
-                if (i + 2 < n_chunk_cnt) {
+                if (!direct_f8 && i + 2 < n_chunk_cnt) {
                     const uint32_t height_p2 = is_quant ? (n_cols_p2 / 32) * n_k_tiles : n_cols_p2;
                     dma_queue_push(ctx->dma[0], dma_make_ptr(curr_raw, weight + nc_p2 * weight_stride),
                                    dma_dst_stride, dma_src_stride, dma_width_bytes, height_p2);
@@ -2600,21 +2707,27 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
 
                 // 4. submit C_i
                 hmx_matmul_job_init(&job_slots[i % 2], (__fp16 *) vtcm_output_bufs[i % 2],
-                                    (__fp16 *) vtcm_f16_act, (__fp16 *) vtcm_weight_bufs[i % 2],
+                                    (__fp16 *) vtcm_f16_act, (__fp16 *) (direct_f8 ? curr_raw : vtcm_weight_bufs[i % 2]),
                                     vtcm_scales, hmx_ceil_div(n_rows, HTP_MM_HMX_TILE_N_ROWS),
-                                    hmx_ceil_div(n_cols, HTP_MM_HMX_TILE_N_COLS), k / HTP_MM_HMX_TILE_N_ROWS);
+                                    hmx_ceil_div(n_cols, HTP_MM_HMX_TILE_N_COLS), k / HTP_MM_HMX_TILE_N_ROWS, weight_type);
                 hmx_queue_push(ctx->hmx_queue, hmx_queue_make_desc(hmx_matmul_worker_fn, &job_slots[i % 2]));
 
                 // 5. wait C_{i-1} and store D_{i-1} (multi-thread HVX, parallel with C_i)
                 if (i > 0) {
                     hmx_queue_pop(ctx->hmx_queue);
+                    if (direct_f8 && i + 1 < n_chunk_cnt) {
+                        const size_t nc_next = (i + 1) * n_chunk_n_cols;
+                        const size_t n_cols_next = hex_smin(n - nc_next, n_chunk_n_cols);
+                        const uint32_t height_next = (n_cols_next / 32) * n_k_tiles;
+                        dma_queue_push(ctx->dma[0], dma_make_ptr(vtcm_weight_raw[(i - 1) % 2], weight + nc_next * weight_stride), dma_dst_stride, dma_src_stride, dma_width_bytes, height_next);
+                    }
                     const size_t nc_prev = (i - 1) * n_chunk_n_cols;
                     const size_t n_cols_prev = hex_smin(n - nc_prev, n_chunk_n_cols);
                     float *output_chunk = dst + (mr * dst_stride + nc_prev);
                     const float *src2_chunk = src2 ? (src2 + mr * src2_stride + nc_prev) : NULL;
                     int chunk_dst_cols = dst_cols - (int)nc_prev;
                     if (chunk_dst_cols > 0) {
-                        transfer_output_chunk_threaded(ctx, output_chunk, src2_chunk, vtcm_output_bufs[(i - 1) % 2], n_rows, n_cols_prev, dst_stride, src2_stride, chunk_dst_cols, n_threads);
+                        transfer_output_chunk_threaded_scaled(ctx, output_chunk, src2_chunk, vtcm_output_bufs[(i - 1) % 2], n_rows, n_cols_prev, dst_stride, src2_stride, chunk_dst_cols, n_threads, output_scale);
                     }
                 }
             }
@@ -2627,7 +2740,7 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
             const float *src2_chunk = src2 ? (src2 + mr * src2_stride + nc_last) : NULL;
             int chunk_dst_cols = dst_cols - (int)nc_last;
             if (chunk_dst_cols > 0) {
-                transfer_output_chunk_threaded(ctx, output_chunk, src2_chunk, vtcm_output_bufs[(n_chunk_cnt - 1) % 2], n_rows, n_cols_last, dst_stride, src2_stride, chunk_dst_cols, n_threads);
+                transfer_output_chunk_threaded_scaled(ctx, output_chunk, src2_chunk, vtcm_output_bufs[(n_chunk_cnt - 1) % 2], n_rows, n_cols_last, dst_stride, src2_stride, chunk_dst_cols, n_threads, output_scale);
             }
         }
     } else {
@@ -2650,7 +2763,9 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
                 .vtcm_f32_act = vtcm_f32_act,
                 .vtcm_f32_act_bytes = L.act_f32_bytes,
             };
-            transfer_activation_chunk_threaded(&act_params);
+            transfer_matmul_activation_chunk(&act_params,
+                                             activation1 ? activation1 + mr * act_stride1 : NULL,
+                                             k0, act_stride1);
 
             // A0: Pre-fetch the first weight chunk (nc = 0)
             if (n > 0) {
@@ -2668,30 +2783,35 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
                 void * curr_raw = dma_queue_pop(ctx->dma[0]).dst;
 
                 // B: Weight Dequantize (Threaded)
-                dequantize_tiled_weight_chunk_to_fp16_tiles(
-                    ctx, vtcm_scratch0, curr_raw,
-                    n_cols, k, row_stride, weight_type,
-                    n_k_tiles, n_k_tiles_div, dequant_worker_fn, n_threads);
+                if (!direct_f8) {
+                    dequantize_tiled_weight_chunk_to_fp16_tiles(ctx, vtcm_scratch0, curr_raw, n_cols, k, row_stride, weight_type, n_k_tiles, n_k_tiles_div, dequant_worker_fn, n_threads);
+                }
 
                 // Start weight DMA for the next chunk early
                 const size_t nc_next = nc + n_chunk_n_cols;
-                if (nc_next < n) {
+                if (!direct_f8 && nc_next < n) {
                     const size_t n_cols_next = hex_smin(n - nc_next, n_chunk_n_cols);
                     const uint32_t height_next = is_quant ? (n_cols_next / 32) * n_k_tiles : n_cols_next;
                     dma_queue_push(ctx->dma[0], dma_make_ptr(curr_raw, weight + nc_next * weight_stride), dma_dst_stride, dma_src_stride, dma_width_bytes, height_next);
                 }
 
                 // C: HMX Compute (Queue-based)
-                hmx_matmul_job_init(&job, vtcm_output, vtcm_f16_act, vtcm_scratch0, vtcm_scales, n_row_tiles, n_col_tiles, k / HTP_MM_HMX_TILE_N_ROWS);
+                hmx_matmul_job_init(&job, vtcm_output, vtcm_f16_act, (__fp16 *) (direct_f8 ? curr_raw : vtcm_scratch0), vtcm_scales, n_row_tiles, n_col_tiles, k / HTP_MM_HMX_TILE_N_ROWS, weight_type);
                 hmx_queue_push(ctx->hmx_queue, hmx_queue_make_desc(hmx_matmul_worker_fn, &job));
                 hmx_queue_pop(ctx->hmx_queue);
 
+                if (direct_f8 && nc_next < n) {
+                    const size_t n_cols_next = hex_smin(n - nc_next, n_chunk_n_cols);
+                    const uint32_t height_next = (n_cols_next / 32) * n_k_tiles;
+                    dma_queue_push(ctx->dma[0], dma_make_ptr(curr_raw, weight + nc_next * weight_stride), dma_dst_stride, dma_src_stride, dma_width_bytes, height_next);
+                }
+
                 // D: Output Store
                 float *output_chunk = dst + (mr * dst_stride + nc);
                 const float *src2_chunk = src2 ? (src2 + mr * src2_stride + nc) : NULL;
                 int chunk_dst_cols = dst_cols - (int)nc;
                 if (chunk_dst_cols > 0) {
-                    transfer_output_chunk_threaded(ctx, output_chunk, src2_chunk, vtcm_output, n_rows, n_cols, dst_stride, src2_stride, chunk_dst_cols, n_threads);
+                    transfer_output_chunk_threaded_scaled(ctx, output_chunk, src2_chunk, vtcm_output, n_rows, n_cols, dst_stride, src2_stride, chunk_dst_cols, n_threads, output_scale);
                 }
             }
         }
@@ -2881,7 +3001,7 @@ static int hmx_mm_nx_2d_f32(struct htp_ops_context * octx, const struct htp_mm_k
                     hmx_matmul_job_init(&job_slots[i % 2], (__fp16 *) vtcm_output_bufs[i % 2],
                                         (__fp16 *) vtcm_f16_act, (__fp16 *) vtcm_weight_bufs[i % 2],
                                         vtcm_scales, hmx_ceil_div(n_rows, HTP_MM_HMX_TILE_N_ROWS),
-                                        hmx_ceil_div(n_cols, HTP_MM_HMX_TILE_N_COLS), k / HTP_MM_HMX_TILE_N_ROWS);
+                                        hmx_ceil_div(n_cols, HTP_MM_HMX_TILE_N_COLS), k / HTP_MM_HMX_TILE_N_ROWS, weight_type);
                     hmx_queue_push(ctx->hmx_queue, hmx_queue_make_desc(hmx_matmul_worker_fn, &job_slots[i % 2]));
 
                     if (i > 0) {
@@ -2967,7 +3087,7 @@ static int hmx_mm_nx_2d_f32(struct htp_ops_context * octx, const struct htp_mm_k
                         dma_queue_push(ctx->dma[0], dma_make_ptr(curr_raw, weight + nc_next * weight_stride), dma_dst_stride, dma_src_stride, dma_width_bytes, height_next);
                     }
 
-                    hmx_matmul_job_init(&job, vtcm_output, vtcm_f16_act, vtcm_scratch0, vtcm_scales, n_row_tiles, n_col_tiles, k / HTP_MM_HMX_TILE_N_ROWS);
+                    hmx_matmul_job_init(&job, vtcm_output, vtcm_f16_act, vtcm_scratch0, vtcm_scales, n_row_tiles, n_col_tiles, k / HTP_MM_HMX_TILE_N_ROWS, weight_type);
                     hmx_queue_push(ctx->hmx_queue, hmx_queue_make_desc(hmx_matmul_worker_fn, &job));
                     hmx_queue_pop(ctx->hmx_queue);
 
@@ -3029,7 +3149,7 @@ static int hmx_mm_f16_f32_batched_simple(struct htp_context *ctx,
                                            params->act_stride, params->weight_stride * (int)sizeof(__fp16),
                                            HTP_TYPE_F16, params->k, params->dst_stride, params->src2_stride, params->n,
                                            m_chunk, n_chunk, pipeline, n_threads, act_threads,
-                                           act_threads_div, k_div, 0, 0, vtcm_size);
+                                           act_threads_div, k_div, 0, 0, vtcm_size, 1.0f, NULL, 0, 0);
         }
     }
     return ret;
@@ -3170,7 +3290,7 @@ static int hmx_mm_f16_f32_batched(struct htp_context *ctx, const hmx_mm_f16_f32_
                     for (int g = 0; g < group_size; ++g) {
                         {
                             const __fp16 * vtcm_act_g = vtcm_f16_act + (size_t) g * L.act_head_stride;
-                            hmx_matmul_job_init(&job, vtcm_output, vtcm_act_g, vtcm_weight, vtcm_scales, n_row_tiles, n_col_tiles, params->k / 32);
+                            hmx_matmul_job_init(&job, vtcm_output, vtcm_act_g, vtcm_weight, vtcm_scales, n_row_tiles, n_col_tiles, params->k / 32, HTP_TYPE_F16);
                             hmx_queue_push(ctx->hmx_queue, hmx_queue_make_desc(hmx_matmul_worker_fn, &job));
                             hmx_queue_pop(ctx->hmx_queue);
                         }
@@ -3431,7 +3551,7 @@ static int hmx_mm_id_2d_f32(struct htp_context *ctx,
             }
 
             // C: HMX Compute (Queue-based)
-            hmx_matmul_job_init(&job, vtcm_output, vtcm_f16_act, vtcm_scratch0, vtcm_scales, n_row_tiles, n_col_tiles, k / HTP_MM_HMX_TILE_N_ROWS);
+            hmx_matmul_job_init(&job, vtcm_output, vtcm_f16_act, vtcm_scratch0, vtcm_scales, n_row_tiles, n_col_tiles, k / HTP_MM_HMX_TILE_N_ROWS, weight_type);
             hmx_queue_push(ctx->hmx_queue, hmx_queue_make_desc(hmx_matmul_worker_fn, &job));
             hmx_queue_pop(ctx->hmx_queue);
 
@@ -3447,6 +3567,34 @@ static int hmx_mm_id_2d_f32(struct htp_context *ctx,
 
 // --- Dispatchers and Public Entry Points ---
 
+static int hmx_mm_apply_output_scales(
+        const struct htp_ops_context * octx,
+        const struct htp_mm_kernel_params * kparams,
+        uint32_t scale_tensor_index,
+        float * output_scale) {
+    if (kparams->scale_flags & HTP_MM_SCALE_PARAM) {
+        float scale;
+        float bias;
+        memcpy(&scale, &octx->op_params[0], sizeof(scale));
+        memcpy(&bias, &octx->op_params[1], sizeof(bias));
+        if (bias != 0.0f) {
+            return HTP_STATUS_INVAL_PARAMS;
+        }
+        *output_scale *= scale;
+    }
+
+    if (kparams->scale_flags & HTP_MM_SCALE_TENSOR) {
+        const struct htp_tensor * scale = octx->src[scale_tensor_index];
+        if (!scale || scale->type != HTP_TYPE_F32 ||
+            scale->ne[0] * scale->ne[1] * scale->ne[2] * scale->ne[3] != 1) {
+            return HTP_STATUS_INVAL_PARAMS;
+        }
+        *output_scale *= *(const float *) scale->data;
+    }
+
+    return HTP_STATUS_OK;
+}
+
 static int hmx_mm_op_matmul(struct htp_ops_context * octx, const struct htp_mm_kernel_params * kparams) {
     htp_matmul_tensors_preamble;
 
@@ -3473,7 +3621,13 @@ static int hmx_mm_op_matmul(struct htp_ops_context * octx, const struct htp_mm_k
     uint32_t src2_stride = 0;
     size_t src2_nb2 = 0;
     size_t src2_nb3 = 0;
-    if (src2) {
+    float output_scale = src0->type == HTP_TYPE_F8_E4M3 ? 256.0f : 1.0f;
+    if (kparams->scale_flags != 0) {
+        const int status = hmx_mm_apply_output_scales(octx, kparams, 2, &output_scale);
+        if (status != HTP_STATUS_OK) {
+            return status;
+        }
+    } else if (src2) {
         src2_stride = (src2->ne[1] == 1) ? 0 : (uint32_t) (src2->nb[1] / sizeof(float));
         src2_ptr = (const float *) src2->data + m_start * src2_stride;
         src2_nb2 = (src2->ne[2] == 1) ? 0 : src2->nb[2];
@@ -3532,7 +3686,8 @@ static int hmx_mm_op_matmul(struct htp_ops_context * octx, const struct htp_mm_k
             kparams->n_act_threads,
             &kparams->div_n_act_threads,
             &kparams->div_ne00_padded,
-            kparams->tile_size, kparams->aligned_tile_size, kparams->vtcm_size
+            kparams->tile_size, kparams->aligned_tile_size, kparams->vtcm_size, output_scale,
+            NULL, 0, 0
         );
     }
 
@@ -3558,6 +3713,44 @@ int op_matmul(struct htp_ops_context * octx) {
     return hvx_mm_matmul(octx);
 }
 
+int op_matmul_segmented(struct htp_ops_context * octx) {
+    const struct htp_tensor * weight = octx->src[0];
+    const struct htp_tensor * src0 = octx->src[1];
+    const struct htp_tensor * src1 = octx->src[2];
+    const struct htp_tensor * dst = octx->dst;
+    const struct htp_mm_kernel_params * kparams = (const struct htp_mm_kernel_params *) octx->kernel_params;
+
+    if (!weight || !src0 || !src1 || !dst || !kparams->n_hmx ||
+        weight->type != HTP_TYPE_F8_E4M3 || src0->type != HTP_TYPE_F32 ||
+        src1->type != HTP_TYPE_F32 || dst->type != HTP_TYPE_F32) {
+        return HTP_STATUS_INVAL_PARAMS;
+    }
+
+    const int k = (int) weight->ne[0];
+    const int k0 = (int) src0->ne[0];
+    if (k0 + (int) src1->ne[0] != k || src0->ne[1] != src1->ne[1]) {
+        return HTP_STATUS_INVAL_PARAMS;
+    }
+
+    float output_scale = 256.0f;
+    const int scale_status = hmx_mm_apply_output_scales(octx, kparams, 3, &output_scale);
+    if (scale_status != HTP_STATUS_OK) {
+        return scale_status;
+    }
+
+    const int ret = hmx_mm_2d_f32(
+        octx->ctx, (float *) dst->data, NULL, (const float *) src0->data,
+        (const uint8_t *) weight->data, (int) src0->ne[1], k, (int) weight->ne[1],
+        (int) (src0->nb[1] / sizeof(float)), (int) weight->nb[1], (int) weight->type, k,
+        (int) (dst->nb[1] / sizeof(float)), 0, (int) dst->ne[0],
+        kparams->m_chunk, kparams->n_chunk, kparams->pipeline,
+        kparams->n_threads, kparams->n_act_threads,
+        &kparams->div_n_act_threads, &kparams->div_ne00_padded,
+        kparams->tile_size, kparams->aligned_tile_size, kparams->vtcm_size, output_scale,
+        (const float *) src1->data, (int) (src1->nb[1] / sizeof(float)), k0);
+    return ret == 0 ? HTP_STATUS_OK : HTP_STATUS_INTERNAL_ERR;
+}
+
 static int hmx_mm_op_matmul_id(
     struct htp_ops_context * octx,
     struct htp_mm_context * mmctx
diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.h b/ggml/src/ggml-hexagon/htp/matmul-ops.h
index 2dbcb0c2e51e..b78f2ef64d22 100644
--- a/ggml/src/ggml-hexagon/htp/matmul-ops.h
+++ b/ggml/src/ggml-hexagon/htp/matmul-ops.h
@@ -25,6 +25,7 @@ extern "C" {
 #define HTP_MM_WEIGHT_TILE_SIZE_Q8_0   1088
 #define HTP_MM_WEIGHT_TILE_SIZE_IQ4_NL 576
 #define HTP_MM_WEIGHT_TILE_SIZE_MXFP4  544
+#define HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3 1024
 
 // --- Weight Repacked Aligned Tile Sizes ---
 #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q4_0   640
@@ -32,6 +33,7 @@ extern "C" {
 #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q8_0   1152
 #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_IQ4_NL 640
 #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_MXFP4  640
+#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_F8_E4M3 1024
 
 // --- Activation Tiled Block Sizes (including padding) ---
 #define HTP_MM_ACT_TILE_SIZE_Q8_0      1152
@@ -69,6 +71,11 @@ enum htp_mm_kernel_type {
     HTP_MM_KERNEL_HVX_QUANT_ROW_FLAT, // row-wise fallback flat quantization
 };
 
+enum htp_mm_scale_flags {
+    HTP_MM_SCALE_PARAM  = 1u << 0,
+    HTP_MM_SCALE_TENSOR = 1u << 1,
+};
+
 // Op-specific struct for precomputed matmul params
 struct htp_mm_kernel_params {
     int32_t  kernel_type;        // enum htp_mm_kernel_type
@@ -88,7 +95,8 @@ struct htp_mm_kernel_params {
     int32_t  vtcm_src2_size;     // src2 scratchpad size in VTCM (fused only)
     int32_t  vtcm_src3_size;     // src3 scratchpad size in VTCM (fused only)
     int32_t  vtcm_dst_size;      // dst scratchpad size in VTCM
-    int32_t  n_weights;          // Number of weights for fused NX
+    uint16_t n_weights;          // Number of weights for fused NX
+    uint16_t scale_flags;        // enum htp_mm_scale_flags
 
     // Precomputed division values
     struct fastdiv_values div_ne12_ne1;
@@ -200,6 +208,8 @@ static inline uint32_t htp_mm_get_weight_tile_size(int weight_type) {
             return HTP_MM_WEIGHT_TILE_SIZE_Q8_0;
         case HTP_TYPE_MXFP4:
             return HTP_MM_WEIGHT_TILE_SIZE_MXFP4;
+        case HTP_TYPE_F8_E4M3:
+            return HTP_MM_WEIGHT_TILE_SIZE_F8_E4M3;
         default:
             return 0;
     }
@@ -216,6 +226,8 @@ static inline uint32_t htp_mm_get_weight_aligned_tile_size(int weight_type) {
             return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q8_0;
         case HTP_TYPE_MXFP4:
             return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_MXFP4;
+        case HTP_TYPE_F8_E4M3:
+            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_F8_E4M3;
         default:
             return 0;
     }
@@ -257,6 +269,8 @@ static inline size_t htp_mm_get_tiled_row_stride(int weight_type, uint32_t k) {
         case HTP_TYPE_Q8_0:
         case HTP_TYPE_MXFP4:
             return (size_t) nb * htp_mm_get_weight_tile_size(weight_type);
+        case HTP_TYPE_F8_E4M3:
+            return (size_t) k;
         case HTP_TYPE_F16:
             return (size_t) k * sizeof(__fp16);
         case HTP_TYPE_F32:
@@ -279,13 +293,14 @@ static inline void htp_mm_hmx_get_2d_chunk_costs(
     size_t * size_per_n_out, size_t * size_per_m_out, size_t * size_per_mn_out
 ) {
     const bool is_quant = (wtype != HTP_TYPE_F16 && wtype != HTP_TYPE_F32);
+    const bool direct_f8 = (wtype == HTP_TYPE_F8_E4M3);
     const size_t row_stride = htp_mm_get_tiled_row_stride(wtype, k);
     const size_t vec_dot_size = k * sizeof(uint16_t);
     const uint32_t n_k_tiles = k / HTP_MM_HMX_TILE_N_COLS;
     const size_t qweight_row_stride = is_quant ? (size_t)(n_k_tiles * aligned_tile_size) / 32 : 0;
 
     *size_per_n_out = (pipeline ? 2 : 1) * (is_quant ? qweight_row_stride : row_stride) +
-                      (pipeline ? 2 * vec_dot_size : vec_dot_size);
+                      (direct_f8 ? 0 : (pipeline ? 2 * vec_dot_size : vec_dot_size));
     *size_per_m_out = vec_dot_size;
     *size_per_mn_out = (pipeline ? 2 : 1) * sizeof(uint16_t);
 }
@@ -409,6 +424,7 @@ static inline void htp_mm_hmx_vtcm_layout_build(
     } else {
         // HTP_MM_KERNEL_HMX_2D
         const bool is_quant = (wtype != HTP_TYPE_F16 && wtype != HTP_TYPE_F32);
+        const bool direct_f8 = (wtype == HTP_TYPE_F8_E4M3);
         const size_t row_stride = htp_mm_get_tiled_row_stride(wtype, k);
         const size_t vec_dot_size = k * sizeof(uint16_t);
         const uint32_t n_k_tiles = k / HTP_MM_HMX_TILE_N_COLS;
@@ -420,7 +436,7 @@ static inline void htp_mm_hmx_vtcm_layout_build(
         const size_t act_area_size    = hex_align_up(mc * vec_dot_size, HTP_MM_HMX_TILE_SIZE);
         const size_t output_area_size = hex_align_up(mc * nc * sizeof(__fp16), HTP_MM_HMX_TILE_SIZE);
 
-        const size_t scratch0_size = hex_align_up(nc * vec_dot_size, HTP_MM_HMX_TILE_SIZE);
+        const size_t scratch0_size = direct_f8 ? 0 : hex_align_up(nc * vec_dot_size, HTP_MM_HMX_TILE_SIZE);
         const size_t scratch1_size = pipeline ? scratch0_size : 0;
 
         // Group A:  Scales and activation tiles (must not overlap with Group B or C)
diff --git a/ggml/src/ggml-hexagon/htp/qknorm-rope-ops.c b/ggml/src/ggml-hexagon/htp/qknorm-rope-ops.c
new file mode 100644
index 000000000000..292f935fd88b
--- /dev/null
+++ b/ggml/src/ggml-hexagon/htp/qknorm-rope-ops.c
@@ -0,0 +1,136 @@
+#include <HAP_farf.h>
+
+#include <math.h>
+#include <stdint.h>
+#include <string.h>
+
+#include "htp-ctx.h"
+#include "hvx-norm.h"
+
+struct qknorm_rope_context {
+    struct htp_context * ctx;
+    const float * src;
+    const float * weight;
+    const float * theta;
+    float * dst;
+    uint32_t head_dim;
+    uint32_t n_heads;
+    uint32_t n_tokens;
+    uint32_t n_batches;
+    uint32_t src_head_stride;
+    uint32_t src_token_stride;
+    uint32_t src_batch_stride;
+    float epsilon;
+    uint32_t n_threads;
+};
+
+static inline void qknorm_rope_pack_theta(float * restrict packed, const float * restrict src, uint32_t head_dim) {
+    for (uint32_t i = 0; i < head_dim / 2; ++i) {
+        packed[2 * i + 0] = src[4 * i + 0];
+        packed[2 * i + 1] = src[4 * i + 2];
+    }
+}
+
+static inline void qknorm_rope_apply_interleaved(float * restrict dst, const float * restrict src,
+                                                  const float * restrict theta, uint32_t head_dim) {
+    const HVX_Vector * vx = (const HVX_Vector *) src;
+    const HVX_Vector * vt = (const HVX_Vector *) theta;
+    HVX_Vector * vo = (HVX_Vector *) dst;
+
+    for (uint32_t i = 0; i < head_dim / 64; ++i) {
+        const HVX_Vector x0 = *vx++;
+        const HVX_Vector x1 = *vx++;
+        const HVX_Vector t0 = *vt++;
+        const HVX_Vector t1 = *vt++;
+        const HVX_VectorPair x_even_odd = Q6_W_vdeal_VVR(x1, x0, -4);
+        const HVX_VectorPair cos_sin = Q6_W_vdeal_VVR(t1, t0, -4);
+
+        const HVX_Vector x0_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(x_even_odd), Q6_V_lo_W(cos_sin));
+        const HVX_Vector x0_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(x_even_odd), Q6_V_hi_W(cos_sin));
+        const HVX_Vector x1_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(x_even_odd), Q6_V_lo_W(cos_sin));
+        const HVX_Vector x1_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(x_even_odd), Q6_V_hi_W(cos_sin));
+        const HVX_Vector even = Q6_Vqf32_vsub_Vqf32Vqf32(x0_c, x1_s);
+        const HVX_Vector odd = Q6_Vqf32_vadd_Vqf32Vqf32(x0_s, x1_c);
+        const HVX_VectorPair out = Q6_W_vshuff_VVR(Q6_Vsf_equals_Vqf32(odd), Q6_Vsf_equals_Vqf32(even), -4);
+
+        *vo++ = Q6_V_lo_W(out);
+        *vo++ = Q6_V_hi_W(out);
+    }
+}
+
+static void qknorm_rope_worker(unsigned int n, unsigned int i, void * data) {
+    struct qknorm_rope_context * c = (struct qknorm_rope_context *) data;
+    struct htp_thread_trace * tr = &c->ctx->trace[i];
+    const uint32_t token_first = (uint32_t) (((uint64_t) c->n_tokens * i) / n);
+    const uint32_t token_last = (uint32_t) (((uint64_t) c->n_tokens * (i + 1)) / n);
+    float theta_packed[256] __attribute__((aligned(128)));
+    float normalized[256] __attribute__((aligned(128)));
+
+    htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_COMP, token_first);
+    for (uint32_t token = token_first; token < token_last; ++token) {
+        qknorm_rope_pack_theta(theta_packed, c->theta + (size_t) token * 2 * c->head_dim, c->head_dim);
+        for (uint32_t batch = 0; batch < c->n_batches; ++batch) {
+            for (uint32_t head = 0; head < c->n_heads; ++head) {
+                const float * src = c->src + (size_t) batch * c->src_batch_stride +
+                                    (size_t) token * c->src_token_stride +
+                                    (size_t) head * c->src_head_stride;
+                float * dst = c->dst + (((size_t) batch * c->n_heads + head) * c->n_tokens + token) * c->head_dim;
+                hvx_fast_rms_norm_mul_f32((const uint8_t *) src, (const uint8_t *) c->weight,
+                                           (uint8_t *) normalized, c->head_dim, c->epsilon);
+                qknorm_rope_apply_interleaved(dst, normalized, theta_packed, c->head_dim);
+            }
+        }
+    }
+    htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_COMP, token_first);
+}
+
+int op_qknorm_rope(struct htp_ops_context * octx) {
+    const struct htp_tensor * src = octx->src[0];
+    const struct htp_tensor * weight = octx->src[1];
+    const struct htp_tensor * theta = octx->src[2];
+    const struct htp_tensor * dst = octx->dst;
+
+    if (!src || !weight || !theta || !dst || src->type != HTP_TYPE_F32 || weight->type != HTP_TYPE_F32 ||
+        theta->type != HTP_TYPE_F32 || dst->type != HTP_TYPE_F32) {
+        return HTP_STATUS_NO_SUPPORT;
+    }
+
+    const uint32_t head_dim = src->ne[0];
+    const uint32_t n_heads = src->ne[1];
+    const uint32_t n_tokens = src->ne[2];
+    const uint32_t n_batches = src->ne[3];
+    const uint64_t output_elements = (uint64_t) head_dim * n_heads * n_tokens * n_batches;
+    const uint64_t weight_elements = (uint64_t) weight->ne[0] * weight->ne[1] * weight->ne[2] * weight->ne[3];
+    const uint64_t theta_elements = (uint64_t) theta->ne[0] * theta->ne[1] * theta->ne[2] * theta->ne[3];
+    const uint64_t dst_elements = (uint64_t) dst->ne[0] * dst->ne[1] * dst->ne[2] * dst->ne[3];
+
+    if (head_dim == 0 || head_dim > 256 || head_dim % 64 != 0 || n_heads == 0 || n_tokens == 0 || n_batches == 0 ||
+        weight_elements != head_dim || theta_elements < (uint64_t) n_tokens * 2 * head_dim ||
+        dst_elements != output_elements || src->nb[0] != sizeof(float) ||
+        src->nb[1] % sizeof(float) || src->nb[2] % sizeof(float) || src->nb[3] % sizeof(float)) {
+        return HTP_STATUS_INVAL_PARAMS;
+    }
+
+    float epsilon = 0.0f;
+    memcpy(&epsilon, &octx->op_params[0], sizeof(epsilon));
+    const uint32_t n_threads = n_tokens < octx->n_threads ? n_tokens : octx->n_threads;
+    struct qknorm_rope_context c = {
+        .ctx = octx->ctx,
+        .src = (const float *) src->data,
+        .weight = (const float *) weight->data,
+        .theta = (const float *) theta->data,
+        .dst = (float *) dst->data,
+        .head_dim = head_dim,
+        .n_heads = n_heads,
+        .n_tokens = n_tokens,
+        .n_batches = n_batches,
+        .src_head_stride = src->nb[1] / sizeof(float),
+        .src_token_stride = src->nb[2] / sizeof(float),
+        .src_batch_stride = src->nb[3] / sizeof(float),
+        .epsilon = epsilon,
+        .n_threads = n_threads,
+    };
+
+    work_queue_run(octx->ctx->work_queue, qknorm_rope_worker, &c, n_threads);
+    return HTP_STATUS_OK;
+}
diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c
index 1ebc50a763f1..505f621575e9 100644
--- a/ggml/src/ggml-quants.c
+++ b/ggml/src/ggml-quants.c
@@ -25,6 +25,27 @@
 
 #define UNUSED GGML_UNUSED
 
+static float fp8_e4m3_to_fp32(uint8_t x) {
+    const uint32_t sign     = x >> 7;
+    const uint32_t exponent = (x >> 3) & 0x0f;
+    const uint32_t mantissa = x & 0x07;
+
+    if (exponent == 0x0f && mantissa == 0x07) {
+        return sign ? -NAN : NAN;
+    }
+
+    const float value = exponent == 0
+                            ? ldexpf((float) mantissa, -9)
+                            : ldexpf(1.0f + (float) mantissa / 8.0f, (int) exponent - 7);
+    return sign ? -value : value;
+}
+
+void dequantize_row_f8_e4m3(const uint8_t * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
+    for (int64_t i = 0; i < k; ++i) {
+        y[i] = fp8_e4m3_to_fp32(x[i]);
+    }
+}
+
 static inline int best_index_int8(int n, const int8_t * val, float x) {
     if (x <= val[0]) return 0;
     if (x >= val[n-1]) return n-1;
diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h
index 75188f1af180..daeae35f6ff1 100644
--- a/ggml/src/ggml-quants.h
+++ b/ggml/src/ggml-quants.h
@@ -43,6 +43,7 @@ GGML_API void quantize_row_iq3_s_ref  (const float * GGML_RESTRICT x, block_iq3_
 GGML_API void quantize_row_iq2_s_ref  (const float * GGML_RESTRICT x, block_iq2_s   * GGML_RESTRICT y, int64_t k);
 
 // Dequantization
+GGML_API void dequantize_row_f8_e4m3(const uint8_t * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_q1_0(const block_q1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_q2_0(const block_q2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_q4_0(const block_q4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c
index 5ef03e190e34..2c8cbe93f502 100644
--- a/ggml/src/ggml.c
+++ b/ggml/src/ggml.c
@@ -943,6 +943,13 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = {
         .type_size                = 0,
         .is_quantized             = false,
     },
+    [GGML_TYPE_F8_E4M3] = {
+        .type_name                = "f8_e4m3",
+        .blck_size                = 1,
+        .type_size                = sizeof(uint8_t),
+        .is_quantized             = false,
+        .to_float                 = (ggml_to_float_t) dequantize_row_f8_e4m3,
+    },
 };
 
 const struct ggml_type_traits * ggml_get_type_traits(enum ggml_type type) {
@@ -1099,9 +1106,15 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = {
     "OPT_STEP_SGD",
 
     "GLU",
+
+    "MUL_MAT_SEGMENTED",
+    "QKNORM_ROPE",
+    "GROUP_NORM_AFFINE_SILU",
+    "CONV_2D_BIAS",
+    "CONV_2D_UPSCALE",
 };
 
-static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101");
+static_assert(GGML_OP_COUNT == 106, "GGML_OP_COUNT != 106");
 
 static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
     "none",
@@ -1214,9 +1227,15 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = {
     "sgd(x)",
 
     "glu(x)",
+
+    "X*concat(Y0,Y1)",
+    "qknorm_rope(x,w,theta)",
+    "silu(group_norm(x)*w+b)",
+    "conv2d(x)+b",
+    "conv2d(upscale(x))",
 };
 
-static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101");
+static_assert(GGML_OP_COUNT == 106, "GGML_OP_COUNT != 106");
 
 static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2");
 
@@ -3193,6 +3212,30 @@ struct ggml_tensor * ggml_rms_norm_inplace(
     return ggml_rms_norm_impl(ctx, a, eps, true);
 }
 
+struct ggml_tensor * ggml_qknorm_rope(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * a,
+        struct ggml_tensor  * weight,
+        struct ggml_tensor  * theta,
+        float                 eps) {
+    GGML_ASSERT(a->type == GGML_TYPE_F32);
+    GGML_ASSERT(weight->type == GGML_TYPE_F32);
+    GGML_ASSERT(theta->type == GGML_TYPE_F32);
+    GGML_ASSERT(ggml_nelements(weight) == a->ne[0]);
+    GGML_ASSERT(ggml_nelements(theta) >= a->ne[2] * 2 * a->ne[0]);
+
+    const int64_t ne[3] = { a->ne[0], a->ne[2], a->ne[1] * a->ne[3] };
+    struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 3, ne);
+
+    ggml_set_op_params(result, &eps, sizeof(eps));
+    result->op     = GGML_OP_QKNORM_ROPE;
+    result->src[0] = a;
+    result->src[1] = weight;
+    result->src[2] = theta;
+
+    return result;
+}
+
 // ggml_rms_norm_back
 
 struct ggml_tensor * ggml_rms_norm_back(
@@ -3246,6 +3289,52 @@ struct ggml_tensor * ggml_group_norm_inplace(
     return ggml_group_norm_impl(ctx, a, n_groups, eps, true);
 }
 
+static struct ggml_tensor * ggml_group_norm_affine_silu_impl(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * a,
+        struct ggml_tensor  * weight,
+        struct ggml_tensor  * bias,
+        int                   n_groups,
+        float                 eps,
+        bool                  inplace) {
+    GGML_ASSERT(weight->type == GGML_TYPE_F32);
+    GGML_ASSERT(bias->type == GGML_TYPE_F32);
+    GGML_ASSERT(ggml_nelements(weight) == a->ne[2]);
+    GGML_ASSERT(ggml_nelements(bias) == a->ne[2]);
+
+    struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
+
+    ggml_set_op_params_i32(result, 0, n_groups);
+    ggml_set_op_params_f32(result, 1, eps);
+
+    result->op     = GGML_OP_GROUP_NORM_AFFINE_SILU;
+    result->src[0] = a;
+    result->src[1] = weight;
+    result->src[2] = bias;
+
+    return result;
+}
+
+struct ggml_tensor * ggml_group_norm_affine_silu(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * a,
+        struct ggml_tensor  * weight,
+        struct ggml_tensor  * bias,
+        int                   n_groups,
+        float                 eps) {
+    return ggml_group_norm_affine_silu_impl(ctx, a, weight, bias, n_groups, eps, false);
+}
+
+struct ggml_tensor * ggml_group_norm_affine_silu_inplace(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * a,
+        struct ggml_tensor  * weight,
+        struct ggml_tensor  * bias,
+        int                   n_groups,
+        float                 eps) {
+    return ggml_group_norm_affine_silu_impl(ctx, a, weight, bias, n_groups, eps, true);
+}
+
 // ggml_l2_norm
 
 static struct ggml_tensor * ggml_l2_norm_impl(
@@ -3355,10 +3444,34 @@ struct ggml_tensor * ggml_mul_mat(
     return result;
 }
 
+struct ggml_tensor * ggml_mul_mat_segmented(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * a,
+        struct ggml_tensor  * b0,
+        struct ggml_tensor  * b1) {
+    GGML_ASSERT(a->ne[0] == b0->ne[0] + b1->ne[0]);
+    GGML_ASSERT(b0->ne[1] == b1->ne[1]);
+    GGML_ASSERT(b0->ne[2] == b1->ne[2]);
+    GGML_ASSERT(b0->ne[3] == b1->ne[3]);
+    GGML_ASSERT(b0->ne[2] % a->ne[2] == 0);
+    GGML_ASSERT(b0->ne[3] % a->ne[3] == 0);
+    GGML_ASSERT(!ggml_is_transposed(a));
+
+    const int64_t ne[4] = { a->ne[1], b0->ne[1], b0->ne[2], b0->ne[3] };
+    struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne);
+
+    result->op     = GGML_OP_MUL_MAT_SEGMENTED;
+    result->src[0] = a;
+    result->src[1] = b0;
+    result->src[2] = b1;
+
+    return result;
+}
+
 void ggml_mul_mat_set_prec(
         struct ggml_tensor * a,
         enum ggml_prec       prec) {
-    GGML_ASSERT(a->op == GGML_OP_MUL_MAT);
+    GGML_ASSERT(a->op == GGML_OP_MUL_MAT || a->op == GGML_OP_MUL_MAT_SEGMENTED);
 
     const int32_t prec_i32 = (int32_t) prec;
 
@@ -4980,6 +5093,58 @@ struct ggml_tensor * ggml_conv_2d_direct(
     return result;
 }
 
+struct ggml_tensor * ggml_conv_2d_direct_bias(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * a,
+        struct ggml_tensor  * b,
+        struct ggml_tensor  * bias,
+        int s0, int s1, int p0, int p1, int d0, int d1) {
+    GGML_ASSERT(bias != NULL);
+    GGML_ASSERT(bias->type == GGML_TYPE_F32);
+    GGML_ASSERT(ggml_nelements(bias) == a->ne[3]);
+
+    struct ggml_tensor * result = ggml_conv_2d_direct(ctx, a, b, s0, s1, p0, p1, d0, d1);
+    result->op     = GGML_OP_CONV_2D_BIAS;
+    result->src[2] = bias;
+    return result;
+}
+
+struct ggml_tensor * ggml_conv_2d_direct_upscale(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * a,
+        struct ggml_tensor  * b,
+        struct ggml_tensor  * bias,
+        int upscale_factor,
+        int s0, int s1, int p0, int p1, int d0, int d1) {
+    GGML_ASSERT(upscale_factor > 1);
+    GGML_ASSERT(a->ne[2] == b->ne[2]);
+    GGML_ASSERT(bias == NULL || bias->type == GGML_TYPE_F32);
+    GGML_ASSERT(bias == NULL || ggml_nelements(bias) == a->ne[3]);
+
+    const int64_t ne[4] = {
+        ggml_calc_conv_output_size(b->ne[0] * upscale_factor, a->ne[0], s0, p0, d0),
+        ggml_calc_conv_output_size(b->ne[1] * upscale_factor, a->ne[1], s1, p1, d1),
+        a->ne[3],
+        b->ne[3],
+    };
+    struct ggml_tensor * result = ggml_new_tensor(ctx, b->type, 4, ne);
+
+    ggml_set_op_params_i32(result, 0, s0);
+    ggml_set_op_params_i32(result, 1, s1);
+    ggml_set_op_params_i32(result, 2, p0);
+    ggml_set_op_params_i32(result, 3, p1);
+    ggml_set_op_params_i32(result, 4, d0);
+    ggml_set_op_params_i32(result, 5, d1);
+    ggml_set_op_params_i32(result, 6, upscale_factor);
+
+    result->op     = GGML_OP_CONV_2D_UPSCALE;
+    result->src[0] = a;
+    result->src[1] = b;
+    result->src[2] = bias;
+
+    return result;
+}
+
 // ggml_conv_3d_direct
 
 struct ggml_tensor * ggml_conv_3d_direct(