getrows: move to a separate file
This commit is contained in:
Akarshan Biswas
parent
04d8b038b8
commit
98f5fd2fd1
4 changed files with 174 additions and 191 deletions
|
|
@ -33,6 +33,7 @@
|
||||||
#include "argmax.hpp"
|
#include "argmax.hpp"
|
||||||
#include "argsort.hpp"
|
#include "argsort.hpp"
|
||||||
#include "cpy.hpp"
|
#include "cpy.hpp"
|
||||||
|
#include "getrows.hpp"
|
||||||
#include "gla.hpp"
|
#include "gla.hpp"
|
||||||
|
|
||||||
#endif // GGML_SYCL_BACKEND_HPP
|
#endif // GGML_SYCL_BACKEND_HPP
|
||||||
|
|
|
||||||
165
ggml/src/ggml-sycl/getrows.cpp
Normal file
165
ggml/src/ggml-sycl/getrows.cpp
Normal file
|
|
@ -0,0 +1,165 @@
|
||||||
|
#include "getrows.hpp"
|
||||||
|
#include "dequantize.hpp"
|
||||||
|
|
||||||
|
template <int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
|
||||||
|
static void k_get_rows(const void * src0, const int32_t * src1, dst_t * dst,
|
||||||
|
int64_t ne00, /*int64_t ne01, int64_t ne02, int64_t ne03,*/
|
||||||
|
/*int64_t ne10, int64_t ne11,*/ int64_t ne12, /*int64_t ne13,*/
|
||||||
|
/*size_t s0,*/ size_t s1, size_t s2, size_t s3,
|
||||||
|
/*size_t nb00,*/ size_t nb01, size_t nb02, size_t nb03, size_t s10, size_t s11, size_t s12,
|
||||||
|
const sycl::nd_item<3> & item_ct1 /*, size_t s13*/) {
|
||||||
|
const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) + item_ct1.get_local_id(2)) * 2;
|
||||||
|
const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) + item_ct1.get_local_id(1);
|
||||||
|
const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + item_ct1.get_local_id(0)) / ne12;
|
||||||
|
const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + item_ct1.get_local_id(0)) % ne12;
|
||||||
|
|
||||||
|
if (i00 >= ne00) {
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
const int i01 = src1[i10 * s10 + i11 * s11 + i12 * s12];
|
||||||
|
|
||||||
|
dst_t * dst_row = dst + i10 * s1 + i11 * s2 + i12 * s3;
|
||||||
|
const void * src0_row = (const char *) src0 + i01 * nb01 + i11 * nb02 + i12 * nb03;
|
||||||
|
|
||||||
|
const int ib = i00 / qk; // block index
|
||||||
|
const int iqs = (i00 % qk) / qr; // quant index
|
||||||
|
const int iybs = i00 - i00 % qk; // dst block start index
|
||||||
|
const int y_offset = qr == 1 ? 1 : qk / 2;
|
||||||
|
|
||||||
|
// dequantize
|
||||||
|
dfloat2 v;
|
||||||
|
dequantize_kernel(src0_row, ib, iqs, v);
|
||||||
|
|
||||||
|
dst_row[iybs + iqs + 0] = v.x();
|
||||||
|
dst_row[iybs + iqs + y_offset] = v.y();
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename src0_t, typename dst_t>
|
||||||
|
static void k_get_rows_float(const src0_t * src0, const int32_t * src1, dst_t * dst,
|
||||||
|
int64_t ne00, /*int64_t ne01, int64_t ne02, int64_t ne03,*/
|
||||||
|
/*int64_t ne10, int64_t ne11,*/ int64_t ne12, /*int64_t ne13,*/
|
||||||
|
/*size_t s0,*/ size_t s1, size_t s2, size_t s3,
|
||||||
|
/*size_t nb00,*/ size_t nb01, size_t nb02, size_t nb03, size_t s10, size_t s11, size_t s12,
|
||||||
|
const sycl::nd_item<3> & item_ct1 /*, size_t s13*/) {
|
||||||
|
const int i00 = item_ct1.get_group(2) * item_ct1.get_local_range(2) + item_ct1.get_local_id(2);
|
||||||
|
const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) + item_ct1.get_local_id(1);
|
||||||
|
const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + item_ct1.get_local_id(0)) / ne12;
|
||||||
|
const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) + item_ct1.get_local_id(0)) % ne12;
|
||||||
|
|
||||||
|
if (i00 >= ne00) {
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
const int i01 = src1[i10 * s10 + i11 * s11 + i12 * s12];
|
||||||
|
|
||||||
|
dst_t * dst_row = dst + i10 * s1 + i11 * s2 + i12 * s3;
|
||||||
|
const src0_t * src0_row = (const src0_t *) ((const char *) src0 + i01 * nb01 + i11 * nb02 + i12 * nb03);
|
||||||
|
|
||||||
|
dst_row[i00] = src0_row[i00];
|
||||||
|
}
|
||||||
|
|
||||||
|
template <int qk, int qr, dequantize_kernel_t dq>
|
||||||
|
static void get_rows_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
|
||||||
|
GGML_SYCL_TENSOR_BINARY_OP_LOCALS
|
||||||
|
|
||||||
|
const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
|
||||||
|
const int block_num_x = (ne00 + 2 * SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2 * SYCL_GET_ROWS_BLOCK_SIZE);
|
||||||
|
const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
|
||||||
|
|
||||||
|
// strides in elements
|
||||||
|
//const size_t s0 = nb0 / ggml_element_size(dst);
|
||||||
|
const size_t s1 = nb1 / ggml_element_size(dst);
|
||||||
|
const size_t s2 = nb2 / ggml_element_size(dst);
|
||||||
|
const size_t s3 = nb3 / ggml_element_size(dst);
|
||||||
|
|
||||||
|
const size_t s10 = nb10 / ggml_element_size(dst->src[1]);
|
||||||
|
const size_t s11 = nb11 / ggml_element_size(dst->src[1]);
|
||||||
|
const size_t s12 = nb12 / ggml_element_size(dst->src[1]);
|
||||||
|
//const size_t s13 = nb13 / ggml_element_size(dst->src[1]);
|
||||||
|
|
||||||
|
GGML_ASSERT(ne00 % 2 == 0);
|
||||||
|
const void * src0_dd = dst->src[0]->data;
|
||||||
|
const int32_t * src1_dd = static_cast<const int32_t *>(dst->src[1]->data);
|
||||||
|
float * dst_dd = static_cast<float *>(dst->data);
|
||||||
|
|
||||||
|
dpct::queue_ptr stream = ctx.stream();
|
||||||
|
|
||||||
|
stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) {
|
||||||
|
k_get_rows<qk, qr, dq>(src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2, s3, nb01, nb02, nb03, s10, s11, s12,
|
||||||
|
item_ct1);
|
||||||
|
});
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename src0_t> static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
|
||||||
|
GGML_SYCL_TENSOR_BINARY_OP_LOCALS
|
||||||
|
|
||||||
|
const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
|
||||||
|
const int block_num_x = (ne00 + SYCL_GET_ROWS_BLOCK_SIZE - 1) / SYCL_GET_ROWS_BLOCK_SIZE;
|
||||||
|
const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
|
||||||
|
|
||||||
|
// strides in elements
|
||||||
|
//const size_t s0 = nb0 / ggml_element_size(dst);
|
||||||
|
const size_t s1 = nb1 / ggml_element_size(dst);
|
||||||
|
const size_t s2 = nb2 / ggml_element_size(dst);
|
||||||
|
const size_t s3 = nb3 / ggml_element_size(dst);
|
||||||
|
|
||||||
|
const size_t s10 = nb10 / ggml_element_size(dst->src[1]);
|
||||||
|
const size_t s11 = nb11 / ggml_element_size(dst->src[1]);
|
||||||
|
const size_t s12 = nb12 / ggml_element_size(dst->src[1]);
|
||||||
|
//const size_t s13 = nb13 / ggml_element_size(dst->src[1]);
|
||||||
|
const src0_t * src0_dd = static_cast<const src0_t *>(dst->src[0]->data);
|
||||||
|
const int32_t * src1_dd = static_cast<const int32_t *>(dst->src[1]->data);
|
||||||
|
float * dst_dd = static_cast<float *>(dst->data);
|
||||||
|
|
||||||
|
dpct::queue_ptr stream = ctx.stream();
|
||||||
|
|
||||||
|
{
|
||||||
|
dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
|
||||||
|
|
||||||
|
stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) {
|
||||||
|
k_get_rows_float(src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2, s3, nb01, nb02, nb03, s10, s11, s12,
|
||||||
|
item_ct1);
|
||||||
|
});
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
|
||||||
|
GGML_ASSERT(dst->src[1]->type == GGML_TYPE_I32);
|
||||||
|
GGML_ASSERT(dst->type == GGML_TYPE_F32);
|
||||||
|
|
||||||
|
GGML_ASSERT(dst->src[0]->nb[0] == ggml_type_size(dst->src[0]->type));
|
||||||
|
GGML_ASSERT(dst->src[1]->nb[0] == ggml_type_size(dst->src[1]->type));
|
||||||
|
GGML_ASSERT(dst->nb[0] == ggml_type_size(dst->type));
|
||||||
|
GGML_ASSERT(strcmp(dst->src[1]->buffer->buft->iface.get_name(dst->src[1]->buffer->buft), GGML_SYCL_NAME "_Split") != 0);
|
||||||
|
GGML_ASSERT(strcmp(dst->buffer->buft->iface.get_name(dst->buffer->buft), GGML_SYCL_NAME "_Split") != 0);
|
||||||
|
|
||||||
|
switch (dst->src[0]->type) {
|
||||||
|
case GGML_TYPE_F16:
|
||||||
|
get_rows_sycl_float<sycl::half>(ctx, dst);
|
||||||
|
break;
|
||||||
|
case GGML_TYPE_F32:
|
||||||
|
get_rows_sycl_float<float>(ctx, dst);
|
||||||
|
break;
|
||||||
|
case GGML_TYPE_Q4_0:
|
||||||
|
get_rows_sycl<QK4_0, QR4_0, dequantize_q4_0>(ctx, dst);
|
||||||
|
break;
|
||||||
|
case GGML_TYPE_Q4_1:
|
||||||
|
get_rows_sycl<QK4_1, QR4_1, dequantize_q4_1>(ctx, dst);
|
||||||
|
break;
|
||||||
|
case GGML_TYPE_Q5_0:
|
||||||
|
get_rows_sycl<QK5_0, QR5_0, dequantize_q5_0>(ctx, dst);
|
||||||
|
break;
|
||||||
|
case GGML_TYPE_Q5_1:
|
||||||
|
get_rows_sycl<QK5_1, QR5_1, dequantize_q5_1>(ctx, dst);
|
||||||
|
break;
|
||||||
|
case GGML_TYPE_Q8_0:
|
||||||
|
get_rows_sycl<QK8_0, QR8_0, dequantize_q8_0>(ctx, dst);
|
||||||
|
break;
|
||||||
|
default:
|
||||||
|
// TODO: k-quants
|
||||||
|
GGML_LOG_ERROR("%s: unsupported type: %s\n", __func__, ggml_type_name(dst->src[0]->type));
|
||||||
|
GGML_ABORT("fatal error");
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
}
|
||||||
8
ggml/src/ggml-sycl/getrows.hpp
Normal file
8
ggml/src/ggml-sycl/getrows.hpp
Normal file
|
|
@ -0,0 +1,8 @@
|
||||||
|
#ifndef GGML_SYCL_GETROWS_HPP
|
||||||
|
#define GGML_SYCL_GETROWS_HPP
|
||||||
|
|
||||||
|
#include "common.hpp"
|
||||||
|
|
||||||
|
void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
|
||||||
|
|
||||||
|
#endif // GGML_SYCL_GETROWS_HPP
|
||||||
|
|
@ -1336,83 +1336,6 @@ static void quantize_q8_1(const float * __restrict__ x, void * __restrict__ vy,
|
||||||
reinterpret_cast<sycl::half &>(y[ib].ds.y()) = sum;
|
reinterpret_cast<sycl::half &>(y[ib].ds.y()) = sum;
|
||||||
}
|
}
|
||||||
|
|
||||||
template<int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
|
|
||||||
static void k_get_rows(
|
|
||||||
const void * src0, const int32_t * src1, dst_t * dst,
|
|
||||||
int64_t ne00, /*int64_t ne01, int64_t ne02, int64_t ne03,*/
|
|
||||||
/*int64_t ne10, int64_t ne11,*/ int64_t ne12, /*int64_t ne13,*/
|
|
||||||
/*size_t s0,*/ size_t s1, size_t s2, size_t s3,
|
|
||||||
/*size_t nb00,*/ size_t nb01, size_t nb02, size_t nb03,
|
|
||||||
size_t s10, size_t s11, size_t s12,
|
|
||||||
const sycl::nd_item<3> &item_ct1/*, size_t s13*/) {
|
|
||||||
|
|
||||||
const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) +
|
|
||||||
item_ct1.get_local_id(2)) *
|
|
||||||
2;
|
|
||||||
const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
|
|
||||||
item_ct1.get_local_id(1);
|
|
||||||
const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
|
|
||||||
item_ct1.get_local_id(0)) /
|
|
||||||
ne12;
|
|
||||||
const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
|
|
||||||
item_ct1.get_local_id(0)) %
|
|
||||||
ne12;
|
|
||||||
|
|
||||||
if (i00 >= ne00) {
|
|
||||||
return;
|
|
||||||
}
|
|
||||||
|
|
||||||
const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
|
|
||||||
|
|
||||||
dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
|
|
||||||
const void * src0_row = (const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03;
|
|
||||||
|
|
||||||
const int ib = i00/qk; // block index
|
|
||||||
const int iqs = (i00%qk)/qr; // quant index
|
|
||||||
const int iybs = i00 - i00%qk; // dst block start index
|
|
||||||
const int y_offset = qr == 1 ? 1 : qk/2;
|
|
||||||
|
|
||||||
// dequantize
|
|
||||||
dfloat2 v;
|
|
||||||
dequantize_kernel(src0_row, ib, iqs, v);
|
|
||||||
|
|
||||||
dst_row[iybs + iqs + 0] = v.x();
|
|
||||||
dst_row[iybs + iqs + y_offset] = v.y();
|
|
||||||
}
|
|
||||||
|
|
||||||
template<typename src0_t, typename dst_t>
|
|
||||||
static void k_get_rows_float(
|
|
||||||
const src0_t * src0, const int32_t * src1, dst_t * dst,
|
|
||||||
int64_t ne00, /*int64_t ne01, int64_t ne02, int64_t ne03,*/
|
|
||||||
/*int64_t ne10, int64_t ne11,*/ int64_t ne12, /*int64_t ne13,*/
|
|
||||||
/*size_t s0,*/ size_t s1, size_t s2, size_t s3,
|
|
||||||
/*size_t nb00,*/ size_t nb01, size_t nb02, size_t nb03,
|
|
||||||
size_t s10, size_t s11, size_t s12,
|
|
||||||
const sycl::nd_item<3> &item_ct1/*, size_t s13*/) {
|
|
||||||
|
|
||||||
const int i00 = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
|
|
||||||
item_ct1.get_local_id(2);
|
|
||||||
const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
|
|
||||||
item_ct1.get_local_id(1);
|
|
||||||
const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
|
|
||||||
item_ct1.get_local_id(0)) /
|
|
||||||
ne12;
|
|
||||||
const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
|
|
||||||
item_ct1.get_local_id(0)) %
|
|
||||||
ne12;
|
|
||||||
|
|
||||||
if (i00 >= ne00) {
|
|
||||||
return;
|
|
||||||
}
|
|
||||||
|
|
||||||
const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
|
|
||||||
|
|
||||||
dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
|
|
||||||
const src0_t * src0_row = (const src0_t *)((const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03);
|
|
||||||
|
|
||||||
dst_row[i00] = src0_row[i00];
|
|
||||||
}
|
|
||||||
|
|
||||||
static void mul_mat_p021_f16_f32(
|
static void mul_mat_p021_f16_f32(
|
||||||
const void * __restrict__ vx, const float * __restrict__ y, float * __restrict__ dst,
|
const void * __restrict__ vx, const float * __restrict__ y, float * __restrict__ dst,
|
||||||
const int ncols_x, const int nrows_x, const int nchannels_x, const int nchannels_y,
|
const int ncols_x, const int nrows_x, const int nchannels_x, const int nchannels_y,
|
||||||
|
|
@ -1644,80 +1567,6 @@ static void pool2d_nchw_kernel(
|
||||||
o_ptr[cur_oh * ow + cur_ow] = res;
|
o_ptr[cur_oh * ow + cur_ow] = res;
|
||||||
}
|
}
|
||||||
|
|
||||||
template <int qk, int qr, dequantize_kernel_t dq>
|
|
||||||
static void get_rows_sycl(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
|
|
||||||
|
|
||||||
GGML_SYCL_TENSOR_BINARY_OP_LOCALS
|
|
||||||
|
|
||||||
const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
|
|
||||||
const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE);
|
|
||||||
const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
|
|
||||||
|
|
||||||
// strides in elements
|
|
||||||
//const size_t s0 = nb0 / ggml_element_size(dst);
|
|
||||||
const size_t s1 = nb1 / ggml_element_size(dst);
|
|
||||||
const size_t s2 = nb2 / ggml_element_size(dst);
|
|
||||||
const size_t s3 = nb3 / ggml_element_size(dst);
|
|
||||||
|
|
||||||
const size_t s10 = nb10 / ggml_element_size(dst->src[1]);
|
|
||||||
const size_t s11 = nb11 / ggml_element_size(dst->src[1]);
|
|
||||||
const size_t s12 = nb12 / ggml_element_size(dst->src[1]);
|
|
||||||
//const size_t s13 = nb13 / ggml_element_size(dst->src[1]);
|
|
||||||
|
|
||||||
GGML_ASSERT(ne00 % 2 == 0);
|
|
||||||
const void * src0_dd = dst->src[0]->data;
|
|
||||||
const int32_t * src1_dd = static_cast<const int32_t *>(dst->src[1]->data);
|
|
||||||
float * dst_dd = static_cast<float *>(dst->data);
|
|
||||||
|
|
||||||
dpct::queue_ptr stream = ctx.stream();
|
|
||||||
|
|
||||||
stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
|
|
||||||
[=](sycl::nd_item<3> item_ct1) {
|
|
||||||
k_get_rows<qk, qr, dq>(
|
|
||||||
src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2,
|
|
||||||
s3, nb01, nb02, nb03, s10, s11, s12, item_ct1);
|
|
||||||
});
|
|
||||||
|
|
||||||
}
|
|
||||||
|
|
||||||
template <typename src0_t>
|
|
||||||
static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
|
|
||||||
|
|
||||||
GGML_SYCL_TENSOR_BINARY_OP_LOCALS
|
|
||||||
|
|
||||||
const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
|
|
||||||
const int block_num_x = (ne00 + SYCL_GET_ROWS_BLOCK_SIZE - 1) / SYCL_GET_ROWS_BLOCK_SIZE;
|
|
||||||
const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
|
|
||||||
|
|
||||||
// strides in elements
|
|
||||||
//const size_t s0 = nb0 / ggml_element_size(dst);
|
|
||||||
const size_t s1 = nb1 / ggml_element_size(dst);
|
|
||||||
const size_t s2 = nb2 / ggml_element_size(dst);
|
|
||||||
const size_t s3 = nb3 / ggml_element_size(dst);
|
|
||||||
|
|
||||||
const size_t s10 = nb10 / ggml_element_size(dst->src[1]);
|
|
||||||
const size_t s11 = nb11 / ggml_element_size(dst->src[1]);
|
|
||||||
const size_t s12 = nb12 / ggml_element_size(dst->src[1]);
|
|
||||||
//const size_t s13 = nb13 / ggml_element_size(dst->src[1]);
|
|
||||||
const src0_t * src0_dd = static_cast<const src0_t *>(dst->src[0]->data);
|
|
||||||
const int32_t * src1_dd = static_cast<const int32_t *>(dst->src[1]->data);
|
|
||||||
float * dst_dd = static_cast<float *>(dst->data);
|
|
||||||
|
|
||||||
dpct::queue_ptr stream = ctx.stream();
|
|
||||||
|
|
||||||
{
|
|
||||||
dpct::has_capability_or_fail(stream->get_device(),
|
|
||||||
{sycl::aspect::fp16});
|
|
||||||
|
|
||||||
stream->parallel_for(
|
|
||||||
sycl::nd_range<3>(block_nums * block_dims, block_dims),
|
|
||||||
[=](sycl::nd_item<3> item_ct1) {
|
|
||||||
k_get_rows_float(src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2,
|
|
||||||
s3, nb01, nb02, nb03, s10, s11, s12, item_ct1);
|
|
||||||
});
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
static void quantize_row_q8_1_sycl(const float *x, void *vy, const int kx,
|
static void quantize_row_q8_1_sycl(const float *x, void *vy, const int kx,
|
||||||
const int ky, const int kx_padded,
|
const int ky, const int kx_padded,
|
||||||
queue_ptr stream) {
|
queue_ptr stream) {
|
||||||
|
|
@ -1912,46 +1761,6 @@ catch (sycl::exception const &exc) {
|
||||||
std::exit(1);
|
std::exit(1);
|
||||||
}
|
}
|
||||||
|
|
||||||
static void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
|
|
||||||
|
|
||||||
GGML_ASSERT(dst->src[1]->type == GGML_TYPE_I32);
|
|
||||||
GGML_ASSERT(dst->type == GGML_TYPE_F32);
|
|
||||||
|
|
||||||
GGML_ASSERT(dst->src[0]->nb[0] == ggml_type_size(dst->src[0]->type));
|
|
||||||
GGML_ASSERT(dst->src[1]->nb[0] == ggml_type_size(dst->src[1]->type));
|
|
||||||
GGML_ASSERT(dst->nb[0] == ggml_type_size(dst->type));
|
|
||||||
GGML_ASSERT(!ggml_backend_buffer_is_sycl_split(dst->src[1]->buffer));
|
|
||||||
GGML_ASSERT(!ggml_backend_buffer_is_sycl_split(dst->buffer));
|
|
||||||
|
|
||||||
switch (dst->src[0]->type) {
|
|
||||||
case GGML_TYPE_F16:
|
|
||||||
get_rows_sycl_float<sycl::half>(ctx, dst);
|
|
||||||
break;
|
|
||||||
case GGML_TYPE_F32:
|
|
||||||
get_rows_sycl_float<float>(ctx, dst);
|
|
||||||
break;
|
|
||||||
case GGML_TYPE_Q4_0:
|
|
||||||
get_rows_sycl<QK4_0, QR4_0, dequantize_q4_0>(ctx, dst);
|
|
||||||
break;
|
|
||||||
case GGML_TYPE_Q4_1:
|
|
||||||
get_rows_sycl<QK4_1, QR4_1, dequantize_q4_1>(ctx, dst);
|
|
||||||
break;
|
|
||||||
case GGML_TYPE_Q5_0:
|
|
||||||
get_rows_sycl<QK5_0, QR5_0, dequantize_q5_0>(ctx, dst);
|
|
||||||
break;
|
|
||||||
case GGML_TYPE_Q5_1:
|
|
||||||
get_rows_sycl<QK5_1, QR5_1, dequantize_q5_1>(ctx, dst);
|
|
||||||
break;
|
|
||||||
case GGML_TYPE_Q8_0:
|
|
||||||
get_rows_sycl<QK8_0, QR8_0, dequantize_q8_0>(ctx, dst);
|
|
||||||
break;
|
|
||||||
default:
|
|
||||||
// TODO: k-quants
|
|
||||||
GGML_LOG_ERROR("%s: unsupported type: %s\n", __func__, ggml_type_name(dst->src[0]->type));
|
|
||||||
GGML_ABORT("fatal error");
|
|
||||||
break;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
inline void ggml_sycl_op_mul_mat_sycl(
|
inline void ggml_sycl_op_mul_mat_sycl(
|
||||||
ggml_backend_sycl_context & ctx,
|
ggml_backend_sycl_context & ctx,
|
||||||
|
|
|
||||||
Loading…
Add table
Add a link
Reference in a new issue