opencl: add depthwise conv2d (#30900)

* opencl: add Adreno depthwise conv2d Assisted-by: OpenCode * opencl: use local program for depthwise conv2d Assisted-by: OpenCode * opencl: use singed type for args --------- Co-authored-by: Li He <lih@qti.qualcomm.com>
2026-10-22 10:54:24 -07:00
parent a08ab75ae7
commit 970b51ce7c
3 changed files with 253 additions and 0 deletions
@@ -241,6 +241,7 @@ set(GGML_OPENCL_KERNELS
    mul_mm_f16_f32_kq_kqv
    conv2d
    conv2d_f16_f32
+    conv2_dw
    flash_attn_repack
    flash_attn_pre_f16
    flash_attn_f32_f16
@@ -1005,6 +1005,7 @@ struct ggml_backend_opencl_context {
    cl_kernel kernel_conv_2d_f16;
    cl_kernel kernel_conv_2d_f32;
    cl_kernel kernel_conv_2d_f16_f32;
+    cl_kernel kernel_conv2d_dw_f32_f32, kernel_conv2d_dwf16_f32;
    cl_kernel kernel_ssm_conv_f32_f32, kernel_ssm_conv_f32_f32_4;
    // [size_idx][kda][tgpp] where size_idx: 0=S_V=16, 1=32, 2=64, 3=128; kda: 0 or 1.
    // tgpp 0 = TG variant (COLS_PER_LANE_GROUP=1), tgpp 1 = prefill variant (COLS_PER_LANE_GROUP=4).
@@ -3761,6 +3762,23 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
                }
    }

+    // conv2_dw
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "conv2_dw.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("conv2_dw.cl");
+#endif
+        cl_program prog =
+            build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
+        CL_CHECK((backend_ctx->kernel_conv2d_dw_f32_f32 = clCreateKernel(prog, "kernel_conv2d_dw_f32_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_conv2d_dwf16_f32 = clCreateKernel(prog, "kernel_conv2d_dwf16_f32", &err), err));
+        CL_CHECK(clReleaseProgram(prog));
+        GGML_LOG_CONT(".");
+    }
+
    // ssm_conv
    {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -9572,6 +9590,11 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
            return (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F16 && op->type == GGML_TYPE_F16) ||
                   (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) ||
                   (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
+        case GGML_OP_CONV_2D_DW:
+            return op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&
+                   op->src[0]->nb[0] == ggml_type_size(op->src[0]->type) &&
+                   op->src[1]->nb[0] == sizeof(float) && op->nb[0] == sizeof(float) &&
+                   (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16);
        case GGML_OP_SSM_CONV:
            return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
        case GGML_OP_SSM_SCAN: {
@@ -19715,6 +19738,90 @@ static void ggml_cl_conv_2d(ggml_backend_t backend, const ggml_tensor * src0, co
    backend_ctx->enqueue_ndrange_kernel(kernel, 2, global_work_size, local_work_size, dst);
 }

+static void ggml_cl_conv2_dw(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    GGML_TENSOR_BINARY_OP_LOCALS;
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *) src0->extra;
+    ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *) src1->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) dst->extra;
+
+    const cl_long offset0 = extra0->offset + src0->view_offs;
+    const cl_long offset1 = extra1->offset + src1->view_offs;
+    const cl_long offsetd = extrad->offset + dst->view_offs;
+
+    const cl_long C  = ne12;
+    const cl_long N  = ne13;
+    const cl_long KW = ne00;
+    const cl_long KH = ne01;
+    const cl_long W  = ne10;
+    const cl_long H  = ne11;
+    const cl_long OW = ne0;
+    const cl_long OH = ne1;
+
+    const cl_int s0 = dst->op_params[0];
+    const cl_int s1 = dst->op_params[1];
+    const cl_int p0 = dst->op_params[2];
+    const cl_int p1 = dst->op_params[3];
+    const cl_int d0 = dst->op_params[4];
+    const cl_int d1 = dst->op_params[5];
+
+    const cl_long cl_nb01 = nb01/ggml_type_size(src0->type);
+    const cl_long cl_nb03 = nb03/ggml_type_size(src0->type);
+    const cl_long cl_nb11 = nb11/ggml_type_size(src1->type);
+    const cl_long cl_nb12 = nb12/ggml_type_size(src1->type);
+    const cl_long cl_nb13 = nb13/ggml_type_size(src1->type);
+    const cl_long cl_nb1  = nb1/ggml_type_size(dst->type);
+    const cl_long cl_nb2  = nb2/ggml_type_size(dst->type);
+    const cl_long cl_nb3  = nb3/ggml_type_size(dst->type);
+
+    cl_kernel kernel;
+    if (src0->type == GGML_TYPE_F32) {
+        kernel = backend_ctx->kernel_conv2d_dw_f32_f32;
+    } else if (src0->type == GGML_TYPE_F16) {
+        kernel = backend_ctx->kernel_conv2d_dwf16_f32;
+    } else {
+        GGML_ASSERT(false && "Unsupported data type combination for conv2_dw");
+    }
+
+    cl_uint idx = 0;
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extra0->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &offset0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extra1->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &offset1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extrad->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &offsetd));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &C));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &N));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &KW));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &KH));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &W));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &H));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &OW));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &OH));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &s0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &s1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &p0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &p1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &d0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_int),   &d1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb01));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb03));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb11));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb12));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb13));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb2));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_long),  &cl_nb3));
+
+    const size_t total = (size_t) N * C * OH * OW;
+    const size_t local_size = MIN((size_t) 256, backend_ctx->get_kernel_workgroup_size(kernel));
+    size_t global_work_size[] = { CEIL_DIV(total, local_size) * local_size };
+    size_t local_work_size[] = { local_size };
+
+    backend_ctx->enqueue_ndrange_kernel(kernel, 1, global_work_size, local_work_size, dst);
+}
+
 // is_kq selects which of the two products this call is, and it is decided by the
 // CALLER -- the two admission arms in ggml_cl_mul_mat, each of which knows which
 // one it matched. It used to be re-derived here from nb01 > nb02, i.e. "K is
