diff --git a/ggml/src/ggml-sycl/conv2d-dw.cpp b/ggml/src/ggml-sycl/conv2d-dw.cpp index 0a52b7917..8755a4c95 100644 --- a/ggml/src/ggml-sycl/conv2d-dw.cpp +++ b/ggml/src/ggml-sycl/conv2d-dw.cpp @@ -71,8 +71,8 @@ struct dw_cwhn_layout { } }; -template -static void conv2d_dw_kernel(const float * input, const float * kernel, float * output, +template +static void conv2d_dw_kernel(const float * input, const KernelT * kernel, float * output, const conv2d_dw_params p, const sycl::nd_item<3> & item_ct1) { const int global_idx = item_ct1.get_local_id(2) + item_ct1.get_group(2) * item_ct1.get_local_range(2); @@ -93,15 +93,15 @@ static void conv2d_dw_kernel(const float * input, const float * kernel, float * for (int kx = bounds.x_min; kx < bounds.x_max; ++kx) { const int in_x = dw_calculate_input_coord(out_x, kx, p.stride_x, p.dilation_x, p.padding_x); acc += input[Layout::input_index(n, c, in_y, in_x, p)] * - kernel[Layout::kernel_index(c, ky, kx, p)]; + static_cast(kernel[Layout::kernel_index(c, ky, kx, p)]); } } output[Layout::output_index(n, c, out_y, out_x, p)] = acc; } -template -static void conv2d_dw_sycl(const float * x_d, const float * w_d, float * y_d, +template +static void conv2d_dw_sycl(const float * x_d, const KernelT * w_d, float * y_d, const conv2d_dw_params p, const queue_ptr & stream) { const int total = p.batches * p.channels * p.out_h * p.out_w; const int num_blocks = (total + SYCL_CONV2D_DW_BLOCK_SIZE - 1) / SYCL_CONV2D_DW_BLOCK_SIZE; @@ -109,7 +109,7 @@ static void conv2d_dw_sycl(const float * x_d, const float * w_d, float * y_d, const sycl::range<3> block_nums(1, 1, num_blocks); stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) { - conv2d_dw_kernel(x_d, w_d, y_d, p, item_ct1); + conv2d_dw_kernel(x_d, w_d, y_d, p, item_ct1); }); } @@ -119,9 +119,9 @@ void ggml_sycl_op_conv2d_dw(ggml_backend_sycl_context & ctx, ggml_tensor * dst) const ggml_tensor * kernel = dst->src[0]; const ggml_tensor * input = dst->src[1]; - GGML_ASSERT(kernel->type == GGML_TYPE_F32 && input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32); + GGML_ASSERT((kernel->type == GGML_TYPE_F32 || kernel->type == GGML_TYPE_F16) && + input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32); - const float * w_d = (const float *) kernel->data; const float * x_d = (const float *) input->data; float * y_d = (float *) dst->data; @@ -148,11 +148,23 @@ void ggml_sycl_op_conv2d_dw(ggml_backend_sycl_context & ctx, ggml_tensor * dst) const queue_ptr stream = ctx.stream(); - if (ggml_is_contiguous(input)) { - conv2d_dw_sycl(x_d, w_d, y_d, params, stream); - } else if (ggml_is_contiguous_channels(input)) { - conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + if (kernel->type == GGML_TYPE_F16) { + const sycl::half * w_d = (const sycl::half *) kernel->data; + if (ggml_is_contiguous(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else if (ggml_is_contiguous_channels(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else { + GGML_ABORT("Unsupported memory layout for conv2d_dw"); + } } else { - GGML_ABORT("Unsupported memory layout for conv2d_dw"); + const float * w_d = (const float *) kernel->data; + if (ggml_is_contiguous(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else if (ggml_is_contiguous_channels(input)) { + conv2d_dw_sycl(x_d, w_d, y_d, params, stream); + } else { + GGML_ABORT("Unsupported memory layout for conv2d_dw"); + } } }