Skip to content

Commit 2330de7

Browse files
safranowithsafranowith
authored
SYCL: Add support for FLOOR,CEIL,ROUND and TRUNC unary operators (#16613)
* SYCL: Add support for FLOOR,CEIL,ROUND and TRUNC unary operators Clean up unrelated changes from previous commit * Chore: remove empty lines and fix indentation * Clean up: remove leftover blank lines and fix spacing * chore: fix trailing whitespace and ensure final newline * Cleanup: remove redundant declarations already defined in header * Sync docs/ops.md with updated backend operation support * docs: update ops.md after rebase * docs: update ops.md - Vulkan supports SSM_CONV and SSM_SCAN
1 parent 7062dd8 commit 2330de7

File tree

7 files changed

+313
-25
lines changed

7 files changed

+313
-25
lines changed

‎docs/ops.md‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ Legend:
2222
| ARANGE | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ | ❌ |
2323
| ARGMAX | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ |
2424
| ARGSORT | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ |
25-
| CEIL | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
25+
| CEIL | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ |
2626
| CLAMP | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ❌ |
2727
| CONCAT | ❌ | ✅ | ✅ | 🟡 | ✅ | 🟡 | 🟡 | ✅ | ❌ |
2828
| CONT | ❌ | 🟡 | ✅ | ✅ | ✅ | 🟡 | 🟡 | 🟡 | ❌ |
@@ -42,7 +42,7 @@ Legend:
4242
| ELU | ❌ | ✅ | ✅ | 🟡 | 🟡 | ❌ | 🟡 | ❌ | ❌ |
4343
| EXP | ❌ | ✅ | ✅ | 🟡 | 🟡 | ❌ | 🟡 | ❌ | ❌ |
4444
| FLASH_ATTN_EXT | ❌ | 🟡 | ✅ | 🟡 | 🟡 | ❌ | ❌ | 🟡 | ❌ |
45-
| FLOOR | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
45+
| FLOOR | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ |
4646
| GATED_LINEAR_ATTN | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ✅ | ❌ | ❌ |
4747
| GEGLU | ❌ | ✅ | ✅ | ✅ | 🟡 | ✅ | ✅ | 🟡 | ❌ |
4848
| GEGLU_ERF | ❌ | ✅ | ✅ | ✅ | 🟡 | ✅ | ✅ | 🟡 | ❌ |
@@ -84,7 +84,7 @@ Legend:
8484
| ROLL | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ |
8585
| ROPE | ❌ | 🟡 | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ |
8686
| ROPE_BACK | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ | ✅ | ❌ |
87-
| ROUND | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
87+
| ROUND | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ |
8888
| RWKV_WKV6 | ❌ | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ |
8989
| RWKV_WKV7 | ❌ | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ |
9090
| SCALE | ❌ | 🟡 | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ |
@@ -111,6 +111,6 @@ Legend:
111111
| TANH | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | 🟡 | ❌ |
112112
| TIMESTEP_EMBEDDING | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ |
113113
| TOPK_MOE | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ |
114-
| TRUNC | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
114+
| TRUNC | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ |
115115
| UPSCALE | ❌ | 🟡 | ✅ | ✅ | 🟡 | ✅ | 🟡 | ✅ | ❌ |
116116
| XIELU | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |

‎docs/ops/SYCL.csv‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,14 @@
3131
"SYCL0","GELU_ERF","type=f16,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
3232
"SYCL0","XIELU","type=f16,ne_a=[128,2,2,2],v=0","support","0","no","SYCL"
3333
"SYCL0","XIELU","type=f16,ne_a=[5,7,11,13],v=0","support","0","no","SYCL"
34+
"SYCL0","FLOOR","type=f16,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
35+
"SYCL0","FLOOR","type=f16,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
36+
"SYCL0","CEIL","type=f16,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
37+
"SYCL0","CEIL","type=f16,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
38+
"SYCL0","ROUND","type=f16,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
39+
"SYCL0","ROUND","type=f16,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
40+
"SYCL0","TRUNC","type=f16,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
41+
"SYCL0","TRUNC","type=f16,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
3442
"SYCL0","ABS","type=f16,ne_a=[128,2,2,2],v=1","support","0","no","SYCL"
3543
"SYCL0","ABS","type=f16,ne_a=[5,7,11,13],v=1","support","0","no","SYCL"
3644
"SYCL0","SGN","type=f16,ne_a=[128,2,2,2],v=1","support","0","no","SYCL"
@@ -95,6 +103,14 @@
95103
"SYCL0","GELU_ERF","type=f32,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
96104
"SYCL0","XIELU","type=f32,ne_a=[128,2,2,2],v=0","support","0","no","SYCL"
97105
"SYCL0","XIELU","type=f32,ne_a=[5,7,11,13],v=0","support","0","no","SYCL"
106+
"SYCL0","FLOOR","type=f32,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
107+
"SYCL0","FLOOR","type=f32,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
108+
"SYCL0","CEIL","type=f32,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
109+
"SYCL0","CEIL","type=f32,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
110+
"SYCL0","ROUND","type=f32,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
111+
"SYCL0","ROUND","type=f32,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
112+
"SYCL0","TRUNC","type=f32,ne_a=[128,2,2,2],v=0","support","1","yes","SYCL"
113+
"SYCL0","TRUNC","type=f32,ne_a=[5,7,11,13],v=0","support","1","yes","SYCL"
98114
"SYCL0","ABS","type=f32,ne_a=[128,2,2,2],v=1","support","0","no","SYCL"
99115
"SYCL0","ABS","type=f32,ne_a=[5,7,11,13],v=1","support","0","no","SYCL"
100116
"SYCL0","SGN","type=f32,ne_a=[128,2,2,2],v=1","support","0","no","SYCL"

‎docs/ops/Vulkan.csv‎

Lines changed: 21 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -3263,27 +3263,27 @@
32633263
"Vulkan0","RMS_NORM_MUL_ADD","type=f32,ne=[64,5,4,3],eps=1.000000,broadcast=0","support","1","yes","Vulkan"
32643264
"Vulkan0","RMS_NORM_MUL_ADD","type=f32,ne=[64,5,4,3],eps=1.000000,broadcast=1","support","1","yes","Vulkan"
32653265
"Vulkan0","L2_NORM","type=f32,ne=[64,5,4,3]","support","1","yes","Vulkan"
3266-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,1,1],ne_b=[3,1024,1,1]","support","0","no","Vulkan"
3267-
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1024,1,1],ne_b=[3,1024,1,1]","support","0","no","Vulkan"
3268-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,4,1],ne_b=[3,1024,1,1]","support","0","no","Vulkan"
3269-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,1,1],ne_b=[3,1536,1,1]","support","0","no","Vulkan"
3270-
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1536,1,1],ne_b=[3,1536,1,1]","support","0","no","Vulkan"
3271-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,4,1],ne_b=[3,1536,1,1]","support","0","no","Vulkan"
3272-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,1,1],ne_b=[3,2048,1,1]","support","0","no","Vulkan"
3273-
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,2048,1,1],ne_b=[3,2048,1,1]","support","0","no","Vulkan"
3274-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,4,1],ne_b=[3,2048,1,1]","support","0","no","Vulkan"
3275-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,1,1],ne_b=[4,1024,1,1]","support","0","no","Vulkan"
3276-
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1024,1,1],ne_b=[4,1024,1,1]","support","0","no","Vulkan"
3277-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,4,1],ne_b=[4,1024,1,1]","support","0","no","Vulkan"
3278-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,1,1],ne_b=[4,1536,1,1]","support","0","no","Vulkan"
3279-
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1536,1,1],ne_b=[4,1536,1,1]","support","0","no","Vulkan"
3280-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,4,1],ne_b=[4,1536,1,1]","support","0","no","Vulkan"
3281-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,1,1],ne_b=[4,2048,1,1]","support","0","no","Vulkan"
3282-
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,2048,1,1],ne_b=[4,2048,1,1]","support","0","no","Vulkan"
3283-
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,4,1],ne_b=[4,2048,1,1]","support","0","no","Vulkan"
3284-
"Vulkan0","SSM_SCAN","type=f32,d_state=16,head_dim=1,n_head=1024,n_group=1,n_seq_tokens=32,n_seqs=4","support","0","no","Vulkan"
3285-
"Vulkan0","SSM_SCAN","type=f32,d_state=128,head_dim=64,n_head=16,n_group=2,n_seq_tokens=32,n_seqs=4","support","0","no","Vulkan"
3286-
"Vulkan0","SSM_SCAN","type=f32,d_state=256,head_dim=64,n_head=8,n_group=2,n_seq_tokens=32,n_seqs=4","support","0","no","Vulkan"
3266+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,1,1],ne_b=[3,1024,1,1]","support","1","yes","Vulkan"
3267+
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1024,1,1],ne_b=[3,1024,1,1]","support","1","yes","Vulkan"
3268+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,4,1],ne_b=[3,1024,1,1]","support","1","yes","Vulkan"
3269+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,1,1],ne_b=[3,1536,1,1]","support","1","yes","Vulkan"
3270+
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1536,1,1],ne_b=[3,1536,1,1]","support","1","yes","Vulkan"
3271+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,4,1],ne_b=[3,1536,1,1]","support","1","yes","Vulkan"
3272+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,1,1],ne_b=[3,2048,1,1]","support","1","yes","Vulkan"
3273+
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,2048,1,1],ne_b=[3,2048,1,1]","support","1","yes","Vulkan"
3274+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,4,1],ne_b=[3,2048,1,1]","support","1","yes","Vulkan"
3275+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,1,1],ne_b=[4,1024,1,1]","support","1","yes","Vulkan"
3276+
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1024,1,1],ne_b=[4,1024,1,1]","support","1","yes","Vulkan"
3277+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1024,4,1],ne_b=[4,1024,1,1]","support","1","yes","Vulkan"
3278+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,1,1],ne_b=[4,1536,1,1]","support","1","yes","Vulkan"
3279+
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,1536,1,1],ne_b=[4,1536,1,1]","support","1","yes","Vulkan"
3280+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,1536,4,1],ne_b=[4,1536,1,1]","support","1","yes","Vulkan"
3281+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,1,1],ne_b=[4,2048,1,1]","support","1","yes","Vulkan"
3282+
"Vulkan0","SSM_CONV","type=f32,ne_a=[8,2048,1,1],ne_b=[4,2048,1,1]","support","1","yes","Vulkan"
3283+
"Vulkan0","SSM_CONV","type=f32,ne_a=[4,2048,4,1],ne_b=[4,2048,1,1]","support","1","yes","Vulkan"
3284+
"Vulkan0","SSM_SCAN","type=f32,d_state=16,head_dim=1,n_head=1024,n_group=1,n_seq_tokens=32,n_seqs=4","support","1","yes","Vulkan"
3285+
"Vulkan0","SSM_SCAN","type=f32,d_state=128,head_dim=64,n_head=16,n_group=2,n_seq_tokens=32,n_seqs=4","support","1","yes","Vulkan"
3286+
"Vulkan0","SSM_SCAN","type=f32,d_state=256,head_dim=64,n_head=8,n_group=2,n_seq_tokens=32,n_seqs=4","support","1","yes","Vulkan"
32873287
"Vulkan0","RWKV_WKV6","type=f32,head_count=32,head_size=64,n_seq_tokens=1,n_seqs=1","support","1","yes","Vulkan"
32883288
"Vulkan0","RWKV_WKV6","type=f32,head_count=32,head_size=64,n_seq_tokens=32,n_seqs=1","support","1","yes","Vulkan"
32893289
"Vulkan0","RWKV_WKV6","type=f32,head_count=32,head_size=64,n_seq_tokens=32,n_seqs=4","support","1","yes","Vulkan"

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

