upstream-archive/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-07-23 11:10:57 -05:00
Code Issues Packages Projects Releases Wiki Activity

fix unsupport ACC UT cases for noncontiguous (llama/25124)

Browse Source 
* fix unsupport ACC UT cases for noncontiguous

* update ops.md
This commit is contained in:
Neo Zhang
2026-07-07 15:40:38 +08:00
committed by Georgi Gerganov
parent bfd86f4304
commit 120fc3729d
2 changed files with 41 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

56
ggml/src/ggml-sycl/element_wise.cpp
View File

@@ -9,9 +9,12 @@
#define SYCL_LOCAL_ID_CALC(ITEM, IDX) \
(ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX))
static void acc_f32(const float * x, const float * y, float * dst, const int64_t ne,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
static void acc_f32(const char * x, const char * y, float * dst, const int64_t ne,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
const int64_t i = SYCL_LOCAL_ID_CALC(item_ct1, 2);
@@ -30,9 +33,18 @@ static void acc_f32(const float * x, const float * y, float * dst, const int64_t
tmp -= i11 * s11;
const int64_t i10 = tmp;
float val = x[i];
int64_t tmp_dst = i;
const int64_t i3 = tmp_dst / (ne2*ne1*ne0);
tmp_dst -= i3 * (ne2*ne1*ne0);
const int64_t i2 = tmp_dst / (ne1*ne0);
tmp_dst -= i2 * (ne1*ne0);
const int64_t i1 = tmp_dst / ne0;
tmp_dst -= i1 * ne0;
const int64_t i0 = tmp_dst;
float val = *(const float *) (x + i0*nb00 + i1*nb01 + i2*nb02 + i3*nb03);
if (src1_idx >= 0 && i10 < ne10 && i11 < ne11 && i12 < ne12 && i13 < ne13) {
val += y[((i13*ne12 + i12) * ne11 + i11) * ne10 + i10];
val += *(const float *) (y + i10*nb10 + i11*nb11 + i12*nb12 + i13*nb13);
}
dst[i] = val;
}
@@ -422,15 +434,24 @@ static void gated_op_fused_geglu_quick(const T * x, const T * g, T * dst, const
}
namespace ggml_sycl_detail {
static void acc_f32_sycl(const float *x, const float *y, float *dst,
const int64_t n_elements, const int64_t ne10, const int64_t ne11,
const int64_t ne12, const int64_t ne13, const int64_t s1, const int64_t s2, const int64_t s3,
static void acc_f32_sycl(const char *x, const char *y, float *dst,
const int64_t n_elements,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
const int64_t s1, const int64_t s2, const int64_t s3,
const int64_t offset, queue_ptr stream) {
const int num_blocks = (n_elements + SYCL_ACC_BLOCK_SIZE - 1) / SYCL_ACC_BLOCK_SIZE;
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE),
sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE)),
[=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
acc_f32(x, y, dst, n_elements, ne10, ne11, ne12, ne13, s1, s2, s3, offset);
acc_f32(x, y, dst, n_elements,
ne0, ne1, ne2, ne3,
nb00, nb01, nb02, nb03,
ne10, ne11, ne12, ne13,
nb10, nb11, nb12, nb13,
s1, s2, s3, offset);
});
}
@@ -843,8 +864,8 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
const float * src0_d = (const float *) src0->data;
const float * src1_d = (const float *) src1->data;
const char * src0_d = (const char *) src0->data;
const char * src1_d = (const char *) src1->data;
float * dst_d = (float *) dst->data;
dpct::queue_ptr stream = ctx.stream();
@@ -853,17 +874,20 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
GGML_ASSERT(src1->type == GGML_TYPE_F32);
GGML_ASSERT( dst->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_contiguous(src1));
GGML_ASSERT(dst->nb[0] == ggml_element_size(dst));
GGML_ASSERT(ggml_is_contiguously_allocated(dst));
GGML_ASSERT(ggml_are_same_shape(src0, dst));
const int64_t s1 = dst->op_params[0] / sizeof(float);
const int64_t s2 = dst->op_params[1] / sizeof(float);
const int64_t s3 = dst->op_params[2] / sizeof(float);
const int64_t offset = dst->op_params[3] / sizeof(float);
const int64_t s1 = (int64_t) ((const int32_t *) dst->op_params)[0] / (int64_t) sizeof(float);
const int64_t s2 = (int64_t) ((const int32_t *) dst->op_params)[1] / (int64_t) sizeof(float);
const int64_t s3 = (int64_t) ((const int32_t *) dst->op_params)[2] / (int64_t) sizeof(float);
const int64_t offset = (int64_t) ((const int32_t *) dst->op_params)[3] / (int64_t) sizeof(float);
ggml_sycl_detail::acc_f32_sycl(src0_d, src1_d, dst_d, ggml_nelements(dst),
dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
s1, s2, s3, offset, stream);
}

3
ggml/src/ggml-sycl/ggml-sycl.cpp
View File

@@ -5852,9 +5852,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
}
case GGML_OP_POOL_2D:
case GGML_OP_POOL_1D:
return true;
case GGML_OP_ACC:
return ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]);
return true;
case GGML_OP_PAD:
if (ggml_get_op_params_i32(op, 8) != 0) {
return false;