@@ -31030,6 +31137,12 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
            }
            func = ggml_cl_conv_2d;
            break;
+        case GGML_OP_CONV_2D_DW:
+            if (!any_on_device) {
+                return false;
+            }
+            func = ggml_cl_conv2_dw;
+            break;
        case GGML_OP_SSM_CONV:
            if (!any_on_device) {
                return false;
@@ -0,0 +1,139 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void kernel_conv2d_dw_f32_f32(
+        global const float * weights,
+        long                 weights_offset,
+        global const float * input,
+        long                 input_offset,
+        global       float * output,
+        long                 output_offset,
+        long C,
+        long N,
+        long KW,
+        long KH,
+        long W,
+        long H,
+        long OW,
+        long OH,
+        int s0,
+        int s1,
+        int p0,
+        int p1,
+        int d0,
+        int d1,
+        long nb01,
+        long nb03,
+        long nb11,
+        long nb12,
+        long nb13,
+        long nb1,
+        long nb2,
+        long nb3
+) {
+    weights = (global const float *) ((global const char *) weights + weights_offset);
+    input = (global const float *) ((global const char *) input + input_offset);
+    output = (global float *) ((global char *) output + output_offset);
+
+    const long gid = get_global_id(0);
+    const long OHW = OH * OW;
+    const long COHW = C * OHW;
+    const long total = N * COHW;
+    if (gid >= total) {
+        return;
+    }
+
+    const long n = gid / COHW;
+    const long c_rem = gid - n * COHW;
+    const long c = c_rem / OHW;
+    const long oh = (c_rem - c * OHW) / OW;
+    const long ow = c_rem - c * OHW - oh * OW;
+
+    float sum = 0.0f;
+    for (long kh = 0; kh < KH; ++kh) {
+        const long ih = oh * s1 + kh * d1 - p1;
+        if (ih < 0 || ih >= H) {
+            continue;
+        }
+        for (long kw = 0; kw < KW; ++kw) {
+            const long iw = ow * s0 + kw * d0 - p0;
+            if (iw < 0 || iw >= W) {
+                continue;
+            }
+            const long weight_idx = kw + kh * nb01 + c * nb03;
+            const long input_idx = iw + ih * nb11 + c * nb12 + n * nb13;
+            sum += weights[weight_idx] * input[input_idx];
+        }
+    }
+
+    const long output_idx = ow + oh * nb1 + c * nb2 + n * nb3;
+    output[output_idx] = sum;
+}
+
+kernel void kernel_conv2d_dwf16_f32(
+        global const half * weights,
+        long                weights_offset,
+        global const float * input,
+        long                 input_offset,
+        global       float * output,
+        long                 output_offset,
+        long C,
+        long N,
+        long KW,
+        long KH,
+        long W,
+        long H,
+        long OW,
+        long OH,
+        int s0,
+        int s1,
+        int p0,
+        int p1,
+        int d0,
+        int d1,
+        long nb01,
+        long nb03,
+        long nb11,
+        long nb12,
+        long nb13,
+        long nb1,
+        long nb2,
+        long nb3
+) {
+    weights = (global const half *) ((global const char *) weights + weights_offset);
+    input = (global const float *) ((global const char *) input + input_offset);
+    output = (global float *) ((global char *) output + output_offset);
+
+    const long gid = get_global_id(0);
+    const long OHW = OH * OW;
+    const long COHW = C * OHW;
+    const long total = N * COHW;
+    if (gid >= total) {
+        return;
+    }
+
+    const long n = gid / COHW;
+    const long c_rem = gid - n * COHW;
+    const long c = c_rem / OHW;
+    const long oh = (c_rem - c * OHW) / OW;
+    const long ow = c_rem - c * OHW - oh * OW;
+
+    float sum = 0.0f;
+    for (long kh = 0; kh < KH; ++kh) {
+        const long ih = oh * s1 + kh * d1 - p1;
+        if (ih < 0 || ih >= H) {
+            continue;
+        }
+        for (long kw = 0; kw < KW; ++kw) {
+            const long iw = ow * s0 + kw * d0 - p0;
+            if (iw < 0 || iw >= W) {
+                continue;
+            }
+            const long weight_idx = kw + kh * nb01 + c * nb03;
+            const long input_idx = iw + ih * nb11 + c * nb12 + n * nb13;
+            sum += (float) weights[weight_idx] * input[input_idx];
+        }
+    }
+
+    const long output_idx = ow + oh * nb1 + c * nb2 + n * nb3;
+    output[output_idx] = sum;
+}