Lines changed: 120 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -150,6 +150,26 @@ static __dpct_inline__ T op_clamp(T x, float min_val, float max_val) {
150150
return x < static_cast<T>(min_val) ? static_cast<T>(min_val) : (x > static_cast<T>(max_val) ? static_cast<T>(max_val) : x);
151151
}
152152

153+
template<typename T>
154+
static __dpct_inline__ T op_floor(T x) {
155+
return sycl::floor(x);
156+
}
157+
158+
template<typename T>
159+
static __dpct_inline__ T op_ceil(T x) {
160+
return sycl::ceil(x);
161+
}
162+
163+
template<typename T>
164+
static __dpct_inline__ T op_round(T x) {
165+
return sycl::round(x);
166+
}
167+
168+
template<typename T>
169+
static __dpct_inline__ T op_trunc(T x) {
170+
return sycl::trunc(x);
171+
}
172+
153173
template<typename T>
154174
static void unary_op_sgn_kernel(const T * x, T * dst, const int k, const sycl::nd_item<1> &item_ct1) {
155175
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
@@ -304,6 +324,34 @@ static void unary_op_clamp_kernel(const T * x, T * dst, const int k, const sycl:
304324
}
305325
}
306326

327+
template<typename T>
328+
static void unary_op_floor_kernel(const T * x, T * dst, const int k, const sycl::nd_item<1> &item_ct1) {
329+
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
330+
dst[i] = op_floor(x[i]);
331+
}
332+
}
333+
334+
template<typename T>
335+
static void unary_op_ceil_kernel(const T * x, T * dst, const int k, const sycl::nd_item<1> &item_ct1) {
336+
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
337+
dst[i] = op_ceil(x[i]);
338+
}
339+
}
340+
341+
template<typename T>
342+
static void unary_op_round_kernel(const T * x, T * dst, const int k, const sycl::nd_item<1> &item_ct1) {
343+
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
344+
dst[i] = op_round(x[i]);
345+
}
346+
}
347+
348+
template<typename T>
349+
static void unary_op_trunc_kernel(const T * x, T * dst, const int k, const sycl::nd_item<1> &item_ct1) {
350+
SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
351+
dst[i] = op_trunc(x[i]);
352+
}
353+
}
354+
307355
template<typename T>
308356
static void upscale(const T *x, T *dst, const int nb00, const int nb01,
309357
const int nb02, const int nb03, const int ne10, const int ne11,
@@ -897,6 +945,58 @@ static inline void ggml_sycl_op_clamp(ggml_backend_sycl_context & ctx, ggml_tens
897945
}, min_val, max_val);
898946
}
899947

