Update ggml_sycl_mul_mat_batched_sycl
This commit is contained in:
Aidan
parent
9d05e6a0aa
commit
88c272a622
1 changed files with 94 additions and 99 deletions
127
ggml-sycl.cpp
127
ggml-sycl.cpp
|
|
@ -2560,8 +2560,7 @@ namespace dpct
|
|||
library_data_t a_type, int lda, const void *b[],
|
||||
library_data_t b_type, int ldb, const void *beta,
|
||||
void *c[], library_data_t c_type, int ldc,
|
||||
int batch_size, library_data_t scaling_type)
|
||||
{
|
||||
int batch_size, library_data_t scaling_type) try {
|
||||
#ifdef DPCT_USM_LEVEL_NONE
|
||||
throw std::runtime_error("this API is unsupported when USM level is none");
|
||||
#else
|
||||
|
|
@ -2700,6 +2699,10 @@ namespace dpct
|
|||
throw std::runtime_error("the combination of data type is unsupported");
|
||||
}
|
||||
#endif
|
||||
} catch (sycl::exception const &exc) {
|
||||
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
|
||||
<< ", line:" << __LINE__ << std::endl;
|
||||
std::exit(1);
|
||||
}
|
||||
|
||||
/// Computes a batch of matrix-matrix product with general matrices.
|
||||
|
|
@ -2732,7 +2735,7 @@ namespace dpct
|
|||
library_data_t b_type, int ldb, long long int stride_b,
|
||||
const void *beta, void *c, library_data_t c_type,
|
||||
int ldc, long long int stride_c, int batch_size,
|
||||
library_data_t scaling_type)
|
||||
library_data_t scaling_type) try
|
||||
{
|
||||
if (scaling_type == library_data_t::real_float &&
|
||||
c_type == library_data_t::complex_float)
|
||||
|
|
@ -2865,6 +2868,10 @@ namespace dpct
|
|||
default:
|
||||
throw std::runtime_error("the combination of data type is unsupported");
|
||||
}
|
||||
} catch (sycl::exception const &exc) {
|
||||
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
|
||||
<< ", line:" << __LINE__ << std::endl;
|
||||
std::exit(1);
|
||||
}
|
||||
|
||||
static inline void
|
||||
|
|
@ -13330,11 +13337,11 @@ static void k_compute_batched_ptrs(const sycl::half *src0_as_f16,
|
|||
int64_t i02 = i12 / r2;
|
||||
|
||||
ptrs_src[0*ne23 + i12 + i13*ne12] = (const char *) src0_as_f16 + i02*nb02 + i03*nb03;
|
||||
ptrs_src[1*ne23 + i12 + i13*ne12] = (const char *) src1_as_f16 + i12*nb12/2 + i13*nb13/2;
|
||||
ptrs_src[1*ne23 + i12 + i13*ne12] = (const char *) src1_as_f16 + i12*nb12 + i13*nb13;
|
||||
ptrs_dst[0*ne23 + i12 + i13*ne12] = ( char *) dst + i12*nbd2 + i13*nbd3;
|
||||
}
|
||||
|
||||
static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
||||
static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0,
|
||||
const ggml_tensor *src1,
|
||||
ggml_tensor *dst) try {
|
||||
GGML_ASSERT(!ggml_is_transposed(src0));
|
||||
|
|
@ -13342,18 +13349,10 @@ static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
|||
|
||||
GGML_ASSERT(src0->backend != GGML_BACKEND_GPU_SPLIT);
|
||||
GGML_ASSERT(src0->type == GGML_TYPE_F16);
|
||||
GGML_ASSERT(src1->type == GGML_TYPE_F32);
|
||||
|
||||
GGML_TENSOR_LOCALS(int64_t, ne0, src0, ne);
|
||||
GGML_TENSOR_BINARY_OP_LOCALS
|
||||
|
||||
GGML_TENSOR_LOCALS(int64_t, nb0, src0, nb);
|
||||
|
||||
GGML_TENSOR_LOCALS(int64_t, ne1, src1, ne);
|
||||
|
||||
GGML_TENSOR_LOCALS(int64_t, nb1, src1, nb);
|
||||
|
||||
const int64_t ne1 = ggml_nelements(src1);
|
||||
const int64_t ne = ggml_nelements(dst);
|
||||
const int64_t ne_dst = ggml_nelements(dst);
|
||||
|
||||
SYCL_CHECK(ggml_sycl_set_device(g_main_device));
|
||||
dpct::queue_ptr main_stream = g_syclStreams[g_main_device_index][0];
|
||||
|
|
@ -13361,25 +13360,31 @@ static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
|||
SYCL_CHECK(
|
||||
CHECK_TRY_ERROR(g_sycl_handles[g_main_device_index] = main_stream));
|
||||
|
||||
ggml_tensor_extra_gpu * src0_extra = (ggml_tensor_extra_gpu *) src0->extra;
|
||||
void * src0_ddq = src0_extra->data_device[g_main_device_index];
|
||||
sycl::half *src0_as_f16 = (sycl::half *)src0_ddq;
|
||||
ggml_tensor_extra_gpu *src0_extra = (ggml_tensor_extra_gpu *)src0->extra;
|
||||
void *src0_ddq = src0_extra->data_device[g_main_device_index];
|
||||
sycl::half *src0_f16 = (sycl::half *)src0_ddq;
|
||||
|
||||
ggml_tensor_extra_gpu * src1_extra = (ggml_tensor_extra_gpu *) src1->extra;
|
||||
float * src1_ddf = (float *) src1_extra->data_device[g_main_device_index];
|
||||
ggml_tensor_extra_gpu *src1_extra = (ggml_tensor_extra_gpu *)src1->extra;
|
||||
float *src1_ddf = (float *)src1_extra->data_device[g_main_device_index];
|
||||
|
||||
ggml_tensor_extra_gpu * dst_extra = (ggml_tensor_extra_gpu *) dst->extra;
|
||||
float * dst_ddf = (float *) dst_extra->data_device[g_main_device_index];
|
||||
ggml_tensor_extra_gpu *dst_extra = (ggml_tensor_extra_gpu *)dst->extra;
|
||||
float *dst_ddf = (float *)dst_extra->data_device[g_main_device_index];
|
||||
|
||||
// convert src1 to fp16
|
||||
sycl_pool_alloc<sycl::half> src1_f16_alloc;
|
||||
if (src1->type != GGML_TYPE_F16) {
|
||||
const to_fp16_sycl_t to_fp16_sycl = ggml_get_to_fp16_sycl(src1->type);
|
||||
const int64_t ne_src1 = ggml_nelements(src1);
|
||||
src1_f16_alloc.alloc(ne_src1);
|
||||
GGML_ASSERT(to_fp16_sycl != nullptr);
|
||||
|
||||
sycl_pool_alloc<sycl::half> src1_as_f16(ne1);
|
||||
to_fp16_sycl(src1_ddf, src1_as_f16.get(), ne1, main_stream);
|
||||
to_fp16_sycl(src1_ddf, src1_f16_alloc.get(), ne_src1, main_stream);
|
||||
main_stream->wait_and_throw();
|
||||
}
|
||||
sycl::half *src1_f16 = src1->type == GGML_TYPE_F16 ? (sycl::half *)src1_ddf
|
||||
: src1_f16_alloc.get();
|
||||
|
||||
sycl_pool_alloc<sycl::half> dst_f16;
|
||||
char * dst_t;
|
||||
char *dst_t;
|
||||
|
||||
dpct::library_data_t cu_compute_type = dpct::library_data_t::real_half;
|
||||
dpct::library_data_t cu_data_type = dpct::library_data_t::real_half;
|
||||
|
|
@ -13394,16 +13399,16 @@ static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
|||
const float alpha_f32 = 1.0f;
|
||||
const float beta_f32 = 0.0f;
|
||||
|
||||
const void * alpha = &alpha_f16;
|
||||
const void * beta = &beta_f16;
|
||||
const void *alpha = &alpha_f16;
|
||||
const void *beta = &beta_f16;
|
||||
|
||||
if (dst->op_params[0] == GGML_PREC_DEFAULT) {
|
||||
dst_t = (char *) dst_f16.alloc(ne);
|
||||
dst_t = (char *)dst_f16.alloc(ne_dst);
|
||||
|
||||
nbd2 /= sizeof(float) / sizeof(sycl::half);
|
||||
nbd3 /= sizeof(float) / sizeof(sycl::half);
|
||||
} else {
|
||||
dst_t = (char *) dst_ddf;
|
||||
dst_t = (char *)dst_ddf;
|
||||
|
||||
cu_compute_type = dpct::library_data_t::real_float;
|
||||
cu_data_type = dpct::library_data_t::real_float;
|
||||
|
|
@ -13416,8 +13421,8 @@ static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
|||
GGML_ASSERT(ne13 % ne03 == 0);
|
||||
|
||||
// broadcast factors
|
||||
const int64_t r2 = ne12/ne02;
|
||||
const int64_t r3 = ne13/ne03;
|
||||
const int64_t r2 = ne12 / ne02;
|
||||
const int64_t r3 = ne13 / ne03;
|
||||
|
||||
#if 0
|
||||
// use syclGemmEx
|
||||
|
|
@ -13441,22 +13446,20 @@ static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
|||
#else
|
||||
if (r2 == 1 && r3 == 1 && src0->nb[2]*src0->ne[2] == src0->nb[3] && src1->nb[2]*src1->ne[2] == src1->nb[3]) {
|
||||
// there is no broadcast and src0, src1 are contiguous across dims 2, 3
|
||||
// use syclGemmStridedBatchedEx
|
||||
// use dpct::gemm_batch
|
||||
SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm_batch(
|
||||
*g_sycl_handles[g_main_device_index], oneapi::mkl::transpose::trans,
|
||||
oneapi::mkl::transpose::nontrans, ne01, ne11, ne10, alpha,
|
||||
(const char *)src0_as_f16, dpct::library_data_t::real_half,
|
||||
nb01 / sizeof(sycl::half), src0->nb[2] / sizeof(sycl::half),
|
||||
(const char *)src1_as_f16.get(), dpct::library_data_t::real_half,
|
||||
nb11 / sizeof(float), src1->nb[2] / sizeof(float), beta,
|
||||
(char *)dst_t, cu_data_type, ne01, dst->nb[2] / sizeof(float),
|
||||
ne12 * ne13, cu_compute_type)));
|
||||
(const char *)src0_f16, dpct::library_data_t::real_half, nb01 / nb00,
|
||||
nb02 / nb00, (const char *)src1_f16, dpct::library_data_t::real_half,
|
||||
nb11 / nb10, nb12 / nb10, beta, (char *)dst_t, cu_data_type, ne01,
|
||||
nb2 / nb0, ne12 * ne13, cu_compute_type)));
|
||||
} else {
|
||||
// use syclGemmBatchedEx
|
||||
const int ne23 = ne12*ne13;
|
||||
// use dpct::gemm_batch
|
||||
const int ne23 = ne12 * ne13;
|
||||
|
||||
sycl_pool_alloc<const void *> ptrs_src(2*ne23);
|
||||
sycl_pool_alloc< void *> ptrs_dst(1*ne23);
|
||||
sycl_pool_alloc<const void *> ptrs_src(2 * ne23);
|
||||
sycl_pool_alloc<void *> ptrs_dst(1 * ne23);
|
||||
|
||||
sycl::range<3> block_dims(1, ne12, ne13);
|
||||
/*
|
||||
|
|
@ -13469,35 +13472,28 @@ static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
|||
{sycl::aspect::fp16});
|
||||
|
||||
main_stream->submit([&](sycl::handler &cgh) {
|
||||
const sycl::half *src1_as_f16_get_ct1 = src1_as_f16.get();
|
||||
const void **ptrs_src_get_ct3 = ptrs_src.get();
|
||||
void **ptrs_dst_get_ct4 = ptrs_dst.get();
|
||||
|
||||
const void **ptrs_src_get = ptrs_src.get();
|
||||
void **ptrs_dst_get = ptrs_dst.get();
|
||||
auto nb12_scaled = src1->type == GGML_TYPE_F16 ? nb12 : nb12 / 2;
|
||||
auto nb13_scaled = src1->type == GGML_TYPE_F16 ? nb13 : nb13 / 2;
|
||||
cgh.parallel_for(sycl::nd_range<3>(block_dims, block_dims),
|
||||
[=](sycl::nd_item<3> item_ct1) {
|
||||
k_compute_batched_ptrs(
|
||||
src0_as_f16, src1_as_f16_get_ct1,
|
||||
dst_t, ptrs_src_get_ct3,
|
||||
ptrs_dst_get_ct4, ne12, ne13, ne23,
|
||||
nb02, nb03, nb12, nb13, nbd2, nbd3, r2,
|
||||
r3, item_ct1);
|
||||
src0_f16, src1_f16, dst_t, ptrs_src_get,
|
||||
ptrs_dst_get, ne12, ne13, ne23, nb02, nb03,
|
||||
nb12_scaled, nb13_scaled, nbd2, nbd3, r2, r3,
|
||||
item_ct1);
|
||||
});
|
||||
});
|
||||
}
|
||||
/*
|
||||
DPCT1010:95: SYCL uses exceptions to report errors and does not use the
|
||||
error codes. The call was replaced with 0. You need to rewrite this
|
||||
code.
|
||||
*/
|
||||
SYCL_CHECK(0);
|
||||
|
||||
SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm_batch(
|
||||
*g_sycl_handles[g_main_device_index], oneapi::mkl::transpose::trans,
|
||||
oneapi::mkl::transpose::nontrans, ne01, ne11, ne10, alpha,
|
||||
(const void **)(ptrs_src.get() + 0 * ne23),
|
||||
dpct::library_data_t::real_half, nb01 / sizeof(sycl::half),
|
||||
dpct::library_data_t::real_half, nb01 / nb00,
|
||||
(const void **)(ptrs_src.get() + 1 * ne23),
|
||||
dpct::library_data_t::real_half, nb11 / sizeof(float), beta,
|
||||
dpct::library_data_t::real_half, nb11 / nb10, beta,
|
||||
(void **)(ptrs_dst.get() + 0 * ne23), cu_data_type, ne01, ne23,
|
||||
cu_compute_type)));
|
||||
}
|
||||
|
|
@ -13505,10 +13501,9 @@ static void ggml_sycl_mul_mat_mat_batched_sycl(const ggml_tensor *src0,
|
|||
|
||||
if (dst->op_params[0] == GGML_PREC_DEFAULT) {
|
||||
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16);
|
||||
to_fp32_sycl(dst_f16.get(), dst_ddf, ne, main_stream);
|
||||
to_fp32_sycl(dst_f16.get(), dst_ddf, ne_dst, main_stream);
|
||||
}
|
||||
}
|
||||
catch (sycl::exception const &exc) {
|
||||
} catch (sycl::exception const &exc) {
|
||||
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
|
||||
<< ", line:" << __LINE__ << std::endl;
|
||||
std::exit(1);
|
||||
|
|
@ -13551,10 +13546,10 @@ static void ggml_sycl_mul_mat(const ggml_tensor * src0, const ggml_tensor * src1
|
|||
// KQV single-batch
|
||||
// GGML_SYCL_DEBUG("ggml_sycl_mul_mat_vec_nc\n");
|
||||
ggml_sycl_mul_mat_vec_nc(src0, src1, dst);
|
||||
} else if (!split && all_on_device && use_xmx && src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && !ggml_is_transposed(src0) && !ggml_is_transposed(src1)) {
|
||||
} else if (!split && all_on_device && use_xmx && src0->type == GGML_TYPE_F16 && !ggml_is_transposed(src0) && !ggml_is_transposed(src1)) {
|
||||
// KQ + KQV multi-batch
|
||||
// GGML_SYCL_DEBUG("ggml_sycl_mul_mat_mat_batched_sycl\n");
|
||||
ggml_sycl_mul_mat_mat_batched_sycl(src0, src1, dst);
|
||||
// GGML_SYCL_DEBUG("ggml_sycl_mul_mat_batched_sycl\n");
|
||||
ggml_sycl_mul_mat_batched_sycl(src0, src1, dst);
|
||||
} else if (src0->type == GGML_TYPE_F32) {
|
||||
// GGML_SYCL_DEBUG("ggml_sycl_op_mul_mat\n");
|
||||
ggml_sycl_op_mul_mat(src0, src1, dst, ggml_sycl_op_mul_mat_sycl, false);
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue