Skip to content

Commit c6860cc

Browse files
qnixsynapseqnixsynapse
authored
SYCL: Refactor ggml_sycl_compute_forward (ggml-org#11121)
* SYCL: refactor ggml_sycl_compute_forward * SYCL: add back GGML_USED(dst) to ggml_sycl_cpy * SYCL: add function name to noop debug * SYCL: Some device info print refactoring and add details of XMX availability
1 parent 1204f97 commit c6860cc

15 files changed

+222
-229
lines changed

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,10 @@ void ggml_sycl_host_free(void* ptr) try {
5151
std::exit(1);
5252
}
5353

54+
bool gpu_has_xmx(sycl::device &dev) {
55+
return dev.has(sycl::aspect::ext_intel_matrix);
56+
}
57+
5458
int64_t downsample_sycl_global_range(int64_t accumulate_block_num, int64_t block_size) {
5559
const int64_t max_range = std::numeric_limits<int>::max();
5660
int64_t sycl_down_blk_size = block_size;

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -662,6 +662,7 @@ inline void ggml_sycl_op_bin_bcast(ggml_backend_sycl_context & ctx, const ggml_t
662662
}
663663
}
664664

665+
bool gpu_has_xmx(sycl::device &dev);
665666

666667
void ggml_sycl_op_flatten(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
667668
const ggml_tensor *src1, ggml_tensor *dst,

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

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -158,8 +158,9 @@ static void concat_f32_sycl_non_cont(
158158
});
159159
}
160160

161-
void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
162-
const ggml_tensor *src1, ggml_tensor *dst) {
161+
void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
162+
const ggml_tensor *src0 = dst->src[0];
163+
const ggml_tensor *src1 = dst->src[1];
163164
queue_ptr stream = ctx.stream();
164165

165166
const int32_t dim = ((int32_t *)dst->op_params)[0];

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,6 @@
1515

1616
#include "common.hpp"
1717

18-
void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
19-
const ggml_tensor *src1, ggml_tensor *dst);
18+
void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst);
2019

2120
#endif // GGML_SYCL_CONCAT_HPP

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

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -71,8 +71,9 @@ static void conv_transpose_1d_f32_f32_sycl(
7171
});
7272
}
7373

74-
void ggml_sycl_op_conv_transpose_1d(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
75-
const ggml_tensor *src1, ggml_tensor *dst) {
74+
void ggml_sycl_op_conv_transpose_1d(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
75+
const ggml_tensor *src0 = dst->src[0];
76+
const ggml_tensor *src1 = dst->src[1];
7677
const float * src0_d = (const float *)src0->data;
7778
const float * src1_d = (const float *)src1->data;
7879

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,6 @@
1515

1616
#include "common.hpp"
1717

18-
void ggml_sycl_op_conv_transpose_1d(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
19-
const ggml_tensor *src1, ggml_tensor *dst);
18+
void ggml_sycl_op_conv_transpose_1d(ggml_backend_sycl_context & ctx, ggml_tensor *dst);
2019

2120
#endif // GGML_SYCL_CONV_HPP

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

Lines changed: 48 additions & 48 deletions
Original file line numberDiff line numberDiff line change
@@ -882,149 +882,149 @@ inline void ggml_sycl_op_div(ggml_backend_sycl_context & ctx, const ggml_tensor
882882
}
883883

884884

885-
void ggml_sycl_sqrt(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
885+
void ggml_sycl_sqrt(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
886886
GGML_SYCL_DEBUG("call %s\n", __func__);
887-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_sqrt);
887+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_sqrt);
888888
GGML_SYCL_DEBUG("call %s done\n", __func__);
889889
}
890890

891-
void ggml_sycl_sin(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
891+
void ggml_sycl_sin(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
892892
GGML_SYCL_DEBUG("call %s\n", __func__);
893-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_sin);
893+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_sin);
894894
GGML_SYCL_DEBUG("call %s done\n", __func__);
895895
}
896896

897-
void ggml_sycl_cos(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
897+
void ggml_sycl_cos(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
898898
GGML_SYCL_DEBUG("call %s\n", __func__);
899-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_cos);
899+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_cos);
900900
GGML_SYCL_DEBUG("call %s done\n", __func__);
901901
}
902902