948+
static inline void ggml_sycl_op_floor(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
949+
ggml_sycl_detail::dispatch_ggml_sycl_op_unary(ctx, dst,
950+
[](const auto* src, auto* dst_ptr, int k_elements, queue_ptr stream) {
951+
const int num_blocks = ceil_div(k_elements, 256);
952+
stream->parallel_for(
953+
sycl::nd_range<1>(sycl::range<1>(num_blocks) * sycl::range<1>(256),
954+
sycl::range<1>(256)),
955+
[=](sycl::nd_item<1> item_ct1) {
956+
unary_op_floor_kernel(src, dst_ptr, k_elements, item_ct1);
957+
});
958+
});
959+
}
960+
961+
static inline void ggml_sycl_op_ceil(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
962+
ggml_sycl_detail::dispatch_ggml_sycl_op_unary(ctx, dst,
963+
[](const auto* src, auto* dst_ptr, int k_elements, queue_ptr stream) {
964+
const int num_blocks = ceil_div(k_elements, 256);
965+
stream->parallel_for(
966+
sycl::nd_range<1>(sycl::range<1>(num_blocks) * sycl::range<1>(256),
967+
sycl::range<1>(256)),
968+
[=](sycl::nd_item<1> item_ct1) {
969+
unary_op_ceil_kernel(src, dst_ptr, k_elements, item_ct1);
970+
});
971+
});
972+
}
973+
974+
static inline void ggml_sycl_op_round(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
975+
ggml_sycl_detail::dispatch_ggml_sycl_op_unary(ctx, dst,
976+
[](const auto* src, auto* dst_ptr, int k_elements, queue_ptr stream) {
977+
const int num_blocks = ceil_div(k_elements, 256);
978+
stream->parallel_for(
979+
sycl::nd_range<1>(sycl::range<1>(num_blocks) * sycl::range<1>(256),
980+
sycl::range<1>(256)),
981+
[=](sycl::nd_item<1> item_ct1) {
982+
unary_op_round_kernel(src, dst_ptr, k_elements, item_ct1);
983+
});
984+
});
985+
}
986+
987+
static inline void ggml_sycl_op_trunc(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
988+
ggml_sycl_detail::dispatch_ggml_sycl_op_unary(ctx, dst,
989+
[](const auto* src, auto* dst_ptr, int k_elements, queue_ptr stream) {
990+
const int num_blocks = ceil_div(k_elements, 256);
991+
stream->parallel_for(
992+
sycl::nd_range<1>(sycl::range<1>(num_blocks) * sycl::range<1>(256),
993+
sycl::range<1>(256)),
994+
[=](sycl::nd_item<1> item_ct1) {
995+
unary_op_trunc_kernel(src, dst_ptr, k_elements, item_ct1);
996+
});
997+
});
998+
}
999+
9001000
static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
9011001
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32);
9021002
GGML_ASSERT(dst->src[1]->type == GGML_TYPE_F32);
@@ -1122,3 +1222,23 @@ void ggml_sycl_arange(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
11221222
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/0);
11231223
ggml_sycl_detail::ggml_sycl_op_arange(ctx, dst);
11241224
}
1225+
1226+
void ggml_sycl_floor(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1227+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
1228+
ggml_sycl_op_floor(ctx, dst);
1229+
}
1230+
1231+
void ggml_sycl_ceil(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1232+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
1233+
ggml_sycl_op_ceil(ctx, dst);
1234+
}
1235+
1236+
void ggml_sycl_round(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1237+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
1238+
ggml_sycl_op_round(ctx, dst);
1239+
}
1240+
1241+
void ggml_sycl_trunc(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
1242+
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
1243+
ggml_sycl_op_trunc(ctx, dst);
1244+
}

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -80,6 +80,10 @@ void ggml_sycl_reglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
8080
void ggml_sycl_swiglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
8181
void ggml_sycl_geglu_erf(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
8282
void ggml_sycl_geglu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
83+
void ggml_sycl_floor(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
84+
void ggml_sycl_ceil(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
85+
void ggml_sycl_round(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
86+
void ggml_sycl_trunc(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
8387

8488
void ggml_sycl_arange(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
8589

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

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3698,6 +3698,18 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
36983698
case GGML_UNARY_OP_ELU:
36993699
ggml_sycl_elu(ctx, dst);
37003700
break;
3701+
case GGML_UNARY_OP_FLOOR:
3702+
ggml_sycl_floor(ctx, dst);
3703+
break;
3704+
case GGML_UNARY_OP_CEIL:
3705+
ggml_sycl_ceil(ctx, dst);
3706+
break;
3707+
case GGML_UNARY_OP_ROUND:
3708+
ggml_sycl_round(ctx, dst);
3709+
break;
3710+
case GGML_UNARY_OP_TRUNC:
3711+
ggml_sycl_trunc(ctx, dst);
3712+
break;
37013713
default:
37023714
return false;
37033715
}
@@ -4262,6 +4274,10 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
42624274
case GGML_UNARY_OP_SGN:
42634275
case GGML_UNARY_OP_ABS:
42644276
case GGML_UNARY_OP_ELU:
4277+
case GGML_UNARY_OP_FLOOR:
4278+
case GGML_UNARY_OP_CEIL:
4279+
case GGML_UNARY_OP_ROUND:
4280+
case GGML_UNARY_OP_TRUNC:
42654281
#if defined (GGML_SYCL_F16)
42664282
return ggml_is_contiguous(op->src[0]) && (op->type == op->src[0]->type);
42674283
#else

0 commit comments

Comments
 (0)