Skip to content

Commit 258b912

Browse files
jpohhhhjpohhhh
committed
Merge branch 'master' of https://github.com/ggml-org/llama.cpp into phi4_tools_clean
2 parents 094f607 + f4c3dd5 commit 258b912

File tree

4 files changed

+76
-38
lines changed

4 files changed

+76
-38
lines changed

‎common/arg.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1889,7 +1889,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
18891889
[](common_params & params, const std::string & value) {
18901890
params.out_file = value;
18911891
}
1892-
).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_CVECTOR_GENERATOR, LLAMA_EXAMPLE_EXPORT_LORA}));
1892+
).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_CVECTOR_GENERATOR, LLAMA_EXAMPLE_EXPORT_LORA, LLAMA_EXAMPLE_TTS}));
18931893
add_opt(common_arg(
18941894
{"-ofreq", "--output-frequency"}, "N",
18951895
string_format("output the imatrix every N iterations (default: %d)", params.n_out_freq),

‎examples/tts/tts.cpp‎

Lines changed: 12 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -87,11 +87,11 @@ struct wav_header {
8787
uint32_t data_size;
8888
};
8989

90-
static void save_wav16(const std::string & fname, const std::vector<float> & data, int sample_rate) {
90+
static bool save_wav16(const std::string & fname, const std::vector<float> & data, int sample_rate) {
9191
std::ofstream file(fname, std::ios::binary);
9292
if (!file) {
93-
LOG_ERR("%s: Failed to open file '%s' for writing", __func__, fname.c_str());
94-
return;
93+
LOG_ERR("%s: Failed to open file '%s' for writing.\n", __func__, fname.c_str());
94+
return false;
9595
}
9696

9797
wav_header header;
@@ -108,7 +108,7 @@ static void save_wav16(const std::string & fname, const std::vector<float> & dat
108108
file.write(reinterpret_cast<const char*>(&pcm_sample), sizeof(pcm_sample));
109109
}
110110

111-
file.close();
111+
return file.good();
112112
}
113113

114114
static void fill_hann_window(int length, bool periodic, float * output) {
@@ -536,6 +536,7 @@ static std::string audio_data_from_speaker(json speaker, const outetts_version t
536536
int main(int argc, char ** argv) {
537537
common_params params;
538538

539+
params.out_file = "output.wav";
539540
params.prompt = "";
540541

541542
params.n_predict = 4096;
@@ -1060,8 +1061,6 @@ lovely<|t_0.56|><|code_start|><|634|><|596|><|1766|><|1556|><|1306|><|1285|><|14
10601061
}
10611062
#endif
10621063

1063-
const std::string fname = "output.wav";
1064-
10651064
const int n_sr = 24000; // sampling rate
10661065

10671066
// zero out first 0.25 seconds
@@ -1072,11 +1071,15 @@ lovely<|t_0.56|><|code_start|><|634|><|596|><|1766|><|1556|><|1306|><|1285|><|14
10721071
LOG_INF("%s: time for spectral ops: %.3f ms\n", __func__, (ggml_time_us() - t_spec_start) / 1000.0f);
10731072
LOG_INF("%s: total time: %.3f ms\n", __func__, (ggml_time_us() - t_main_start) / 1000.0f);
10741073

1075-
save_wav16(fname, audio, n_sr);
1074+
int retval = 0;
10761075

1077-
LOG_INF("%s: audio written to file '%s'\n", __func__, fname.c_str());
1076+
if (save_wav16(params.out_file, audio, n_sr)) {
1077+
LOG_INF("%s: audio written to file '%s'\n", __func__, params.out_file.c_str());
1078+
} else {
1079+
retval = ENOENT;
1080+
}
10781081

10791082
llama_backend_free();
10801083

1081-
return 0;
1084+
return retval;
10821085
}

‎ggml/src/ggml-sycl/common.hpp‎

Lines changed: 61 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -474,6 +474,7 @@ static void k_bin_bcast(const src0_t * src0, const src1_t * src1, dst_t * dst,
474474
int ne0, int ne1, int ne2, int ne3,
475475
int ne10, int ne11, int ne12, int ne13,
476476
/*int s0, */ int s1, int s2, int s3,
477+
/*int s00,*/ int s01, int s02, int s03,
477478
/*int s10,*/ int s11, int s12, int s13,
478479
const sycl::nd_item<3> &item_ct1) {
479480
const int i0s = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
@@ -495,9 +496,9 @@ static void k_bin_bcast(const src0_t * src0, const src1_t * src1, dst_t * dst,
495496
const int i12 = i2 % ne12;
496497
const int i13 = i3 % ne13;
497498

498-
const size_t i_src0 = i3*s3 + i2*s2 + i1*s1;
499+
const size_t i_src0 = i3*s03 + i2*s02 + i1*s01;
499500
const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
500-
const size_t i_dst = i_src0;
501+
const size_t i_dst = i3*s3 + i2*s2 + i1*s1;
501502

502503
const src0_t * src0_row = src0 + i_src0;
503504
const src1_t * src1_row = src1 + i_src1;
@@ -515,6 +516,7 @@ static void k_bin_bcast_unravel(const src0_t * src0, const src1_t * src1, dst_t
515516
int ne0, int ne1, int ne2, int ne3,
516517
int ne10, int ne11, int ne12, int ne13,
517518
/*int s0, */ int s1, int s2, int s3,
519+
/*int s00,*/ int s01, int s02, int s03,
518520
/*int s10,*/ int s11, int s12, int s13,
519521
const sycl::nd_item<3> &item_ct1) {
520522

@@ -534,9 +536,9 @@ static void k_bin_bcast_unravel(const src0_t * src0, const src1_t * src1, dst_t
534536
const int i12 = i2 % ne12;
535537
const int i13 = i3 % ne13;
536538

537-
const size_t i_src0 = i3*s3 + i2*s2 + i1*s1;
539+
const size_t i_src0 = i3*s03 + i2*s02 + i1*s01;
538540
const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
539-
const size_t i_dst = i_src0;
541+
const size_t i_dst = i3*s3 + i2*s2 + i1*s1;
540542

541543
const src0_t * src0_row = src0 + i_src0;
542544
const src1_t * src1_row = src1 + i_src1;
@@ -566,9 +568,11 @@ struct bin_bcast_sycl {
566568
int nr[4] = { nr0, nr1, nr2, nr3 };
567569

568570
// collapse dimensions until first broadcast dimension
569-
int64_t cne0[] = {ne0, ne1, ne2, ne3};
571+
int64_t cne[] = {ne0, ne1, ne2, ne3};
572+
int64_t cne0[] = {ne00, ne01, ne02, ne03};
570573
int64_t cne1[] = {ne10, ne11, ne12, ne13};
571-
size_t cnb0[] = {nb0, nb1, nb2, nb3};
574+
size_t cnb[] = {nb0, nb1, nb2, nb3};
575+
size_t cnb0[] = {nb00, nb01, nb02, nb03};
572576
size_t cnb1[] = {nb10, nb11, nb12, nb13};
573577
auto collapse = [](int64_t cne[]) {
574578
cne[0] *= cne[1];
@@ -583,32 +587,41 @@ struct bin_bcast_sycl {
583587
cnb[3] *= cne[3];
584588
};
585589

586-
for (int i = 0; i < 4; i++) {
587-
if (nr[i] != 1) {
588-
break;
589-
}
590-
if (i > 0) {
591-
collapse_nb(cnb0, cne0);
592-
collapse_nb(cnb1, cne1);
593-
collapse(cne0);
594-
collapse(cne1);
590+
if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst)) {
591+
for (int i = 0; i < 4; i++) {
592+
if (nr[i] != 1) {
593+
break;
594+
}
595+
if (i > 0) {
596+
collapse_nb(cnb, cne);
597+
collapse_nb(cnb0, cne0);
598+
collapse_nb(cnb1, cne1);
599+
collapse(cne);
600+
collapse(cne0);
601+
collapse(cne1);
602+
}
595603
}
596604
}
597605
{
598-
int64_t ne0 = cne0[0];
599-
int64_t ne1 = cne0[1];
600-
int64_t ne2 = cne0[2];
601-
int64_t ne3 = cne0[3];
606+
int64_t ne0 = cne[0];
607+
int64_t ne1 = cne[1];
608+
int64_t ne2 = cne[2];
609+
int64_t ne3 = cne[3];
602610

603611
int64_t ne10 = cne1[0];
604612
int64_t ne11 = cne1[1];
605613
int64_t ne12 = cne1[2];
606614
int64_t ne13 = cne1[3];
607615

608-
size_t nb0 = cnb0[0];
609-
size_t nb1 = cnb0[1];
610-
size_t nb2 = cnb0[2];
611-
size_t nb3 = cnb0[3];
616+
size_t nb0 = cnb[0];
617+
size_t nb1 = cnb[1];
618+
size_t nb2 = cnb[2];
619+
size_t nb3 = cnb[3];
620+
621+
size_t nb00 = cnb0[0];
622+
size_t nb01 = cnb0[1];
623+
size_t nb02 = cnb0[2];
624+
size_t nb03 = cnb0[3];
612625

613626
size_t nb10 = cnb1[0];
614627
size_t nb11 = cnb1[1];
@@ -625,6 +638,28 @@ struct bin_bcast_sycl {
625638
size_t s12 = nb12 / sizeof(src1_t);
626639
size_t s13 = nb13 / sizeof(src1_t);
627640

641+
size_t s00 = nb00 / sizeof(src0_t);
642+
size_t s01 = nb01 / sizeof(src0_t);
643+
size_t s02 = nb02 / sizeof(src0_t);
644+
size_t s03 = nb03 / sizeof(src0_t);
645+
646+
GGML_UNUSED(s00);
647+
648+
GGML_ASSERT(nb0 % sizeof(dst_t) == 0);
649+
GGML_ASSERT(nb1 % sizeof(dst_t) == 0);
650+
GGML_ASSERT(nb2 % sizeof(dst_t) == 0);
651+
GGML_ASSERT(nb3 % sizeof(dst_t) == 0);
652+
653+
GGML_ASSERT(nb00 % sizeof(src0_t) == 0);
654+
GGML_ASSERT(nb01 % sizeof(src0_t) == 0);
655+
GGML_ASSERT(nb02 % sizeof(src0_t) == 0);
656+
GGML_ASSERT(nb03 % sizeof(src0_t) == 0);
657+
658+
GGML_ASSERT(nb10 % sizeof(src1_t) == 0);
659+
GGML_ASSERT(nb11 % sizeof(src1_t) == 0);
660+
GGML_ASSERT(nb12 % sizeof(src1_t) == 0);
661+
GGML_ASSERT(nb13 % sizeof(src1_t) == 0);
662+
628663
GGML_ASSERT(s0 == 1);
629664
GGML_ASSERT(s10 == 1);
630665

@@ -661,8 +696,8 @@ struct bin_bcast_sycl {
661696
[=](sycl::nd_item<3> item_ct1) {
662697
k_bin_bcast_unravel<bin_op>(
663698
src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3,
664-
ne10, ne11, ne12, ne13, s1, s2, s3, s11, s12,
665-
s13, item_ct1);
699+
ne10, ne11, ne12, ne13, s1, s2, s3, s01, s02,
700+
s03, s11, s12, s13, item_ct1);
666701
});
667702
}
668703
} else {
@@ -680,7 +715,7 @@ struct bin_bcast_sycl {
680715
[=](sycl::nd_item<3> item_ct1) {
681716
k_bin_bcast<bin_op>(src0_dd, src1_dd, dst_dd, ne0, ne1,
682717
ne2, ne3, ne10, ne11, ne12, ne13,
683-
s1, s2, s3, s11, s12, s13,
718+
s1, s2, s3, s01, s02, s03, s11, s12, s13,
684719
item_ct1);
685720
});
686721
}

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3113,8 +3113,8 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx,
31133113
const int64_t i2 = i12;
31143114

31153115
src0_row.data = src0_original + i02*nb02;
3116-
src1_row.data = src1_original + + i11*nb11 + i12*nb12;
3117-
dst_row.data = dst_original + i1*nb1 + i2*nb2;
3116+
src1_row.data = src1_original + i11*nb11 + i12*nb12;
3117+
dst_row.data = dst_original + i1*nb1 + i2*nb2;
31183118

31193119
ggml_sycl_mul_mat(ctx, &src0_row, &src1_row, &dst_row);
31203120
}

0 commit comments

Comments
 (0)