Lumpiasty/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

sycl: fix multiple warnings in compiling sycl backend (#26713)

Browse Source 
* Update norm.cpp

* Update helper.hpp

* Update im2col.cpp

* Update fattn-mkl.cpp

* Update element_wise.cpp

* Update fattn-mkl.cpp

* Update set_rows.cpp

* Update element_wise.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update norm.cpp

* Update CMakeLists.txt

* Update CMakeLists.txt

* Update CMakeLists.txt

* Update ggml-sycl.cpp
This commit is contained in:
Ian Faust
2026-08-21 08:15:40 +03:00
committed by GitHub
parent 9e96cf77ff
commit 6602dd3389
8 changed files with 28 additions and 34 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cpu/CMakeLists.txt
+5 -4
View File
@@ -737,8 +737,9 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
set_target_properties(${GGML_CPU_NAME} PROPERTIES COMPILE_FLAGS "-msimd128") set_target_properties(${GGML_CPU_NAME} PROPERTIES COMPILE_FLAGS "-msimd128")
endif() endif()
if (CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM") if (CMAKE_C_COMPILER_ID STREQUAL "IntelLLVM" OR CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM")
# The compiler automatically enables "-ffast-math" which can cause NaNs in tests due to "-fassociative-math" # The compiler automatically enables "-ffast-math" which can cause NaNs in tests due to "-fassociative-math"
target_compile_options(${GGML_CPU_NAME} PRIVATE "-fno-associative-math") target_compile_options(${GGML_CPU_NAME} PRIVATE "$<$<OR:$<COMPILE_LANG_AND_ID:C,IntelLLVM>,$<COMPILE_LANG_AND_ID:CXX,IntelLLVM>>:$<$<BOOL:${WIN32}>:/clang:>-fno-associative-math>")
endif() endif()
endfunction() endfunction()
ggml/src/ggml-sycl/dpct/helper.hpp
+1 -1
View File
@@ -62,7 +62,7 @@
#define DPCT_UNUSED(x) (void)(x) #define DPCT_UNUSED(x) (void)(x)
inline void _abort(const char * str) { [[noreturn]] inline void _abort(const char * str) {
std::cerr << str << std::endl; std::cerr << str << std::endl;
std::abort(); std::abort();
} }
ggml/src/ggml-sycl/element_wise.cpp
+4 -4
View File
@@ -10,7 +10,7 @@
(ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX)) (ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX))
static void acc_f32(const char * x, const char * y, float * dst, const int64_t ne, 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 ne0, const int64_t ne1, const int64_t ne2,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03, 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 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 nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
@@ -455,7 +455,7 @@ static void unary_mul_sycl(const T * x, const T * g, T * dst, const int64_t k, c
namespace ggml_sycl_detail { namespace ggml_sycl_detail {
static void acc_f32_sycl(const char *x, const char *y, float *dst, static void acc_f32_sycl(const char *x, const char *y, float *dst,
const int64_t n_elements, const int64_t n_elements,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, const int64_t ne0, const int64_t ne1, const int64_t ne2,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03, 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 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 nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
@@ -466,7 +466,7 @@ static void acc_f32_sycl(const char *x, const char *y, float *dst,
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)]] { [=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
acc_f32(x, y, dst, n_elements, acc_f32(x, y, dst, n_elements,
ne0, ne1, ne2, ne3, ne0, ne1, ne2,
nb00, nb01, nb02, nb03, nb00, nb01, nb02, nb03,
ne10, ne11, ne12, ne13, ne10, ne11, ne12, ne13,
nb10, nb11, nb12, nb13, nb10, nb11, nb12, nb13,
@@ -970,7 +970,7 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
const int64_t offset = (int64_t) ((const int32_t *) dst->op_params)[3] / (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), 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], dst->ne[0], dst->ne[1], dst->ne[2],
src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[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->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3], src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
ggml/src/ggml-sycl/fattn-mkl.cpp
+4 -5
View File
@@ -43,7 +43,7 @@ static void mkl_fa_pack_q_fp16(
dpct::queue_ptr stream, dpct::queue_ptr stream,
sycl::half * __restrict dst, sycl::half * __restrict dst,
const float * __restrict q_src, const float * __restrict q_src,
int n_queries, int n_query_rows, int DKQ, int n_queries, int DKQ,
int gqa_ratio, int kvh_base_head, int gqa_ratio, int kvh_base_head,
float q_scale, int64_t q_row_stride, int64_t q_head_stride, float q_scale, int64_t q_row_stride, int64_t q_head_stride,
int64_t wg_size) { int64_t wg_size) {
@@ -121,7 +121,7 @@ static void mkl_fa_online_softmax_chunk(
float * __restrict VKQ_accum, float * __restrict VKQ_accum,
int q0, int q_rows, int n_queries, int DV, int q0, int q_rows, int n_queries, int DV,
int chunk_size, int chunk_start, int chunk_size, int chunk_start,
int kvh_head, int gqa_ratio, int kvh_head,
const sycl::half * mask_data, int64_t mask_head_stride, const sycl::half * mask_data, int64_t mask_head_stride,
int64_t mask_row_stride, int mask_n_heads, int64_t mask_row_stride, int mask_n_heads,
float logit_softcap, int64_t wg_size) { float logit_softcap, int64_t wg_size) {
@@ -473,7 +473,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor *
MKL_ACCUM(dequant_time_us, t_deq); MKL_ACCUM(dequant_time_us, t_deq);
// --- Resolve mask pointers --- // --- Resolve mask pointers ---
const sycl::half * mask_data = nullptr;
int64_t mask_head_stride = 0; int64_t mask_head_stride = 0;
int64_t mask_row_stride = 0; int64_t mask_row_stride = 0;
int mask_n_heads = 0; int mask_n_heads = 0;
@@ -547,7 +546,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor *
// 1. Pack all GQA Q heads into fp16 (full n_query_rows) // 1. Pack all GQA Q heads into fp16 (full n_query_rows)
mkl_fa_pack_q_fp16(stream, mkl_fa_pack_q_fp16(stream,
Q_head_f16_ptr, Q_batch, Q_head_f16_ptr, Q_batch,
n_queries, n_query_rows, DKQ, n_queries, DKQ,
gqa_ratio, kvh_base_head, gqa_ratio, kvh_base_head,
q_scale, q_row_stride, q_head_stride, wg_size); q_scale, q_row_stride, q_head_stride, wg_size);
@@ -605,7 +604,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor *
KQ_max_ptr, KQ_sum_ptr, VKQ_accum_ptr, KQ_max_ptr, KQ_sum_ptr, VKQ_accum_ptr,
q0, q_rows, n_queries, DV, q0, q_rows, n_queries, DV,
this_chunk, chunk_start, this_chunk, chunk_start,
kvh_base_head, gqa_ratio, kvh_base_head,
mask_batch, mask_head_stride, mask_batch, mask_head_stride,
mask_row_stride, mask_n_heads, mask_row_stride, mask_n_heads,
logit_softcap, wg_size); logit_softcap, wg_size);
ggml/src/ggml-sycl/ggml-sycl.cpp
+11 -9
View File
@@ -921,16 +921,16 @@ ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft,
void * dev_ptr; void * dev_ptr;
if (use_usm_system) { if (use_usm_system) {
GGML_SYCL_DEBUG("[SYCL] allocating %lu Bytes with USM system\n", size); GGML_SYCL_DEBUG("[SYCL] allocating %zu Bytes with USM system\n", size);
dev_ptr = (void *)aligned_malloc_host(alignment, aligned_size); dev_ptr = (void *)aligned_malloc_host(alignment, aligned_size);
if (!dev_ptr) { if (!dev_ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size); GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on host\n", __func__, size);
return nullptr; return nullptr;
} }
} else { } else {
SYCL_CHECK(CHECK_TRY_ERROR(dev_ptr = (void *)ggml_sycl_malloc_device(size, *stream))); SYCL_CHECK(CHECK_TRY_ERROR(dev_ptr = (void *)ggml_sycl_malloc_device(size, *stream)));
if (!dev_ptr) { if (!dev_ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device\n", __func__, size); GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device\n", __func__, size);
return nullptr; return nullptr;
} }
} }
@@ -1177,7 +1177,7 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer,
SYCL_CHECK(CHECK_TRY_ERROR(buf = (char *)ggml_sycl_malloc_device(size, *stream))); SYCL_CHECK(CHECK_TRY_ERROR(buf = (char *)ggml_sycl_malloc_device(size, *stream)));
if (!buf) { if (!buf) {
char err_buf[1024]; char err_buf[1024];
snprintf(err_buf, 1023, "%s: can't allocate %lu Bytes of memory on device\n", __func__, size); snprintf(err_buf, 1023, "%s: can't allocate %zu Bytes of memory on device\n", __func__, size);
throw std::runtime_error(err_buf); throw std::runtime_error(err_buf);
} }
// set padding to 0 to avoid possible NaN values // set padding to 0 to avoid possible NaN values
@@ -1651,7 +1651,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool {
SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *)ggml_sycl_malloc_device(look_ahead_size, *qptr))); SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *)ggml_sycl_malloc_device(look_ahead_size, *qptr)));
if (!ptr) { if (!ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device/GPU\n", __func__, look_ahead_size); GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device/GPU\n", __func__, look_ahead_size);
return nullptr; return nullptr;
} }
@@ -1663,7 +1663,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool {
(uint32_t)(max_size/1024/1024), (uint32_t)(g_sycl_pool_size[id]/1024/1024), (uint32_t)(size/1024/1024)); (uint32_t)(max_size/1024/1024), (uint32_t)(g_sycl_pool_size[id]/1024/1024), (uint32_t)(size/1024/1024));
#endif #endif
// GGML_SYCL_DEBUG("ggml_sycl_pool_malloc_leg look_ahead_size=%lu, return %p\n", look_ahead_size, ptr); // GGML_SYCL_DEBUG("ggml_sycl_pool_malloc_leg look_ahead_size=%zu, return %p\n", look_ahead_size, ptr);
return ptr; return ptr;
} }
@@ -1843,7 +1843,7 @@ struct ggml_sycl_pool_host : public ggml_sycl_pool {
SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *) sycl::malloc_host(size, *qptr))); SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *) sycl::malloc_host(size, *qptr)));
if (!ptr) { if (!ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size); GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on host\n", __func__, size);
return nullptr; return nullptr;
} }
pool_size += size; pool_size += size;
@@ -2779,9 +2779,9 @@ inline void ggml_sycl_op_mul_mat_sycl(
const float * src1_ddf1_i = src1->type == GGML_TYPE_F32 ? (const float *) src1_ddf_i : src1_ddq_as_f32.get(); const float * src1_ddf1_i = src1->type == GGML_TYPE_F32 ? (const float *) src1_ddf_i : src1_ddq_as_f32.get();
{ {
#if GGML_SYCL_DNNL
const int64_t gemm_flops = (int64_t)row_diff * src1_ncols * ne10; const int64_t gemm_flops = (int64_t)row_diff * src1_ncols * ne10;
const bool use_mkl_direct = gemm_flops < 256 * 256 * 256; const bool use_mkl_direct = gemm_flops < 256 * 256 * 256;
#if GGML_SYCL_DNNL
if (g_ggml_sycl_enable_dnn && !use_mkl_direct) { if (g_ggml_sycl_enable_dnn && !use_mkl_direct) {
DnnlGemmWrapper::row_gemm(ctx, row_diff, src1_ncols, ne10, src0_ddf_i, DnnlGemmWrapper::row_gemm(ctx, row_diff, src1_ncols, ne10, src0_ddf_i,
DnnlGemmWrapper::to_dt<float>(), src1_ddf1_i, DnnlGemmWrapper::to_dt<float>(), DnnlGemmWrapper::to_dt<float>(), src1_ddf1_i, DnnlGemmWrapper::to_dt<float>(),
@@ -3518,7 +3518,9 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
float * dst_ddf = static_cast<float *>(dst->data); float * dst_ddf = static_cast<float *>(dst->data);
const sycl::half * src1_f16 = static_cast<const sycl::half *>(src1->data); const sycl::half * src1_f16 = static_cast<const sycl::half *>(src1->data);
#if GGML_SYCL_DNNL
const size_t type_size_src0 = ggml_type_size(src0->type); const size_t type_size_src0 = ggml_type_size(src0->type);
#endif
const size_t type_size_src1 = ggml_type_size(src1->type); const size_t type_size_src1 = ggml_type_size(src1->type);
bool is_src0_cont_2 = ggml_is_contiguous_2(src0); bool is_src0_cont_2 = ggml_is_contiguous_2(src0);
@@ -3535,6 +3537,7 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
scope_op_debug_print scope_dbg_print(__func__, "/to_fp16_nc_sycl", dst, /*num_src=*/2, scope_op_debug_print scope_dbg_print(__func__, "/to_fp16_nc_sycl", dst, /*num_src=*/2,
" : converting src1 to fp16"); " : converting src1 to fp16");
#if GGML_SYCL_DNNL
// iterate tensor dims and find the slowest moving dim and stride // iterate tensor dims and find the slowest moving dim and stride
int last_dim=0; int last_dim=0;
int last_str=0; int last_str=0;
@@ -3554,7 +3557,6 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
} }
} }
#if GGML_SYCL_DNNL
// oneDNN handles strided data and does not need overhead of ggml_get_to_fp16_nc_sycl // oneDNN handles strided data and does not need overhead of ggml_get_to_fp16_nc_sycl
const int64_t ne_src1 = src1->nb[last_str] * src1->ne[last_dim] / type_size_src1; const int64_t ne_src1 = src1->nb[last_str] * src1->ne[last_dim] / type_size_src1;
src1_f16_alloc.alloc(ne_src1); src1_f16_alloc.alloc(ne_src1);
ggml/src/ggml-sycl/im2col.cpp
+2 -2
View File
@@ -85,7 +85,7 @@ static void im2col_sycl(const float * x,
*/ */
stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE)), stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE)),
sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE))), sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE))),
[=](sycl::nd_item<3> item_ct1) { [=](sycl::nd_item<3>) {
im2col_kernel(x, dst, IC, IW, IH, OH, OW, KW, KH, IC_IH_IW, IH_IW, N_OH, KH_KW, IC_KH_KW, im2col_kernel(x, dst, IC, IW, IH, OH, OW, KW, KH, IC_IH_IW, IH_IW, N_OH, KH_KW, IC_KH_KW,
s0, s1, p0, p1, d0, d1); s0, s1, p0, p1, d0, d1);
}); });
@@ -271,7 +271,7 @@ static void im2col_3d_sycl(const float * src,
*/ */
stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE)), stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE)),
sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE))), sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE))),
[=](sycl::nd_item<3> item_ct1) { [=](sycl::nd_item<3>) {
im2col_3d_kernel(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, OH_OW, KD_KH_KW, im2col_3d_kernel(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, OH_OW, KD_KH_KW,
ID_IH_IW, KH_KW, IH_IW, IC_ID_IH_IW, IC_KD_KH_KW, OW_KD_KH_KW, ID_IH_IW, KH_KW, IH_IW, IC_ID_IH_IW, IC_KD_KH_KW, OW_KD_KH_KW,
OD_OH_OW_IC_KD_KH_KW, OH_OW_IC_KD_KH_KW, OW_IC_KD_KH_KW, N_OD_OH, OD_OH, OD_OH_OW_IC_KD_KH_KW, OH_OW_IC_KD_KH_KW, OW_IC_KD_KH_KW, N_OD_OH, OD_OH,
ggml/src/ggml-sycl/norm.cpp
-8
View File
@@ -7,9 +7,6 @@ static void norm_f32(const float* x, float* dst, const int ncols,
const int64_t dst_stride_col, const int64_t dst_stride_row, const int64_t dst_stride_channel, const int64_t dst_stride_sample, const int64_t dst_stride_col, const int64_t dst_stride_row, const int64_t dst_stride_channel, const int64_t dst_stride_sample,
const float eps, const sycl::nd_item<3>& item_ct1, sycl::float2* s_sum, int block_size) { const float eps, const sycl::nd_item<3>& item_ct1, sycl::float2* s_sum, int block_size) {
const int nrows = item_ct1.get_group_range(2);
const int nchannels = item_ct1.get_group_range(1);
const int nthreads = item_ct1.get_local_range(2); const int nthreads = item_ct1.get_local_range(2);
const int sample = item_ct1.get_group(0); const int sample = item_ct1.get_group(0);
const int channel = item_ct1.get_group(1); const int channel = item_ct1.get_group(1);
@@ -155,9 +152,6 @@ static void rms_norm_f32(const float* x, float* dst, const int ncols,
const float* mul = nullptr, const int64_t mul_stride_row = 0, const int64_t mul_stride_channel = 0, const float* mul = nullptr, const int64_t mul_stride_row = 0, const int64_t mul_stride_channel = 0,
const int64_t mul_stride_sample = 0, const int mul_nrows = 0, const int mul_nchannels = 0, const int mul_nsamples = 0) { const int64_t mul_stride_sample = 0, const int mul_nrows = 0, const int mul_nchannels = 0, const int mul_nsamples = 0) {
const int nrows = item_ct1.get_group_range(2);
const int nchannels = item_ct1.get_group_range(1);
const int sample = item_ct1.get_group(0); const int sample = item_ct1.get_group(0);
const int channel = item_ct1.get_group(1); const int channel = item_ct1.get_group(1);
const int row = item_ct1.get_group(2); const int row = item_ct1.get_group(2);
@@ -225,8 +219,6 @@ static void l2_norm_f32(const float * x, float * dst, const int ncols,
const int64_t src_stride_sample, const int64_t dst_stride_col, const int64_t dst_stride_row, const int64_t src_stride_sample, const int64_t dst_stride_col, const int64_t dst_stride_row,
const int64_t dst_stride_channel, const int64_t dst_stride_sample, const float eps, const int64_t dst_stride_channel, const int64_t dst_stride_sample, const float eps,
const sycl::nd_item<3>& item_ct1, float* s_sum, const int block_size) { const sycl::nd_item<3>& item_ct1, float* s_sum, const int block_size) {
const int nrows = item_ct1.get_group_range(2);
const int nchannels = item_ct1.get_group_range(1);
const int row = item_ct1.get_group(2); const int row = item_ct1.get_group(2);
const int channel = item_ct1.get_group(1); const int channel = item_ct1.get_group(1);
ggml/src/ggml-sycl/set_rows.cpp
+1 -1
View File
@@ -291,7 +291,7 @@ static void set_rows_sycl(
stream->parallel_for( stream->parallel_for(
sycl::nd_range<1>(grid_size * block_size, block_size), sycl::nd_range<1>(grid_size * block_size, block_size),
[=](sycl::nd_item<1> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] { [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
k_set_rows<TIn, TIdx, TOut>( k_set_rows<TIn, TIdx, TOut>(
src0_d, src1_d, dst_d, src0_d, src1_d, dst_d,
ne00, ne01, ne02, ne00, ne01, ne02,