Skip to content

Commit d262ec6

Browse files
committed
applied Romain's comments
1 parent 6ad1ed3 commit d262ec6

File tree

2 files changed

+5
-8
lines changed

2 files changed

+5
-8
lines changed

ggml/src/ggml-sycl/ggml-sycl.cpp

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2824,21 +2824,20 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
28242824
dnn_gemm(src1_f16, src0_f16, dst_ddf, ne12*ne13, ne02 * ne03);
28252825
}
28262826
else {
2827-
28282827
for (int64_t ie03 = 0; ie03 < ne03; ++ie03) {
2829-
const sycl::half* src0_f16_shifted = src0_f16 + ((ie03*nb03)/2); // div2 cuz nb is in bytes and pointer is in f16 (2 bytes)
2828+
const sycl::half* src0_f16_shifted = src0_f16 + ((ie03*nb03)/sizeof(sycl::half)); // nb is in bytes
28302829
const sycl::half* src1_f16_shifted = src1_f16 + ie03*s13;
2831-
float* dst_shifted = dst_ddf + ((ie03*nb3)/4); // div4 cuz nb is in bytes and pointer is float (4 bytes)
2830+
float* dst_shifted = dst_ddf + ((ie03*nb3)/sizeof(float));
28322831
dnn_gemm(src1_f16_shifted, src0_f16_shifted, dst_shifted, ne12, ne02);
28332832
}
28342833
}
28352834
} else {
28362835
// iterate over batches from smaller set of matrices (matrix 0)
28372836
for (int64_t ie02 = 0; ie02 < ne02; ++ie02) {
28382837
for (int64_t ie03 = 0; ie03 < ne03; ++ie03) {
2839-
const sycl::half* src0_f16_shifted = src0_f16 + ((ie02*nb02 + ie03*nb03)/2); // div2 cuz nb is in bytes and pointer is in f16 (2 bytes)
2838+
const sycl::half* src0_f16_shifted = src0_f16 + ((ie02*nb02 + ie03*nb03)/sizeof(sycl::half));
28402839
const sycl::half* src1_f16_shifted = src1_f16 + ie02*s12*r2 + ie03*s13*r3;
2841-
float* dst_shifted = dst_ddf + ((ie02*nb2*r2 + ie03*nb3*r3)/4); // div4 cuz nb is in bytes and pointer is float (4 bytes)
2840+
float* dst_shifted = dst_ddf + ((ie02*nb2*r2 + ie03*nb3*r3)/sizeof(float));
28422841
dnn_gemm(src1_f16_shifted, src0_f16_shifted, dst_shifted, r2*r3, 1);
28432842
}
28442843
}

tests/test-backend-ops.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3919,7 +3919,7 @@ static const ggml_type other_types[] = {
39193919
// Test cases for evaluation: should try to cover edge cases while using small input sizes to keep the runtime low
39203920
static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
39213921
std::vector<std::unique_ptr<test_case>> test_cases;
3922-
[[maybe_unused]] std::default_random_engine rng(0);
3922+
std::default_random_engine rng(0);
39233923

39243924
// unary ops
39253925
for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) {
@@ -4242,8 +4242,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
42424242
test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 16, 256, {2, 3}, {1, 1}, {0, 1, 3, 2}));
42434243
test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 16, 256, {2, 3}, {1, 1}, {0, 3, 2, 1}));
42444244

4245-
test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 16, 256, {2, 1}, {1, 1}, {0, 2, 1, 3}));
4246-
42474245
// test cases with large ne00/ne10 to cover stream-k fixup
42484246
test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 1, 1024, {3, 2}, {1, 1}));
42494247
test_cases.emplace_back(new test_mul_mat(type_a, type_b, 16, 8, 1024, {3, 2}, {1, 1}));

0 commit comments

Comments
 (0)