903-
void ggml_sycl_acc(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
903+
void ggml_sycl_acc(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
904904
GGML_SYCL_DEBUG("call %s\n", __func__);
905-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_acc);
905+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_acc);
906906
GGML_SYCL_DEBUG("call %s done\n", __func__);
907907
}
908908

909-
void ggml_sycl_gelu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
909+
void ggml_sycl_gelu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
910910
GGML_SYCL_DEBUG("call %s\n", __func__);
911-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_gelu);
911+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_gelu);
912912
GGML_SYCL_DEBUG("call %s done\n", __func__);
913913
}
914914

915-
void ggml_sycl_silu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
915+
void ggml_sycl_silu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
916916
GGML_SYCL_DEBUG("call %s\n", __func__);
917-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_silu);
917+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_silu);
918918
GGML_SYCL_DEBUG("call %s done\n", __func__);
919919
}
920920

921-
void ggml_sycl_gelu_quick(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
921+
void ggml_sycl_gelu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
922922
GGML_SYCL_DEBUG("call %s\n", __func__);
923-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_gelu_quick);
923+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_gelu_quick);
924924
GGML_SYCL_DEBUG("call %s done\n", __func__);
925925
}
926926

927-
void ggml_sycl_tanh(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
927+
void ggml_sycl_tanh(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
928928
GGML_SYCL_DEBUG("call %s\n", __func__);
929-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_tanh);
929+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_tanh);
930930
GGML_SYCL_DEBUG("call %s done\n", __func__);
931931
}
932932

933-
void ggml_sycl_relu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
933+
void ggml_sycl_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
934934
GGML_SYCL_DEBUG("call %s\n", __func__);
935-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_relu);
935+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_relu);
936936
GGML_SYCL_DEBUG("call %s done\n", __func__);
937937
}
938938

939-
void ggml_sycl_sigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
939+
void ggml_sycl_sigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
940940
GGML_SYCL_DEBUG("call %s\n", __func__);
941-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_sigmoid);
941+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_sigmoid);
942942
GGML_SYCL_DEBUG("call %s done\n", __func__);
943943
}
944944

945-
void ggml_sycl_hardsigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
945+
void ggml_sycl_hardsigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
946946
GGML_SYCL_DEBUG("call %s\n", __func__);
947-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_hardsigmoid);
947+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_hardsigmoid);
948948
GGML_SYCL_DEBUG("call %s done\n", __func__);
949949
}
950950

951-
void ggml_sycl_hardswish(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
951+
void ggml_sycl_hardswish(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
952952
GGML_SYCL_DEBUG("call %s\n", __func__);
953-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_hardswish);
953+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_hardswish);
954954
GGML_SYCL_DEBUG("call %s done\n", __func__);
955955
}
956956

957957

958-
void ggml_sycl_exp(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
958+
void ggml_sycl_exp(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
959959
GGML_SYCL_DEBUG("call %s\n", __func__);
960-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_exp);
960+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_exp);
961961
GGML_SYCL_DEBUG("call %s done\n", __func__);
962962
}
963963

964-
void ggml_sycl_log(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
964+
void ggml_sycl_log(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
965965
GGML_SYCL_DEBUG("call %s\n", __func__);
966-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_log);
966+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_log);
967967
GGML_SYCL_DEBUG("call %s done\n", __func__);
968968
}
969969

970-
void ggml_sycl_neg(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
970+
void ggml_sycl_neg(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
971971
GGML_SYCL_DEBUG("call %s\n", __func__);
972-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_neg);
972+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_neg);
973973
GGML_SYCL_DEBUG("call %s done\n", __func__);
974974
}
975975

976-
void ggml_sycl_step(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
976+
void ggml_sycl_step(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
977977
GGML_SYCL_DEBUG("call %s\n", __func__);
978-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_step);
978+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_step);
979979
GGML_SYCL_DEBUG("call %s done\n", __func__);
980980
}
981981

982-
void ggml_sycl_leaky_relu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
982+
void ggml_sycl_leaky_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
983983
GGML_SYCL_DEBUG("call %s\n", __func__);
984-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_leaky_relu);
984+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_leaky_relu);
985985
GGML_SYCL_DEBUG("call %s done\n", __func__);
986986
}
987987

988-
void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
988+
void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
989989
GGML_SYCL_DEBUG("call %s\n", __func__);
990-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_sqr);
990+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_sqr);
991991
GGML_SYCL_DEBUG("call %s done\n", __func__);
992992
}
993993

