2
2
#include " ggml-sycl/presets.hpp"
3
3
#include " ggml.h"
4
4
#include " element_wise.hpp"
5
- #include < cstring>
6
-
5
+ #include < cstring>
7
6
#define SYCL_GLOBAL_ID_LOOP (K, ITEM ) \
8
7
for (auto i = ITEM.get_global_id(0 ); i < (size_t )K; i += ITEM.get_global_range(0 ))
9
8
@@ -939,8 +938,7 @@ static inline void ggml_sycl_op_set(ggml_backend_sycl_context & ctx, ggml_tensor
939
938
#else
940
939
GGML_ASSERT (dst->type == GGML_TYPE_F32 || dst->type == GGML_TYPE_I32);
941
940
#endif
942
- const size_t ts = ggml_type_size (dst->type );
943
-
941
+ const size_t ts = ggml_type_size (dst->type );
944
942
dpct::queue_ptr q = ctx.stream ();
945
943
{
946
944
const bool same_type = (src0->type == dst->type );
@@ -1003,9 +1001,8 @@ static inline void ggml_sycl_op_set(ggml_backend_sycl_context & ctx, ggml_tensor
1003
1001
void *p_dst = dst->data ;
1004
1002
1005
1003
const size_t sb0 = src1->nb [0 ], sb1 = src1->nb [1 ], sb2 = src1->nb [2 ], sb3 = src1->nb [3 ];
1006
- const size_t db0 = dst->nb [0 ];
1004
+ const size_t db0 = dst->nb [0 ];
1007
1005
const int64_t ne0 = src1->ne [0 ], ne1 = src1->ne [1 ], ne2 = src1->ne [2 ], ne3 = src1->ne [3 ];
1008
-
1009
1006
if (ggml_is_contiguous (src1) && db0 == ts) {
1010
1007
const size_t row_bytes = (size_t ) ne0 * ts;
1011
1008
const char *s_base = (const char *) p_src1;
@@ -1028,11 +1025,10 @@ static inline void ggml_sycl_op_set(ggml_backend_sycl_context & ctx, ggml_tensor
1028
1025
}
1029
1026
}
1030
1027
} else {
1031
-
1028
+
1032
1029
const size_t N = (size_t ) (ne0 * ne1 * ne2 * ne3);
1033
1030
const size_t WG = 256 ;
1034
1031
const size_t NG = ((N + WG - 1 ) / WG) * WG;
1035
-
1036
1032
const size_t ge0 = (size_t ) ne0;
1037
1033
const size_t ge1 = ge0 * (size_t ) ne1;
1038
1034
const size_t ge2 = ge1 * (size_t ) ne2;
0 commit comments