994-
void ggml_sycl_upscale(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
994+
void ggml_sycl_upscale(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
995995
GGML_SYCL_DEBUG("call %s\n", __func__);
996-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_upscale);
996+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_upscale);
997997
GGML_SYCL_DEBUG("call %s done\n", __func__);
998998
}
999999

1000-
void ggml_sycl_pad(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
1000+
void ggml_sycl_pad(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
10011001
GGML_SYCL_DEBUG("call %s\n", __func__);
1002-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_pad);
1002+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_pad);
10031003
GGML_SYCL_DEBUG("call %s done\n", __func__);
10041004
}
10051005

10061006

10071007

1008-
void ggml_sycl_add(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
1008+
void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
10091009
GGML_SYCL_DEBUG("call %s\n", __func__);
1010-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_add);
1010+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_add);
10111011
GGML_SYCL_DEBUG("call %s done\n", __func__);
10121012
}
10131013

1014-
void ggml_sycl_sub(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
1014+
void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
10151015
GGML_SYCL_DEBUG("call %s\n", __func__);
1016-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_sub);
1016+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_sub);
10171017
GGML_SYCL_DEBUG("call %s done\n", __func__);
10181018
}
10191019

1020-
void ggml_sycl_mul(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
1020+
void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
10211021
GGML_SYCL_DEBUG("call %s\n", __func__);
1022-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_mul);
1022+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_mul);
10231023
GGML_SYCL_DEBUG("call %s done\n", __func__);
10241024
}
10251025

1026-
void ggml_sycl_div(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
1026+
void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
10271027
GGML_SYCL_DEBUG("call %s\n", __func__);
1028-
ggml_sycl_op_flatten(ctx, src0, src1, dst, ggml_sycl_op_div);
1028+
ggml_sycl_op_flatten(ctx, dst->src[0], dst->src[1], dst, ggml_sycl_op_div);
10291029
GGML_SYCL_DEBUG("call %s done\n", __func__);
10301030
}

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

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -25,52 +25,52 @@ static __dpct_inline__ float op_div(const float a, const float b) {
2525
}
2626

2727

28-
void ggml_sycl_sqrt(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
28+
void ggml_sycl_sqrt(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
2929

30-
void ggml_sycl_sin(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
30+
void ggml_sycl_sin(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
3131

32-
void ggml_sycl_cos(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
32+
void ggml_sycl_cos(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
3333

34-
void ggml_sycl_acc(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
34+
void ggml_sycl_acc(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
3535

36-
void ggml_sycl_gelu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
36+
void ggml_sycl_gelu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
3737

38-
void ggml_sycl_silu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
38+
void ggml_sycl_silu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
3939

40-
void ggml_sycl_gelu_quick(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
40+
void ggml_sycl_gelu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
4141

42-
void ggml_sycl_tanh(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
42+
void ggml_sycl_tanh(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
4343

44-
void ggml_sycl_relu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
44+
void ggml_sycl_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
4545

46-
void ggml_sycl_sigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
46+
void ggml_sycl_sigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
4747

48-
void ggml_sycl_hardsigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
48+
void ggml_sycl_hardsigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
4949

50-
void ggml_sycl_hardswish(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
50+
void ggml_sycl_hardswish(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
5151

52-
void ggml_sycl_exp(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
52+
void ggml_sycl_exp(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
5353

54-
void ggml_sycl_log(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
54+
void ggml_sycl_log(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
5555

56-
void ggml_sycl_neg(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
56+
void ggml_sycl_neg(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
5757

58-
void ggml_sycl_step(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
58+
void ggml_sycl_step(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
5959

60-
void ggml_sycl_leaky_relu(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
60+
void ggml_sycl_leaky_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
6161

62-
void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
62+
void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
6363

64-
void ggml_sycl_upscale(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
64+
void ggml_sycl_upscale(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
6565

66-
void ggml_sycl_pad(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
66+
void ggml_sycl_pad(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
6767

68-
void ggml_sycl_add(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
68+
void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
6969

70-
void ggml_sycl_sub(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
70+
void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
7171

72-
void ggml_sycl_mul(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
72+
void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
7373

74-
void ggml_sycl_div(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
74+
void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
7575

7676
#endif // GGML_SYCL_ELEMENTWISE_HPP

0 commit comments

Comments
 (0)