[SYCL] fix unsupported UT cases of CONT & CPY (#25231 )

* fix unsupported UT cases of CONT & CPY * update ops.md * rm unused head file
[SYCL] support op col2im_1d (#25264 )
2026-07-08 13:35:56 +02:00 · 2026-07-07 12:20:52 +03:00 · 2026-07-07 11:07:46 +03:00 · 2026-07-07 10:48:50 +03:00 · 2026-07-07 10:43:41 +03:00 · 2026-07-07 10:40:38 +03:00
12 changed files with 1453 additions and 287 deletions
@@ -24,9 +24,9 @@ Legend:
 |                          ARGSORT | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                             CEIL | ❌ | ❌ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                            CLAMP | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | ✅ | 🟡 | ✅ | ❌ | ❌ |
-|                        COL2IM_1D | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
+|                        COL2IM_1D | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
 |                           CONCAT | ❌ | ✅ | ✅ | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
-|                             CONT | ❌ | 🟡 | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ❌ | ❌ |
+|                             CONT | ❌ | 🟡 | ✅ | ✅ | ✅ | 🟡 | ✅ | ✅ | 🟡 | ❌ | ❌ |
 |                          CONV_2D | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                       CONV_2D_DW | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                          CONV_3D | ❌ | ❌ | ✅ | ❌ | ✅ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
@@ -35,8 +35,8 @@ Legend:
 |                              COS | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | 🟡 | ✅ | ❌ | ❌ |
 |                      COUNT_EQUAL | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
 |                              CPY | ❌ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ❌ | ❌ |
-|               CROSS_ENTROPY_LOSS | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
-|          CROSS_ENTROPY_LOSS_BACK | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
+|               CROSS_ENTROPY_LOSS | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
+|          CROSS_ENTROPY_LOSS_BACK | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
 |                           CUMSUM | ❌ | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                             DIAG | ❌ | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
 |                    DIAG_MASK_INF | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |
@@ -407,6 +407,7 @@
 "SYCL0","GET_ROWS","type=i32,n=256,m=5,r=4,be1=7,be2=1,v=0","support","1","yes","SYCL"
 "SYCL0","GET_ROWS","type=i32,n=256,m=5,r=4,be1=7,be2=1,v=1","support","1","yes","SYCL"
 "SYCL0","GET_ROWS_BACK","type=f32,n=1,m=8,r=2,b=1,v=0","support","0","no","SYCL"
+"SYCL0","GET_ROWS_BACK","type=f32,n=1,m=70000,r=4,b=1,v=0","support","0","no","SYCL"
 "SYCL0","GET_ROWS_BACK","type=f32,n=256,m=5,r=4,b=1,v=0","support","0","no","SYCL"
 "SYCL0","GET_ROWS_BACK","type=f32,n=256,m=5,r=4,b=1,v=1","support","0","no","SYCL"
 "SYCL0","GET_ROWS_BACK","type=f16,n=256,m=5,r=4,b=1,v=0","support","0","no","SYCL"
@@ -5052,39 +5053,39 @@
 "SYCL0","CONV_TRANSPOSE_1D","ne_input=[3,2,1,1],ne_kernel=[3,2,2,1],s0=1,p0=0,d0=1","support","1","yes","SYCL"
 "SYCL0","CONV_TRANSPOSE_1D","ne_input=[3,2,1,1],ne_kernel=[3,1,2,1],s0=1,p0=0,d0=1","support","1","yes","SYCL"
 "SYCL0","CONV_TRANSPOSE_1D","ne_input=[2,1,1,1],ne_kernel=[3,1,1,1],s0=1,p0=0,d0=1","support","1","yes","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=16,OC=32,T_in=197,s0=8,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=7,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=1,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=6,OC=4,T_in=11,s0=3,p0=1","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=2,OC=3,T_in=9,s0=3,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=5,OC=4,T_in=11,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=8,OC=4,T_in=13,s0=4,p0=2","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=1,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=16,OC=1,T_in=197,s0=8,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=3,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f32,K=8,OC=2,T_in=3,s0=2,p0=5","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=16,OC=32,T_in=197,s0=8,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=7,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=1,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=6,OC=4,T_in=11,s0=3,p0=1","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=2,OC=3,T_in=9,s0=3,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=5,OC=4,T_in=11,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=8,OC=4,T_in=13,s0=4,p0=2","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=1,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=16,OC=1,T_in=197,s0=8,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=3,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=f16,K=8,OC=2,T_in=3,s0=2,p0=5","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=16,OC=32,T_in=197,s0=8,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=7,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=1,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=6,OC=4,T_in=11,s0=3,p0=1","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=2,OC=3,T_in=9,s0=3,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=5,OC=4,T_in=11,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=8,OC=4,T_in=13,s0=4,p0=2","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=1,s0=2,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=16,OC=1,T_in=197,s0=8,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=3,p0=0","support","0","no","SYCL"
-"SYCL0","COL2IM_1D","type=bf16,K=8,OC=2,T_in=3,s0=2,p0=5","support","0","no","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=16,OC=32,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=7,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=1,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=6,OC=4,T_in=11,s0=3,p0=1","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=2,OC=3,T_in=9,s0=3,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=5,OC=4,T_in=11,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=8,OC=4,T_in=13,s0=4,p0=2","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=1,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=16,OC=1,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=3,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f32,K=8,OC=2,T_in=3,s0=2,p0=5","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=16,OC=32,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=7,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=1,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=6,OC=4,T_in=11,s0=3,p0=1","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=2,OC=3,T_in=9,s0=3,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=5,OC=4,T_in=11,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=8,OC=4,T_in=13,s0=4,p0=2","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=1,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=16,OC=1,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=3,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=f16,K=8,OC=2,T_in=3,s0=2,p0=5","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=16,OC=32,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=7,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=1,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=6,OC=4,T_in=11,s0=3,p0=1","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=2,OC=3,T_in=9,s0=3,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=5,OC=4,T_in=11,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=8,OC=4,T_in=13,s0=4,p0=2","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=1,s0=2,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=16,OC=1,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=3,p0=0","support","1","yes","SYCL"
+"SYCL0","COL2IM_1D","type=bf16,K=8,OC=2,T_in=3,s0=2,p0=5","support","1","yes","SYCL"
 "SYCL0","CONV_TRANSPOSE_2D","kernel_type=f32,ne_input=[3,2,3,1],ne_kernel=[2,2,1,3],stride=1","support","1","yes","SYCL"
 "SYCL0","CONV_TRANSPOSE_2D","kernel_type=f32,ne_input=[10,10,9,1],ne_kernel=[3,3,1,9],stride=2","support","1","yes","SYCL"
 "SYCL0","CONV_TRANSPOSE_2D","kernel_type=f32,ne_input=[129,63,35,1],ne_kernel=[3,3,48,35],stride=1","support","1","yes","SYCL"
@@ -5166,15 +5167,15 @@
 "SYCL0","CPY","type_src=f16,type_dst=f16,ne_src=[3,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f16,ne_src=[3,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f16,ne_src=[3,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[1,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[1,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[1,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[2,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[2,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[2,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[3,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[3,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[3,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[1,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[1,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[1,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[2,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[2,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[2,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[3,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[3,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[3,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_0,type_dst=q4_0,ne_src=[32,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_0,type_dst=q4_0,ne_src=[32,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_0,type_dst=q4_0,ne_src=[32,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
@@ -5221,180 +5222,180 @@
 "SYCL0","CPY","type_src=q8_0,type_dst=q8_0,ne_src=[96,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q8_0,type_dst=q8_0,ne_src=[96,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[128,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[128,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[128,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[128,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[128,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[384,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[384,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[384,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[384,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q1_0,type_dst=q1_0,ne_src=[384,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[32,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[32,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[32,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[32,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[32,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[64,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[64,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[64,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[64,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[64,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[96,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[96,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[96,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[96,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=mxfp4,ne_src=[96,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[64,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[64,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[64,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[64,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[64,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[128,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[128,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[128,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[128,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[128,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[192,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[192,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[192,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[192,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=nvfp4,type_dst=nvfp4,ne_src=[192,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q2_K,type_dst=q2_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q3_K,type_dst=q3_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q4_K,type_dst=q4_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q5_K,type_dst=q5_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=q6_K,type_dst=q6_K,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xxs,type_dst=iq2_xxs,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_xs,type_dst=iq2_xs,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq2_s,type_dst=iq2_s,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_xxs,type_dst=iq3_xxs,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_s,type_dst=iq1_s,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq1_m,type_dst=iq1_m,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[32,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[32,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[32,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[32,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[32,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[64,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[64,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[64,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[64,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[64,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[96,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[96,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[96,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[96,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq4_nl,type_dst=iq4_nl,ne_src=[96,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq3_s,type_dst=iq3_s,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[256,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[512,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[512,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[512,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[768,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[768,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=iq4_xs,type_dst=iq4_xs,ne_src=[768,2,3,4],permute_src=[0,3,1,2],permute_dst=[0,2,1,3],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=bf16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=bf16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q4_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q4_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q4_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q4_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q5_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q5_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q5_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q5_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q8_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q8_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q1_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=mxfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=mxfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=nvfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=nvfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=bf16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=bf16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q4_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q4_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q4_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q4_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q5_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q5_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q5_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q5_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q8_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q8_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q1_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=mxfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=mxfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=nvfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=nvfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=q2_K,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=q3_K,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
@@ -5417,34 +5418,34 @@
 "SYCL0","CPY","type_src=f16,type_dst=iq1_s,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=iq1_m,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=iq1_m,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=iq4_nl,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f16,type_dst=iq4_nl,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=iq4_nl,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f16,type_dst=iq4_nl,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=iq3_s,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=iq3_s,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=iq4_xs,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=iq4_xs,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=f16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=f16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=f16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=f16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q4_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q4_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q4_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q4_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=q5_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=q5_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=q5_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=q5_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=q5_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=q5_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=q5_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=q5_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q8_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q8_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=q1_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=mxfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=mxfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=nvfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=nvfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=q1_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=mxfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=mxfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=nvfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=nvfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q2_K,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=q3_K,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
@@ -5467,8 +5468,8 @@
 "SYCL0","CPY","type_src=bf16,type_dst=iq1_s,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=iq1_m,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=iq1_m,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=iq4_nl,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=iq4_nl,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=iq4_nl,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=iq4_nl,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=iq3_s,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=iq3_s,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=bf16,type_dst=iq4_xs,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
@@ -5477,8 +5478,8 @@
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=bf16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=bf16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=bf16,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=bf16,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q4_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q4_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q4_1,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
@@ -5489,12 +5490,12 @@
 "SYCL0","CPY","type_src=f32,type_dst=q5_1,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q8_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q8_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=q1_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=mxfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=mxfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=nvfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=nvfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=q1_0,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=q1_0,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=mxfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=mxfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=nvfp4,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=nvfp4,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q2_K,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q2_K,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=q3_K,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
@@ -5527,8 +5528,8 @@
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_0,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_0,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q4_1,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
@@ -5541,8 +5542,8 @@
 "SYCL0","CPY","type_src=q8_0,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=q1_0,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=q1_0,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=mxfp4,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=mxfp4,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=nvfp4,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=nvfp4,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,2,1,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
 "SYCL0","CPY","type_src=q2_K,type_dst=f32,ne_src=[256,4,4,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
@@ -5577,19 +5578,19 @@
 "SYCL0","CPY","type_src=f16,type_dst=f32,ne_src=[256,2,3,4],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f16,ne_src=[256,2,3,4],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[256,2,3,4],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=i32,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=f32,type_dst=i32,ne_src=[256,2,3,4],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=i32,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
-"SYCL0","CPY","type_src=i32,type_dst=f32,ne_src=[256,2,3,4],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=i32,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=f32,type_dst=i32,ne_src=[256,2,3,4],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=i32,type_dst=f32,ne_src=[256,2,3,4],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=i32,type_dst=f32,ne_src=[256,2,3,4],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f16,ne_src=[256,4,3,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[256,4,3,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[256,4,3,3],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,4,3,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,4,3,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f16,type_dst=f16,ne_src=[256,4,1,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[256,4,1,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
-"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,4,1,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","0","no","SYCL"
-"SYCL0","CPY","type_src=i32,type_dst=i32,ne_src=[256,4,1,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","0","no","SYCL"
-"SYCL0","CPY","type_src=i32,type_dst=i32,ne_src=[256,1,4,1],permute_src=[1,2,0,3],permute_dst=[0,0,0,0],_src_transpose=0","support","0","no","SYCL"
+"SYCL0","CPY","type_src=bf16,type_dst=bf16,ne_src=[256,4,1,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=i32,type_dst=i32,ne_src=[256,4,1,1],permute_src=[0,0,0,0],permute_dst=[0,0,0,0],_src_transpose=1","support","1","yes","SYCL"
+"SYCL0","CPY","type_src=i32,type_dst=i32,ne_src=[256,1,4,1],permute_src=[1,2,0,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[256,1,4,1],permute_src=[1,2,0,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[2,2097121,1,1],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
 "SYCL0","CPY","type_src=f32,type_dst=f32,ne_src=[2,2,524281,1],permute_src=[1,0,2,3],permute_dst=[0,0,0,0],_src_transpose=0","support","1","yes","SYCL"
@@ -5723,12 +5724,12 @@
 "SYCL0","CONT","type=f16,ne=[1,4,4,1],use_view_slice=0","support","1","yes","SYCL"
 "SYCL0","CONT","type=f16,ne=[1,8,17,1],use_view_slice=0","support","1","yes","SYCL"
 "SYCL0","CONT","type=f16,ne=[10,10,10,1],use_view_slice=0","support","1","yes","SYCL"
-"SYCL0","CONT","type=bf16,ne=[2,1,1,1],use_view_slice=0","support","0","no","SYCL"
-"SYCL0","CONT","type=bf16,ne=[2,1,3,5],use_view_slice=0","support","0","no","SYCL"
-"SYCL0","CONT","type=bf16,ne=[2,3,5,7],use_view_slice=0","support","0","no","SYCL"
-"SYCL0","CONT","type=bf16,ne=[1,4,4,1],use_view_slice=0","support","0","no","SYCL"
-"SYCL0","CONT","type=bf16,ne=[1,8,17,1],use_view_slice=0","support","0","no","SYCL"
-"SYCL0","CONT","type=bf16,ne=[10,10,10,1],use_view_slice=0","support","0","no","SYCL"
+"SYCL0","CONT","type=bf16,ne=[2,1,1,1],use_view_slice=0","support","1","yes","SYCL"
+"SYCL0","CONT","type=bf16,ne=[2,1,3,5],use_view_slice=0","support","1","yes","SYCL"
+"SYCL0","CONT","type=bf16,ne=[2,3,5,7],use_view_slice=0","support","1","yes","SYCL"
+"SYCL0","CONT","type=bf16,ne=[1,4,4,1],use_view_slice=0","support","1","yes","SYCL"
+"SYCL0","CONT","type=bf16,ne=[1,8,17,1],use_view_slice=0","support","1","yes","SYCL"
+"SYCL0","CONT","type=bf16,ne=[10,10,10,1],use_view_slice=0","support","1","yes","SYCL"
 "SYCL0","ADD","type=f16,ne=[1,1,8,1],nr=[1,1,1,1],nf=1,perm1=0,src_overlap=0","support","1","yes","SYCL"
 "SYCL0","SUB","type=f16,ne=[1,1,8,1],nr=[1,1,1,1],nf=1,perm1=0,src_overlap=0","support","1","yes","SYCL"
 "SYCL0","MUL","type=f16,ne=[1,1,8,1],nr=[1,1,1,1],nf=1,perm1=0,src_overlap=0","support","1","yes","SYCL"
@@ -16747,10 +16748,10 @@ zjy 2
 "SYCL0","FLASH_ATTN_EXT","hsk=128,hsv=64,nh=4,nr23=[1,1],kv=128,nb=2,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=q1_0,type_V=q4_0,permute=[0,1,2,3]","support","0","no","SYCL"
 "SYCL0","FLASH_ATTN_EXT","hsk=64,hsv=128,nh=4,nr23=[1,1],kv=128,nb=2,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=q4_0,type_V=q1_0,permute=[0,1,2,3]","support","0","no","SYCL"
 "SYCL0","FLASH_ATTN_EXT","hsk=128,hsv=64,nh=4,nr23=[1,1],kv=64,nb=2,mask=1,sinks=0,max_bias=0.000000,logit_softcap=0.000000,prec=f32,type_K=q1_0,type_V=f16,permute=[0,1,2,3]","support","0","no","SYCL"
-"SYCL0","CROSS_ENTROPY_LOSS","type=f32,ne=[10,5,4,3]","support","0","no","SYCL"
-"SYCL0","CROSS_ENTROPY_LOSS","type=f32,ne=[30000,1,1,1]","support","0","no","SYCL"
-"SYCL0","CROSS_ENTROPY_LOSS_BACK","type=f32,ne=[10,5,4,3]","support","0","no","SYCL"
-"SYCL0","CROSS_ENTROPY_LOSS_BACK","type=f32,ne=[30000,1,1,1]","support","0","no","SYCL"
+"SYCL0","CROSS_ENTROPY_LOSS","type=f32,ne=[10,5,4,3]","support","1","yes","SYCL"
+"SYCL0","CROSS_ENTROPY_LOSS","type=f32,ne=[30000,1,1,1]","support","1","yes","SYCL"
+"SYCL0","CROSS_ENTROPY_LOSS_BACK","type=f32,ne=[10,5,4,3]","support","1","yes","SYCL"
+"SYCL0","CROSS_ENTROPY_LOSS_BACK","type=f32,ne=[30000,1,1,1]","support","1","yes","SYCL"
 "SYCL0","OPT_STEP_ADAMW","type=f32,ne=[10,5,4,3]","support","0","no","SYCL"
 "SYCL0","OPT_STEP_SGD","type=f32,ne=[10,5,4,3]","support","0","no","SYCL"
 "SYCL0","GATED_DELTA_NET","type=f32,head_count=32,head_size=128,n_seq_tokens=1,n_seqs=1,v_repeat=1,permuted=0,kda=0,K=1","support","1","yes","SYCL"
@@ -14,6 +14,7 @@
 #define GGML_SYCL_BACKEND_HPP

 #include "binbcast.hpp"
+#include "col2im-1d.hpp"
 #include "common.hpp"
 #include "concat.hpp"
 #include "conv.hpp"
@@ -0,0 +1,102 @@
+#include "col2im-1d.hpp"
+
+template <typename T>
+static void col2im_1d_sycl(
+        const T * col,
+        T * dst,
+        const int T_in,
+        const sycl::uint3 T_out_fd,
+        const int K,
+        const int K_OC,
+        const int32_t s0,
+        const int32_t p0,
+        const int total,
+        dpct::queue_ptr stream) {
+
+    const uint32_t block_size = SYCL_COL2IM_1D_BLOCK_SIZE;
+    const uint32_t num_blocks = (uint32_t) ((total + block_size - 1) / block_size);
+
+    stream->parallel_for(
+        sycl::nd_range<3>(
+            sycl::range<3>(1, 1, num_blocks * block_size),
+            sycl::range<3>(1, 1, block_size)),
+        [=](sycl::nd_item<3> item_ct1) {
+            const int idx = (int) item_ct1.get_global_id(2);
+            if (idx >= total) {
+                return;
+            }
+
+            const sycl::uint2 qr = fast_div_modulo((uint32_t) idx, T_out_fd);
+            const int oc    = (int) qr.x();
+            const int t_out = (int) qr.y();
+            const int t_abs = t_out + p0;
+
+            int t_in_min = (t_abs - K + s0) / s0;
+            if (t_in_min < 0) {
+                t_in_min = 0;
+            }
+            int t_in_max = t_abs / s0;
+            if (t_in_max >= T_in) {
+                t_in_max = T_in - 1;
+            }
+
+            float sum = 0.0f;
+            for (int t_in = t_in_min; t_in <= t_in_max; ++t_in) {
+                const int k = t_abs - t_in * s0;
+                sum += static_cast<float>(col[(oc * K + k) + t_in * K_OC]);
+            }
+
+            dst[idx] = static_cast<T>(sum);
+        });
+}
+
+void ggml_sycl_op_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    const ggml_tensor * src0 = dst->src[0];
+
+    GGML_ASSERT(src0 != nullptr);
+    GGML_ASSERT(ggml_is_contiguous(src0));
+    GGML_ASSERT(src0->type == dst->type);
+
+    const int32_t s0 = ((const int32_t *) dst->op_params)[0];
+    const int32_t OC = ((const int32_t *) dst->op_params)[1];
+    const int32_t p0 = ((const int32_t *) dst->op_params)[2];
+
+    const int K_OC  = (int) src0->ne[0];
+    const int T_in  = (int) src0->ne[1];
+    const int K     = K_OC / OC;
+    const int T_out = (int) dst->ne[0];
+
+    GGML_ASSERT(OC > 0);
+    GGML_ASSERT(K_OC % OC == 0);
+
+    const sycl::uint3 T_out_fd = init_fastdiv_values((uint32_t) T_out);
+
+    const int total = T_out * OC;
+
+    dpct::queue_ptr stream = ctx.stream();
+
+    switch (src0->type) {
+        case GGML_TYPE_F32:
+            col2im_1d_sycl<float>(
+                (const float *) src0->data,
+                (float *) dst->data,
+                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
+            break;
+        case GGML_TYPE_F16:
+            col2im_1d_sycl<sycl::half>(
+                (const sycl::half *) src0->data,
+                (sycl::half *) dst->data,
+                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
+            break;
+#ifdef GGML_SYCL_HAS_BF16
+        case GGML_TYPE_BF16:
+            col2im_1d_sycl<sycl::ext::oneapi::bfloat16>(
+                (const sycl::ext::oneapi::bfloat16 *) src0->data,
+                (sycl::ext::oneapi::bfloat16 *) dst->data,
+                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
+            break;
+#endif
+        default:
+            GGML_ABORT("col2im_1d: unsupported type %d", src0->type);
+    }
+}
@@ -0,0 +1,8 @@
+#ifndef GGML_SYCL_COL2IM_1D_HPP
+#define GGML_SYCL_COL2IM_1D_HPP
+
+#include "common.hpp"
+
+void ggml_sycl_op_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+#endif // GGML_SYCL_COL2IM_1D_HPP
@@ -1,6 +1,7 @@
 #include "cpy.hpp"

 #include <float.h>
+#include <vector>

 #include "dequantize.hpp"
 #include "ggml-sycl/common.hpp"
@@ -50,6 +51,57 @@ static void cpy_1_i32_i32(const char * cxi, char * cdsti) {
    *dsti = *xi;
 }

+static void cpy_1_f32_i32(const char * cxi, char * cdsti) {
+    const float * xi     = (const float *) cxi;
+    int32_t *     dsti   = (int32_t *) cdsti;
+
+    *dsti = (int32_t) *xi;
+}
+
+static void cpy_1_i32_f32(const char * cxi, char * cdsti) {
+    const int32_t * xi   = (const int32_t *) cxi;
+    float *         dsti = (float *) cdsti;
+
+    *dsti = (float) *xi;
+}
+
+#ifdef GGML_SYCL_HAS_BF16
+static void cpy_1_f32_bf16(const char * cxi, char * cdsti) {
+    const float * xi   = (const float *) cxi;
+    sycl::ext::oneapi::bfloat16 *  dsti = (sycl::ext::oneapi::bfloat16 *) cdsti;
+
+    *dsti = sycl::ext::oneapi::bfloat16(*xi);
+}
+
+static void cpy_1_bf16_f32(const char * cxi, char * cdsti) {
+    const sycl::ext::oneapi::bfloat16 * xi   = (const sycl::ext::oneapi::bfloat16 *) cxi;
+    float *            dsti = (float *) cdsti;
+
+    *dsti = static_cast<float>(*xi);
+}
+
+static void cpy_1_bf16_bf16(const char * cxi, char * cdsti) {
+    const sycl::ext::oneapi::bfloat16 * xi   = (const sycl::ext::oneapi::bfloat16 *) cxi;
+    sycl::ext::oneapi::bfloat16 *       dsti = (sycl::ext::oneapi::bfloat16 *) cdsti;
+
+    *dsti = *xi;
+}
+
+static void cpy_1_f16_bf16(const char * cxi, char * cdsti) {
+    const sycl::half * xi   = (const sycl::half *) cxi;
+    sycl::ext::oneapi::bfloat16 *  dsti = (sycl::ext::oneapi::bfloat16 *) cdsti;
+
+    *dsti = sycl::ext::oneapi::bfloat16(static_cast<float>(*xi));
+}
+
+static void cpy_1_bf16_f16(const char * cxi, char * cdsti) {
+    const sycl::ext::oneapi::bfloat16 * xi   = (const sycl::ext::oneapi::bfloat16 *) cxi;
+    sycl::half *            dsti = (sycl::half *) cdsti;
+
+    *dsti = sycl::half(static_cast<float>(*xi));
+}
+#endif
+
 template <cpy_kernel_t cpy_1>
 static void cpy_f32_f16(const char * cx, char * cdst, const int ne, const int ne00, const int ne01, const int ne02,
                        const int nb00, const int nb01, const int nb02, const int nb03, const int ne10, const int ne11,
@@ -247,6 +299,38 @@ static void ggml_cpy_f32_f16_sycl(const char * cx, char * cdst, const int ne, co
    }
 }

+static void ggml_cpy_f32_i32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                  const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
+    {
+        stream->parallel_for(
+            sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
+                              sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
+            [=](sycl::nd_item<3> item_ct1) {
+                cpy_f32_f16<cpy_1_f32_i32>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                           nb10, nb11, nb12, nb13, item_ct1);
+            });
+    }
+}
+
+static void ggml_cpy_i32_f32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                  const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
+    {
+        stream->parallel_for(
+            sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
+                              sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
+            [=](sycl::nd_item<3> item_ct1) {
+                cpy_f32_f16<cpy_1_i32_f32>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                           nb10, nb11, nb12, nb13, item_ct1);
+            });
+    }
+}
+
 static void ggml_cpy_f32_q8_0_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
@@ -376,6 +460,19 @@ static void ggml_cpy_q5_1_f32_sycl(const char * cx, char * cdst, const int ne, c
        });
 }

+static void ggml_cpy_mxfp4_f32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                    const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                    const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                    const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ne;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_f32<cpy_blck_q_f32<dequantize_mxfp4, QK_MXFP4>, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
+                                                                             nb01, nb02, nb03, ne10, ne11, ne12,
+                                                                             nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
 static void ggml_cpy_f32_iq4_nl_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
                                     const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
                                     const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
@@ -389,6 +486,269 @@ static void ggml_cpy_f32_iq4_nl_sycl(const char * cx, char * cdst, const int ne,
        });
 }

+static void cpy_blck_f16_q4_0(const char * cxi, char * cdsti) {
+    const sycl::half * xi = (const sycl::half *) cxi;
+    float xf[QK4_0];
+
+    for (int j = 0; j < QK4_0; ++j) {
+        xf[j] = (float) xi[j];
+    }
+
+    cpy_blck_f32_q4_0((const char *) xf, cdsti);
+}
+
+static void cpy_blck_f16_q4_1(const char * cxi, char * cdsti) {
+    const sycl::half * xi = (const sycl::half *) cxi;
+    float xf[QK4_1];
+
+    for (int j = 0; j < QK4_1; ++j) {
+        xf[j] = (float) xi[j];
+    }
+
+    cpy_blck_f32_q4_1((const char *) xf, cdsti);
+}
+
+static void cpy_blck_f16_q5_0(const char * cxi, char * cdsti) {
+    const sycl::half * xi = (const sycl::half *) cxi;
+    float xf[QK5_0];
+
+    for (int j = 0; j < QK5_0; ++j) {
+        xf[j] = (float) xi[j];
+    }
+
+    cpy_blck_f32_q5_0((const char *) xf, cdsti);
+}
+
+static void ggml_cpy_f16_q4_0_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    GGML_ASSERT(ne % QK4_0 == 0);
+    const int num_blocks = ne / QK4_0;
+    stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
+                         [=](sycl::nd_item<3> item_ct1) {
+                             cpy_f32_q<cpy_blck_f16_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02,
+                                                                 nb00, nb01, nb02, nb03,
+                                                                 ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+                         });
+}
+
+static void ggml_cpy_f16_q4_1_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    GGML_ASSERT(ne % QK4_1 == 0);
+    const int num_blocks = ne / QK4_1;
+    stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
+                         [=](sycl::nd_item<3> item_ct1) {
+                             cpy_f32_q<cpy_blck_f16_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02,
+                                                                 nb00, nb01, nb02, nb03,
+                                                                 ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+                         });
+}
+
+static void ggml_cpy_f16_q5_0_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    GGML_ASSERT(ne % QK5_0 == 0);
+    const int num_blocks = ne / QK5_0;
+    stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
+                         [=](sycl::nd_item<3> item_ct1) {
+                             cpy_f32_q<cpy_blck_f16_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02,
+                                                                 nb00, nb01, nb02, nb03,
+                                                                 ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+                         });
+}
+
+static bool ggml_sycl_is_quantized_type(enum ggml_type type) {
+    switch (type) {
+        case GGML_TYPE_Q1_0:
+        case GGML_TYPE_Q4_0:
+        case GGML_TYPE_Q4_1:
+        case GGML_TYPE_Q5_0:
+        case GGML_TYPE_Q5_1:
+        case GGML_TYPE_Q8_0:
+        case GGML_TYPE_MXFP4:
+        case GGML_TYPE_NVFP4:
+        case GGML_TYPE_Q2_K:
+        case GGML_TYPE_Q3_K:
+        case GGML_TYPE_Q4_K:
+        case GGML_TYPE_Q5_K:
+        case GGML_TYPE_Q6_K:
+        case GGML_TYPE_IQ2_XXS:
+        case GGML_TYPE_IQ2_XS:
+        case GGML_TYPE_IQ2_S:
+        case GGML_TYPE_IQ3_XXS:
+        case GGML_TYPE_IQ3_S:
+        case GGML_TYPE_IQ1_S:
+        case GGML_TYPE_IQ1_M:
+        case GGML_TYPE_IQ4_NL:
+        case GGML_TYPE_IQ4_XS:
+            return true;
+        default:
+            return false;
+    }
+}
+
+static bool ggml_sycl_can_quantize_rows_sycl(enum ggml_type type) {
+    switch (type) {
+        case GGML_TYPE_Q1_0:
+        case GGML_TYPE_Q4_0:
+        case GGML_TYPE_Q4_1:
+        case GGML_TYPE_Q5_0:
+        case GGML_TYPE_Q5_1:
+        case GGML_TYPE_Q8_0:
+        case GGML_TYPE_MXFP4:
+        case GGML_TYPE_NVFP4:
+        case GGML_TYPE_Q2_K:
+        case GGML_TYPE_Q3_K:
+        case GGML_TYPE_Q4_K:
+        case GGML_TYPE_Q5_K:
+        case GGML_TYPE_Q6_K:
+        case GGML_TYPE_IQ4_NL:
+        case GGML_TYPE_IQ4_XS:
+            return true;
+        default:
+            return false;
+    }
+}
+
+template <typename SrcScalar>
+static inline float ggml_sycl_src_to_f32(const SrcScalar & x) {
+    return (float) x;
+}
+
+#ifdef GGML_SYCL_HAS_BF16
+template <>
+inline float ggml_sycl_src_to_f32<sycl::ext::oneapi::bfloat16>(const sycl::ext::oneapi::bfloat16 & x) {
+    return static_cast<float>(x);
+}
+
+template <>
+inline float ggml_sycl_src_to_f32<ggml_bf16_t>(const ggml_bf16_t & x) {
+    union {
+        uint32_t u32;
+        float f32;
+    } value;
+
+    value.u32 = (uint32_t) x.bits << 16;
+    return value.f32;
+}
+#endif
+
+template <typename SrcScalar, cpy_kernel_t quantize_block, int qk>
+static void ggml_sycl_quantize_rows_q(const char * cx, char * cdst, const int64_t ne,
+                                      const int64_t ne00, const int64_t ne01, const int64_t ne02,
+                                      const size_t nb00, const size_t nb01, const size_t nb02, const size_t nb03,
+                                      const int64_t ne10, const int64_t ne11, const int64_t ne12,
+                                      const size_t nb10, const size_t nb11, const size_t nb12, const size_t nb13,
+                                      queue_ptr stream) {
+    GGML_ASSERT(ne % qk == 0);
+    GGML_ASSERT(ne00 % qk == 0);
+
+    const int64_t total_blocks = ne / qk;
+    constexpr int block_size = 256;
+    const int64_t grid_size = ceil_div(total_blocks, (int64_t) block_size);
+
+    stream->parallel_for(sycl::nd_range<1>(grid_size * block_size, block_size), [=](sycl::nd_item<1> item_ct1) {
+        const int64_t block_idx = item_ct1.get_global_linear_id();
+        if (block_idx >= total_blocks) {
+            return;
+        }
+
+        const int64_t i = block_idx * qk;
+
+        const int64_t i03 = i / (ne00 * ne01 * ne02);
+        const int64_t i02 = (i - i03 * ne00 * ne01 * ne02) / (ne00 * ne01);
+        const int64_t i01 = (i - i03 * ne00 * ne01 * ne02 - i02 * ne01 * ne00) / ne00;
+        const int64_t i00 = i - i03 * ne00 * ne01 * ne02 - i02 * ne01 * ne00 - i01 * ne00;
+        const size_t x_offset = i00 * nb00 + i01 * nb01 + i02 * nb02 + i03 * nb03;
+
+        const int64_t i13 = i / (ne10 * ne11 * ne12);
+        const int64_t i12 = (i - i13 * ne10 * ne11 * ne12) / (ne10 * ne11);
+        const int64_t i11 = (i - i13 * ne10 * ne11 * ne12 - i12 * ne10 * ne11) / ne10;
+        const int64_t i10 = i - i13 * ne10 * ne11 * ne12 - i12 * ne10 * ne11 - i11 * ne10;
+        const size_t dst_offset = (i10 / qk) * nb10 + i11 * nb11 + i12 * nb12 + i13 * nb13;
+
+        float xf[qk];
+        if (nb00 == sizeof(SrcScalar)) {
+            const SrcScalar * src_row = (const SrcScalar *) (cx + x_offset);
+            for (int j = 0; j < qk; ++j) {
+                xf[j] = ggml_sycl_src_to_f32(src_row[j]);
+            }
+        } else {
+            for (int j = 0; j < qk; ++j) {
+                const SrcScalar * src_val = (const SrcScalar *) (cx + x_offset + j * nb00);
+                xf[j] = ggml_sycl_src_to_f32(*src_val);
+            }
+        }
+
+        quantize_block((const char *) xf, cdst + dst_offset);
+    });
+}
+
+template <typename SrcScalar>
+static void ggml_sycl_quantize_rows_sycl(const char * cx, char * cdst, const ggml_tensor * src0, const ggml_tensor * src1,
+                                         const int64_t ne, const int64_t ne00, const int64_t ne01, const int64_t ne02,
+                                         const size_t nb00, const size_t nb01, const size_t nb02, const size_t nb03,
+                                         const int64_t ne10, const int64_t ne11, const int64_t ne12, const size_t nb10,
+                                         const size_t nb11, const size_t nb12, const size_t nb13, queue_ptr stream) {
+    GGML_UNUSED(src0);
+    GGML_UNUSED(src1);
+
+    switch (src1->type) {
+        case GGML_TYPE_Q8_0:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q8_0, QK8_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
+                                                                            nb02, nb03, ne10, ne11, ne12, nb10, nb11,
+                                                                            nb12, nb13, stream);
+            break;
+        case GGML_TYPE_Q1_0:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q1_0, QK1_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
+                                                                            nb02, nb03, ne10, ne11, ne12, nb10, nb11,
+                                                                            nb12, nb13, stream);
+            break;
+        case GGML_TYPE_Q5_1:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q5_1, QK5_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
+                                                                            nb02, nb03, ne10, ne11, ne12, nb10, nb11,
+                                                                            nb12, nb13, stream);
+            break;
+        case GGML_TYPE_Q5_0:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
+                                                                            nb02, nb03, ne10, ne11, ne12, nb10, nb11,
+                                                                            nb12, nb13, stream);
+            break;
+        case GGML_TYPE_Q4_1:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
+                                                                            nb02, nb03, ne10, ne11, ne12, nb10, nb11,
+                                                                            nb12, nb13, stream);
+            break;
+        case GGML_TYPE_Q4_0:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01,
+                                                                            nb02, nb03, ne10, ne11, ne12, nb10, nb11,
+                                                                            nb12, nb13, stream);
+            break;
+        case GGML_TYPE_IQ4_NL:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_iq4_nl, QK4_NL>(cx, cdst, ne, ne00, ne01, ne02, nb00,
+                                                                               nb01, nb02, nb03, ne10, ne11, ne12,
+                                                                               nb10, nb11, nb12, nb13, stream);
+            break;
+        case GGML_TYPE_MXFP4:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_mxfp4, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
+                                                                                nb01, nb02, nb03, ne10, ne11, ne12,
+                                                                                nb10, nb11, nb12, nb13, stream);
+            break;
+        case GGML_TYPE_NVFP4:
+            ggml_sycl_quantize_rows_q<SrcScalar, cpy_blck_f32_nvfp4, QK_NVFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
+                                                                                nb01, nb02, nb03, ne10, ne11, ne12,
+                                                                                nb10, nb11, nb12, nb13, stream);
+            break;
+        default:
+            GGML_ABORT("unsupported quantized target type in sycl quantizer src1->type=%s\n",
+                       ggml_type_name(src1->type));
+    }
+}
+
 static void ggml_cpy_f16_f16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
@@ -509,8 +869,269 @@ static void ggml_cpy_q4_1_q4_1(const char * cx, char * cdst, const int ne, const
        });
 }

+static void ggml_cpy_q1_0_q1_0(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_q1_0, QK1_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_mxfp4_mxfp4(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_mxfp4, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_nvfp4_nvfp4(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_nvfp4, QK_NVFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_q2_K_q2_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_q2_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_q3_K_q3_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_q3_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_q4_K_q4_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_q4_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_q5_K_q5_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_q5_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_q6_K_q6_K(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_q6_K, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq2_xxs_iq2_xxs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq2_xxs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq2_xs_iq2_xs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq2_xs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq2_s_iq2_s(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq2_s, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq3_xxs_iq3_xxs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq3_xxs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq1_s_iq1_s(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq1_s, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq1_m_iq1_m(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq1_m, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq4_nl_iq4_nl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq4_nl, QK4_NL>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq3_s_iq3_s(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq3_s, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_iq4_xs_iq4_xs(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                   const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                   const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                   const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)), [=](sycl::nd_item<3> item_ct1) {
+            cpy_q_q<block_iq4_xs, QK_K>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+#ifdef GGML_SYCL_HAS_BF16
+static void ggml_cpy_f32_bf16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                  const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
+                          sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
+        [=](sycl::nd_item<3> item_ct1) {
+            cpy_f32_f16<cpy_1_f32_bf16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                           nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_bf16_f32_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                  const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
+                          sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
+        [=](sycl::nd_item<3> item_ct1) {
+            cpy_f32_f16<cpy_1_bf16_f32>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                           nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_bf16_bf16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                  const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
+                          sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
+        [=](sycl::nd_item<3> item_ct1) {
+            cpy_f32_f16<cpy_1_bf16_bf16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                            nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_f16_bf16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                  const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
+                          sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
+        [=](sycl::nd_item<3> item_ct1) {
+            cpy_f32_f16<cpy_1_f16_bf16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                           nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+
+static void ggml_cpy_bf16_f16_sycl(const char * cx, char * cdst, const int ne, const int ne00, const int ne01,
+                                  const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
+                                  const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
+                                  const int nb12, const int nb13, queue_ptr stream) {
+    const int num_blocks = (ne + SYCL_CPY_BLOCK_SIZE - 1) / SYCL_CPY_BLOCK_SIZE;
+    stream->parallel_for(
+        sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
+                          sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
+        [=](sycl::nd_item<3> item_ct1) {
+            cpy_f32_f16<cpy_1_bf16_f16>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                           nb10, nb11, nb12, nb13, item_ct1);
+        });
+}
+#endif
+
 void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1) try {
    // Unlike other operators ggml_sycl_cpy takes 2 distinct tensors instead of a dst ggml_tensor and rely on its src field
+    GGML_SYCL_DEBUG("ggml_sycl_cpy: src0->type=%s, src1->type=%s\n",
+        ggml_type_name(src0->type), ggml_type_name(src1->type));
    scope_op_debug_print scope_dbg_print(__func__, src1, /*num_src=*/0, debug_get_tensor_str("\tsrc0", src0));
    const int64_t ne = ggml_nelements(src0);
    GGML_ASSERT(ne == ggml_nelements(src1));
@@ -525,12 +1146,31 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
    if ((src0->type == src1->type) && (ggml_is_contiguous(src0) && ggml_is_contiguous(src1))) {
        GGML_SYCL_DEBUG("%s: memcpy path\n", __func__);
        main_stream->memcpy(src1_ddc, src0_ddc, ggml_nbytes(src0));
+    } else if (src0->type == GGML_TYPE_F32 && ggml_sycl_is_quantized_type(src1->type)) {
+        GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(src1->type));
+        ggml_sycl_quantize_rows_sycl<float>(src0_ddc, src1_ddc, src0, src1, ne, ne00, ne01, ne02, nb00, nb01,
+                                            nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_F16 && ggml_sycl_is_quantized_type(src1->type)) {
+        GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(src1->type));
+        ggml_sycl_quantize_rows_sycl<sycl::half>(src0_ddc, src1_ddc, src0, src1, ne, ne00, ne01, ne02, nb00,
+                                                 nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
+                                                 main_stream);
+#ifdef GGML_SYCL_HAS_BF16
+    } else if (src0->type == GGML_TYPE_BF16 && ggml_sycl_is_quantized_type(src1->type)) {
+        GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(src1->type));
+        ggml_sycl_quantize_rows_sycl<ggml_bf16_t>(src0_ddc, src1_ddc, src0, src1, ne, ne00, ne01, ne02,
+                                                  nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11,
+                                                  nb12, nb13, main_stream);
+#endif
    } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) {
        ggml_cpy_f32_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                              nb11, nb12, nb13, main_stream);
    } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F16) {
        ggml_cpy_f32_f16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                              nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_I32) {
+        ggml_cpy_f32_i32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
+                              nb11, nb12, nb13, main_stream);
    } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_Q8_0) {
        ggml_cpy_f32_q8_0_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                               nb11, nb12, nb13, main_stream);
@@ -546,12 +1186,24 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
    } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16) {
        ggml_cpy_f16_f16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                              nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_Q4_0) {
+        ggml_cpy_f16_q4_0_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02,
+                               nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_Q4_1) {
+        ggml_cpy_f16_q4_1_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02,
+                               nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_Q5_0) {
+        ggml_cpy_f16_q5_0_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02,
+                               nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
    } else if (src0->type == GGML_TYPE_I16 && src1->type == GGML_TYPE_I16) {
        ggml_cpy_i16_i16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                              nb11, nb12, nb13, main_stream);
    } else if (src0->type == GGML_TYPE_I32 && src1->type == GGML_TYPE_I32) {
        ggml_cpy_i32_i32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                              nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_I32 && src1->type == GGML_TYPE_F32) {
+        ggml_cpy_i32_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
+                              nb11, nb12, nb13, main_stream);
    } else if (src0->type == GGML_TYPE_Q4_0 && src1->type == GGML_TYPE_F32) {
        ggml_cpy_q4_0_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                               nb11, nb12, nb13, main_stream);
@@ -573,6 +1225,9 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
    } else if (src0->type == GGML_TYPE_Q5_1 && src1->type == GGML_TYPE_F32) {
        ggml_cpy_q5_1_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
                               nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_MXFP4 && src1->type == GGML_TYPE_F32) {
+        ggml_cpy_mxfp4_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
+                                nb10, nb11, nb12, nb13, main_stream);
    } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_IQ4_NL) {
        ggml_cpy_f32_iq4_nl_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12,
                                 nb10, nb11, nb12, nb13, main_stream);
@@ -586,6 +1241,57 @@ void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, co
        ggml_cpy_q4_0_q4_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
    } else if (src0->type == GGML_TYPE_Q4_1 && src1->type == GGML_TYPE_Q4_1) {
        ggml_cpy_q4_1_q4_1(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_Q1_0 && src1->type == GGML_TYPE_Q1_0) {
+        ggml_cpy_q1_0_q1_0(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_MXFP4 && src1->type == GGML_TYPE_MXFP4) {
+        ggml_cpy_mxfp4_mxfp4(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_NVFP4 && src1->type == GGML_TYPE_NVFP4) {
+        ggml_cpy_nvfp4_nvfp4(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_Q2_K && src1->type == GGML_TYPE_Q2_K) {
+        ggml_cpy_q2_K_q2_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_Q3_K && src1->type == GGML_TYPE_Q3_K) {
+        ggml_cpy_q3_K_q3_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_Q4_K && src1->type == GGML_TYPE_Q4_K) {
+        ggml_cpy_q4_K_q4_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_Q5_K && src1->type == GGML_TYPE_Q5_K) {
+        ggml_cpy_q5_K_q5_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_Q6_K && src1->type == GGML_TYPE_Q6_K) {
+        ggml_cpy_q6_K_q6_K(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ2_XXS && src1->type == GGML_TYPE_IQ2_XXS) {
+        ggml_cpy_iq2_xxs_iq2_xxs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ2_XS && src1->type == GGML_TYPE_IQ2_XS) {
+        ggml_cpy_iq2_xs_iq2_xs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ2_S && src1->type == GGML_TYPE_IQ2_S) {
+        ggml_cpy_iq2_s_iq2_s(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ3_XXS && src1->type == GGML_TYPE_IQ3_XXS) {
+        ggml_cpy_iq3_xxs_iq3_xxs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ1_S && src1->type == GGML_TYPE_IQ1_S) {
+        ggml_cpy_iq1_s_iq1_s(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ1_M && src1->type == GGML_TYPE_IQ1_M) {
+        ggml_cpy_iq1_m_iq1_m(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ4_NL && src1->type == GGML_TYPE_IQ4_NL) {
+        ggml_cpy_iq4_nl_iq4_nl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ3_S && src1->type == GGML_TYPE_IQ3_S) {
+        ggml_cpy_iq3_s_iq3_s(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_IQ4_XS && src1->type == GGML_TYPE_IQ4_XS) {
+        ggml_cpy_iq4_xs_iq4_xs(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+#ifdef GGML_SYCL_HAS_BF16
+    } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_BF16) {
+        ggml_cpy_f32_bf16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
+                              nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F32) {
+        ggml_cpy_bf16_f32_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
+                              nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_BF16) {
+        ggml_cpy_bf16_bf16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
+                               nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_BF16) {
+        ggml_cpy_f16_bf16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
+                              nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F16) {
+        ggml_cpy_bf16_f16_sycl(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10,
+                              nb11, nb12, nb13, main_stream);
+#endif
    } else {
        GGML_LOG_ERROR("%s: unsupported type combination (%s to %s)\n", __func__, ggml_type_name(src0->type),
                       ggml_type_name(src1->type));
@@ -0,0 +1,255 @@
+#include "cross_entropy_loss.hpp"
+
+#include <cstdint>
+#include <cmath>
+
+template <bool has_shared>
+static __dpct_inline__ void cross_entropy_loss_f32_kernel(
+        const float * __restrict__ logits,
+        const float * __restrict__ labels,
+        float * __restrict__ row_loss,
+        const int nclasses,
+        const int nrows,
+        float * __restrict__ smem,
+        const sycl::nd_item<3> & item) {
+
+    const int row = item.get_group(2);
+    const int tid = item.get_local_id(2);
+
+    logits += (int64_t) row * nclasses;
+    labels += (int64_t) row * nclasses;
+
+    float max_logit = -INFINITY;
+    for (int i = tid; i < nclasses; i += WARP_SIZE) {
+        const float v = logits[i];
+        max_logit = sycl::fmax(max_logit, v);
+        if (has_shared) {
+            smem[i] = v;
+        }
+    }
+    max_logit = warp_reduce_max<WARP_SIZE>(max_logit);
+
+    float sum_exp = 0.0f;
+    for (int i = tid; i < nclasses; i += WARP_SIZE) {
+        const float v = has_shared ? smem[i] : logits[i];
+        sum_exp += sycl::exp(v - max_logit);
+    }
+    sum_exp = warp_reduce_sum<WARP_SIZE>(sum_exp);
+    const float log_sum = sycl::log(sum_exp);
+
+    float loss = 0.0f;
+    for (int i = tid; i < nclasses; i += WARP_SIZE) {
+        const float v = has_shared ? smem[i] : logits[i];
+        loss += (v - max_logit - log_sum) * labels[i];
+    }
+    loss = -warp_reduce_sum<WARP_SIZE>(loss) / (float) nrows;
+
+    if (tid == 0) {
+        row_loss[row] = loss;
+    }
+}
+
+template <bool has_shared>
+static __dpct_inline__ void cross_entropy_loss_back_f32_kernel(
+        const float * __restrict__ grad,
+        const float * __restrict__ logits,
+        const float * __restrict__ labels,
+        float * __restrict__ dst,
+        const int nclasses,
+        const int nrows,
+        float * __restrict__ smem,
+        const sycl::nd_item<3> & item) {
+
+    const int row = item.get_group(2);
+    const int tid = item.get_local_id(2);
+
+    logits += (int64_t) row * nclasses;
+    labels += (int64_t) row * nclasses;
+    dst    += (int64_t) row * nclasses;
+
+    float max_logit = -INFINITY;
+    for (int i = tid; i < nclasses; i += WARP_SIZE) {
+        const float v = logits[i];
+        max_logit = sycl::fmax(max_logit, v);
+        if (has_shared) {
+            smem[i] = v;
+        }
+    }
+    max_logit = warp_reduce_max<WARP_SIZE>(max_logit);
+
+    float sum_exp = 0.0f;
+    for (int i = tid; i < nclasses; i += WARP_SIZE) {
+        const float v = sycl::exp((has_shared ? smem[i] : logits[i]) - max_logit);
+        sum_exp += v;
+        if (has_shared) {
+            smem[i] = v;
+        } else {
+            dst[i] = v;
+        }
+    }
+    sum_exp = warp_reduce_sum<WARP_SIZE>(sum_exp);
+    const float inv_sum = 1.0f / sum_exp;
+
+    const float d_by_nrows = grad[0] / (float) nrows;
+    for (int i = tid; i < nclasses; i += WARP_SIZE) {
+        const float sm_num = has_shared ? smem[i] : dst[i];
+        dst[i] = (sm_num * inv_sum - labels[i]) * d_by_nrows;
+    }
+}
+
+static void cross_entropy_reduce_rows(
+        ggml_backend_sycl_context & ctx,
+        const float * row_loss,
+        float * dst,
+        const int64_t nrows) {
+    if (nrows == 1) {
+        SYCL_CHECK(CHECK_TRY_ERROR(
+            ctx.stream()->memcpy(dst, row_loss, sizeof(float))));
+        return;
+    }
+
+    ggml_sycl_pool_alloc<float> tmp_alloc(ctx.pool(), nrows);
+    float * tmp = tmp_alloc.get();
+    SYCL_CHECK(CHECK_TRY_ERROR(
+        ctx.stream()->memcpy(tmp, row_loss, nrows * sizeof(float))));
+
+    int64_t cur = nrows;
+    while (cur > 1) {
+        const int64_t out = (cur + WARP_SIZE - 1) / WARP_SIZE;
+        const sycl::range<3> block(1, 1, WARP_SIZE);
+        const sycl::range<3> grid(1, 1, out);
+        ctx.stream()->parallel_for(
+            sycl::nd_range<3>(grid * block, block),
+            [=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                const int row = item.get_group(2);
+                const int tid = item.get_local_id(2);
+                const int64_t i = (int64_t) row * WARP_SIZE + tid;
+                float v = i < cur ? tmp[i] : 0.0f;
+                v = warp_reduce_sum<WARP_SIZE>(v);
+                if (tid == 0) {
+                    tmp[row] = v;
+                }
+            });
+        cur = out;
+    }
+
+    SYCL_CHECK(CHECK_TRY_ERROR(
+        ctx.stream()->memcpy(dst, tmp, sizeof(float))));
+}
+
+void ggml_sycl_cross_entropy_loss(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+
+    const ggml_tensor * src0 = dst->src[0];
+    const ggml_tensor * src1 = dst->src[1];
+
+    GGML_ASSERT(src0->type == GGML_TYPE_F32);
+    GGML_ASSERT(src1->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type  == GGML_TYPE_F32);
+    GGML_ASSERT(ggml_is_contiguous(src0));
+    GGML_ASSERT(ggml_is_contiguous(src1));
+    GGML_ASSERT(ggml_is_contiguous(dst));
+    GGML_ASSERT(ggml_are_same_shape(src0, src1));
+    GGML_ASSERT(ggml_is_scalar(dst));
+
+    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
+
+    const int64_t nclasses = src0->ne[0];
+    const int64_t nrows = ggml_nrows(src0);
+
+    const float * logits_d = (const float *) src0->data;
+    const float * labels_d = (const float *) src1->data;
+    float * dst_d = (float *) dst->data;
+
+    ggml_sycl_pool_alloc<float> row_loss_alloc(ctx.pool(), nrows);
+    float * row_loss = row_loss_alloc.get();
+
+    const sycl::range<3> block(1, 1, WARP_SIZE);
+    const sycl::range<3> grid(1, 1, nrows);
+    const size_t nbytes_shared = (size_t) nclasses * sizeof(float);
+    const size_t smpbo = ggml_sycl_info().devices[ctx.device].smpbo;
+
+    if (nbytes_shared <= smpbo) {
+        ctx.stream()->submit([&](sycl::handler & cgh) {
+            sycl::local_accessor<float, 1> smem(sycl::range<1>(nclasses), cgh);
+            cgh.parallel_for(
+                sycl::nd_range<3>(grid * block, block),
+                [=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                    cross_entropy_loss_f32_kernel<true>(
+                        logits_d, labels_d, row_loss,
+                        (int) nclasses, (int) nrows,
+                        get_pointer(smem), item);
+                });
+        });
+    } else {
+        ctx.stream()->parallel_for(
+            sycl::nd_range<3>(grid * block, block),
+            [=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                cross_entropy_loss_f32_kernel<false>(
+                    logits_d, labels_d, row_loss,
+                    (int) nclasses, (int) nrows,
+                    nullptr, item);
+            });
+    }
+
+    cross_entropy_reduce_rows(ctx, row_loss, dst_d, nrows);
+}
+
+void ggml_sycl_cross_entropy_loss_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/3);
+
+    const ggml_tensor * grad  = dst->src[0];
+    const ggml_tensor * src0f = dst->src[1];
+    const ggml_tensor * src1f = dst->src[2];
+
+    GGML_ASSERT(grad->type == GGML_TYPE_F32);
+    GGML_ASSERT(src0f->type == GGML_TYPE_F32);
+    GGML_ASSERT(src1f->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type  == GGML_TYPE_F32);
+
+    GGML_ASSERT(ggml_is_scalar(grad));
+    GGML_ASSERT(ggml_is_contiguous(grad));
+    GGML_ASSERT(ggml_is_contiguous(src0f));
+    GGML_ASSERT(ggml_is_contiguous(src1f));
+    GGML_ASSERT(ggml_is_contiguous(dst));
+    GGML_ASSERT(ggml_are_same_shape(src0f, src1f));
+    GGML_ASSERT(ggml_are_same_shape(src0f, dst));
+
+    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
+
+    const int64_t nclasses = src0f->ne[0];
+    const int64_t nrows = ggml_nrows(src0f);
+
+    const float * grad_d  = (const float *) grad->data;
+    const float * logits_d = (const float *) src0f->data;
+    const float * labels_d = (const float *) src1f->data;
+    float * dst_d = (float *) dst->data;
+
+    const sycl::range<3> block(1, 1, WARP_SIZE);
+    const sycl::range<3> grid(1, 1, nrows);
+    const size_t nbytes_shared = (size_t) nclasses * sizeof(float);
+    const size_t smpbo = ggml_sycl_info().devices[ctx.device].smpbo;
+
+    if (nbytes_shared <= smpbo) {
+        ctx.stream()->submit([&](sycl::handler & cgh) {
+            sycl::local_accessor<float, 1> smem(sycl::range<1>(nclasses), cgh);
+            cgh.parallel_for(
+                sycl::nd_range<3>(grid * block, block),
+                [=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                    cross_entropy_loss_back_f32_kernel<true>(
+                        grad_d, logits_d, labels_d, dst_d,
+                        (int) nclasses, (int) nrows,
+                        get_pointer(smem), item);
+                });
+        });
+    } else {
+        ctx.stream()->parallel_for(
+            sycl::nd_range<3>(grid * block, block),
+            [=](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+                cross_entropy_loss_back_f32_kernel<false>(
+                    grad_d, logits_d, labels_d, dst_d,
+                    (int) nclasses, (int) nrows,
+                    nullptr, item);
+            });
+    }
+}
@@ -0,0 +1,7 @@
+#pragma once
+
+#include "common.hpp"
+
+void ggml_sycl_cross_entropy_loss(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+void ggml_sycl_cross_entropy_loss_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
@@ -680,14 +680,14 @@ static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx,
            q16[2] = q2[0] & 0x0f0f;
            q16[3] = q2[0] & 0xf0f0;

-            float4 s = {0.f, 0.f, 0.f, 0.f};
+            sycl::float4 s = {0.f, 0.f, 0.f, 0.f};
            float smin = 0;
            for (int l = 0; l < 2; ++l) {
-                s.x += y1[l] * q4[l+0]; s.y += y1[l+32] * q4[l+2];
-                s.z += y2[l] * q4[l+4]; s.w += y2[l+32] * q4[l+6];
+                s.x() += y1[l] * q4[l+0]; s.y() += y1[l+32] * q4[l+2];
+                s.z() += y2[l] * q4[l+4]; s.w() += y2[l+32] * q4[l+6];
                smin += y1[l] * sc[2] + y1[l+32] * sc[3] + y2[l] * sc[6] + y2[l+32] * sc[7];
            }
-            tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin;
+            tmp += dall * (s.x() * sc[0] + s.y() * sc[1] * 1.f/16.f + s.z() * sc[4] + s.w() * sc[5] * 1.f/16.f) - dmin * smin;
 #endif
        }

@@ -835,14 +835,14 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx,
            q16[2] = q2[0] & 0x0f0f;
            q16[3] = q2[0] & 0xf0f0;

-            float4 s = {0.f, 0.f, 0.f, 0.f};
+            sycl::float4 s = {0.f, 0.f, 0.f, 0.f};
            float smin = 0;
            for (int l = 0; l < 2; ++l) {
-                s.x += y1[l] * q4[l+0]; s.y += y1[l+32] * q4[l+2];
-                s.z += y2[l] * q4[l+4]; s.w += y2[l+32] * q4[l+6];
+                s.x() += y1[l] * q4[l+0]; s.y() += y1[l+32] * q4[l+2];
+                s.z() += y2[l] * q4[l+4]; s.w() += y2[l+32] * q4[l+6];
                smin += y1[l] * sc[2] + y1[l+32] * sc[3] + y2[l] * sc[6] + y2[l+32] * sc[7];
            }
-            tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin;
+            tmp += dall * (s.x() * sc[0] + s.y() * sc[1] * 1.f/16.f + s.z() * sc[4] + s.w() * sc[5] * 1.f/16.f) - dmin * smin;
 #endif
        }

@@ -1126,7 +1126,7 @@ static void dequantize_mul_mat_vec_q5_k_reorder(const void *__restrict__ vx,

    // sum up partial sums and write back result
 #pragma unroll
-    for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) {
+    for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
        tmp +=
            dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
    }
@@ -1762,10 +1762,13 @@ static void dequantize_mul_mat_vec_q5_K_sycl_reorder(const void *vx, const float
                                                     const int nrows,
                                                     dpct::queue_ptr stream) {
    GGML_ASSERT(ncols % QK_K == 0);
-    const sycl::range<3> block_dims(1, 1, QK_WARP_SIZE);
+    const int ny = 2 / K_QUANTS_PER_ITERATION;
+    const int block_num_y = (nrows + ny - 1) / ny;
+    const sycl::range<3> block_nums(1, 1, block_num_y);
+    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
    stream->parallel_for(
-        sycl::nd_range<3>(sycl::range<3>(1, 1, nrows) * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
+        sycl::nd_range<3>(block_nums * block_dims, block_dims),
+        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
            dequantize_mul_mat_vec_q5_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
        });
 }
@@ -9,9 +9,12 @@
 #define SYCL_LOCAL_ID_CALC(ITEM, IDX) \
    (ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX))

-static void acc_f32(const float * x, const float * y, float * dst, const int64_t ne,
-        const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
-        const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
+static void acc_f32(const char * x, const char * y, float * dst, const int64_t ne,
+    const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
+    const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
+    const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
+    const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
+    const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
    auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
    const int64_t i = SYCL_LOCAL_ID_CALC(item_ct1, 2);

@@ -30,9 +33,18 @@ static void acc_f32(const float * x, const float * y, float * dst, const int64_t
    tmp -= i11 * s11;
    const int64_t i10 = tmp;

-    float val = x[i];
+    int64_t tmp_dst = i;
+    const int64_t i3 = tmp_dst / (ne2*ne1*ne0);
+    tmp_dst -= i3 * (ne2*ne1*ne0);
+    const int64_t i2 = tmp_dst / (ne1*ne0);
+    tmp_dst -= i2 * (ne1*ne0);
+    const int64_t i1 = tmp_dst / ne0;
+    tmp_dst -= i1 * ne0;
+    const int64_t i0 = tmp_dst;
+
+    float val = *(const float *) (x + i0*nb00 + i1*nb01 + i2*nb02 + i3*nb03);
    if (src1_idx >= 0 && i10 < ne10 && i11 < ne11 && i12 < ne12 && i13 < ne13) {
-        val += y[((i13*ne12 + i12) * ne11 + i11) * ne10 + i10];
+        val += *(const float *) (y + i10*nb10 + i11*nb11 + i12*nb12 + i13*nb13);
    }
    dst[i] = val;
 }
@@ -422,15 +434,24 @@ static void gated_op_fused_geglu_quick(const T * x, const T * g, T * dst, const
 }

 namespace ggml_sycl_detail {
-static void acc_f32_sycl(const float *x, const float *y, float *dst,
-                         const int64_t n_elements, const int64_t ne10, const int64_t ne11,
-                         const int64_t ne12, const int64_t ne13, const int64_t s1, const int64_t s2, const int64_t s3,
+static void acc_f32_sycl(const char *x, const char *y, float *dst,
+                         const int64_t n_elements,
+                         const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
+                         const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
+                         const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
+                         const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
+                         const int64_t s1, const int64_t s2, const int64_t s3,
                         const int64_t offset, queue_ptr stream) {
    const int num_blocks = (n_elements + SYCL_ACC_BLOCK_SIZE - 1) / SYCL_ACC_BLOCK_SIZE;
    stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE),
                                           sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE)),
                         [=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                             acc_f32(x, y, dst, n_elements, ne10, ne11, ne12, ne13, s1, s2, s3, offset);
+                            acc_f32(x, y, dst, n_elements,
+                                ne0, ne1, ne2, ne3,
+                                nb00, nb01, nb02, nb03,
+                                ne10, ne11, ne12, ne13,
+                                nb10, nb11, nb12, nb13,
+                                s1, s2, s3, offset);
                         });
 }

@@ -843,8 +864,8 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
    const ggml_tensor * src0 = dst->src[0];
    const ggml_tensor * src1 = dst->src[1];

-    const float * src0_d = (const float *) src0->data;
-    const float * src1_d = (const float *) src1->data;
+    const char  * src0_d = (const char  *) src0->data;
+    const char  * src1_d = (const char  *) src1->data;
    float       * dst_d  = (float       *)  dst->data;

    dpct::queue_ptr stream = ctx.stream();
@@ -853,17 +874,20 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
    GGML_ASSERT(src1->type == GGML_TYPE_F32);
    GGML_ASSERT( dst->type == GGML_TYPE_F32);

-    GGML_ASSERT(ggml_is_contiguous(src1));
    GGML_ASSERT(dst->nb[0] == ggml_element_size(dst));
    GGML_ASSERT(ggml_is_contiguously_allocated(dst));
+    GGML_ASSERT(ggml_are_same_shape(src0, dst));

-    const int64_t s1     = dst->op_params[0] / sizeof(float);
-    const int64_t s2     = dst->op_params[1] / sizeof(float);
-    const int64_t s3     = dst->op_params[2] / sizeof(float);
-    const int64_t offset = dst->op_params[3] / sizeof(float);
+    const int64_t s1     = (int64_t) ((const int32_t *) dst->op_params)[0] / (int64_t) sizeof(float);
+    const int64_t s2     = (int64_t) ((const int32_t *) dst->op_params)[1] / (int64_t) sizeof(float);
+    const int64_t s3     = (int64_t) ((const int32_t *) dst->op_params)[2] / (int64_t) sizeof(float);
+    const int64_t offset = (int64_t) ((const int32_t *) dst->op_params)[3] / (int64_t) sizeof(float);

    ggml_sycl_detail::acc_f32_sycl(src0_d, src1_d, dst_d, ggml_nelements(dst),
+        dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
+        src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
        src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
+        src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
        s1, s2, s3, offset, stream);
 }

@@ -74,6 +74,7 @@
 #include "ggml-sycl/solve_tri.hpp"
 #include "ggml-sycl/gated_delta_net.hpp"
 #include "ggml-sycl/pool.hpp"
+#include "ggml-sycl/cross_entropy_loss.hpp"

 #define MEM_SIZE_2M	0x00200000
 #define MEM_SIZE_1G	0x40000000
@@ -546,7 +547,9 @@ ggml_backend_sycl_buffer_init_tensor(ggml_backend_buffer_t buffer,
        switch (tensor->type) {
            case GGML_TYPE_Q4_0:
            case GGML_TYPE_Q8_0:
+            case GGML_TYPE_Q3_K:
            case GGML_TYPE_Q4_K:
+            case GGML_TYPE_Q5_K:
            case GGML_TYPE_Q6_K:{
                ggml_tensor_extra_gpu * extra = new ggml_tensor_extra_gpu{};
                tensor->extra                 = extra;
@@ -3687,6 +3690,10 @@ inline bool ggml_sycl_supports_reorder_dmmv(enum ggml_type type) {
        case GGML_TYPE_Q1_0:
        case GGML_TYPE_Q4_0:
        case GGML_TYPE_Q8_0:
+        case GGML_TYPE_Q3_K:
+        case GGML_TYPE_Q4_K:
+        case GGML_TYPE_Q5_K:
+        case GGML_TYPE_Q6_K:
            return true;
        default:
            return false;
@@ -4764,6 +4771,11 @@ static void ggml_sycl_im2col_3d(ggml_backend_sycl_context & ctx, ggml_tensor * d
    ggml_sycl_op_im2col_3d(ctx, dst);
 }

+static void ggml_sycl_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
+    ggml_sycl_op_col2im_1d(ctx, dst);
+}
+
 static void ggml_sycl_conv_3d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
    ggml_sycl_op_conv_3d(ctx, dst);
@@ -5072,6 +5084,12 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
        case GGML_OP_SOFT_MAX_BACK:
            ggml_sycl_op_soft_max_back(ctx, dst);
            break;
+        case GGML_OP_CROSS_ENTROPY_LOSS:
+            ggml_sycl_cross_entropy_loss(ctx, dst);
+            break;
+        case GGML_OP_CROSS_ENTROPY_LOSS_BACK:
+            ggml_sycl_cross_entropy_loss_back(ctx, dst);
+            break;
        case GGML_OP_ROPE:
            ggml_sycl_rope(ctx, dst);
            break;
@@ -5084,6 +5102,9 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
        case GGML_OP_IM2COL_3D:
            ggml_sycl_im2col_3d(ctx, dst);
            break;
+        case GGML_OP_COL2IM_1D:
+            ggml_sycl_col2im_1d(ctx, dst);
+            break;
        case GGML_OP_POOL_2D:
            ggml_sycl_pool2d(ctx, dst);
            break;
@@ -5633,7 +5654,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
                // TODO: This specific configuration can fail with oneDNN and needs more debugging
                if (!ggml_is_permuted(a) && ggml_is_permuted(b) && b->ne[2] > 1 && b->ne[3] > 1 &&
                    a->ne[0] > 128 && a->ne[2] == 1 && src0_type == GGML_TYPE_F16) {
-                        printf("zjy 2\n");
                    return false;
                }
                return true;
@@ -5701,70 +5721,99 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
            {
                ggml_type src0_type = op->src[0]->type;
                ggml_type src1_type = op->src[1]->type;
-                if (src0_type == src1_type && (ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1])) && src0_type != GGML_TYPE_BF16) {
-                    return true;
+
+                if (src0_type == GGML_TYPE_F16) {
+                    if (src1_type == GGML_TYPE_Q2_K ||
+                        src1_type == GGML_TYPE_Q3_K ||
+                        src1_type == GGML_TYPE_Q4_K ||
+                        src1_type == GGML_TYPE_Q5_K ||
+                        src1_type == GGML_TYPE_Q6_K ||
+                        src1_type == GGML_TYPE_IQ2_XXS ||
+                        src1_type == GGML_TYPE_IQ2_XS ||
+                        src1_type == GGML_TYPE_IQ2_S ||
+                        src1_type == GGML_TYPE_IQ3_XXS ||
+                        src1_type == GGML_TYPE_IQ1_S ||
+                        src1_type == GGML_TYPE_IQ1_M ||
+                        src1_type == GGML_TYPE_IQ3_S ||
+                        src1_type == GGML_TYPE_IQ4_XS) {
+                        return false;
+                    }
                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_F32) {
-                    return true;
+
+                if (src0_type == GGML_TYPE_BF16) {
+                    if (src1_type == GGML_TYPE_Q4_0 || //big error in ut
+                        src1_type == GGML_TYPE_Q4_1 || //big error in ut
+                        src1_type == GGML_TYPE_Q8_0 || //big error in ut
+                        src1_type == GGML_TYPE_Q2_K ||
+                        src1_type == GGML_TYPE_Q3_K ||
+                        src1_type == GGML_TYPE_Q4_K ||
+                        src1_type == GGML_TYPE_Q5_K ||
+                        src1_type == GGML_TYPE_Q6_K ||
+                        src1_type == GGML_TYPE_IQ2_XXS ||
+                        src1_type == GGML_TYPE_IQ2_XS ||
+                        src1_type == GGML_TYPE_IQ2_S ||
+                        src1_type == GGML_TYPE_IQ3_XXS ||
+                        src1_type == GGML_TYPE_IQ1_S ||
+                        src1_type == GGML_TYPE_IQ1_M ||
+                        src1_type == GGML_TYPE_IQ3_S ||
+                        src1_type == GGML_TYPE_IQ4_XS) {
+                        return false;
+                    }
                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_F16) {
-                    return true;
+
+                if (src0_type == GGML_TYPE_F32) {
+                    if (src1_type == GGML_TYPE_Q2_K ||
+                        src1_type == GGML_TYPE_Q3_K ||
+                        src1_type == GGML_TYPE_Q4_K ||
+                        src1_type == GGML_TYPE_Q5_K ||
+                        src1_type == GGML_TYPE_Q6_K ||
+                        src1_type == GGML_TYPE_IQ2_XXS ||
+                        src1_type == GGML_TYPE_IQ2_XS ||
+                        src1_type == GGML_TYPE_IQ2_S ||
+                        src1_type == GGML_TYPE_IQ3_XXS ||
+                        src1_type == GGML_TYPE_IQ1_S ||
+                        src1_type == GGML_TYPE_IQ1_M ||
+                        src1_type == GGML_TYPE_IQ3_S ||
+                        src1_type == GGML_TYPE_IQ4_XS) {
+                        return false;
+                    }
                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q8_0) {
-                    return true;
+
+                if (src1_type == GGML_TYPE_F32) {
+                    if (src0_type == GGML_TYPE_Q1_0 ||
+                        src0_type == GGML_TYPE_NVFP4 ||
+                        src0_type == GGML_TYPE_Q2_K ||
+                        src0_type == GGML_TYPE_Q3_K ||
+                        src0_type == GGML_TYPE_Q4_K ||
+                        src0_type == GGML_TYPE_Q5_K ||
+                        src0_type == GGML_TYPE_Q6_K ||
+                        src0_type == GGML_TYPE_IQ2_XXS ||
+                        src0_type == GGML_TYPE_IQ2_XS ||
+                        src0_type == GGML_TYPE_IQ2_S ||
+                        src0_type == GGML_TYPE_IQ3_XXS ||
+                        src0_type == GGML_TYPE_IQ1_S ||
+                        src0_type == GGML_TYPE_IQ1_M ||
+                        src0_type == GGML_TYPE_IQ3_S ||
+                        src0_type == GGML_TYPE_IQ4_NL ||
+                        src0_type == GGML_TYPE_IQ4_XS
+                    ) {
+                        return false;
+                    }
                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q4_0) {
-                    return true;
+
+                if (src0_type == src1_type) {
+                    if (src1_type == GGML_TYPE_IQ2_XXS ||
+                        src1_type == GGML_TYPE_IQ2_XS ||
+                        src1_type == GGML_TYPE_IQ2_S ||
+                        src1_type == GGML_TYPE_IQ3_XXS ||
+                        src1_type == GGML_TYPE_IQ3_S ||
+                        src1_type == GGML_TYPE_IQ1_S ||
+                        src1_type == GGML_TYPE_IQ1_M) {
+                        return false;
+                    }
                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q4_1) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_F16 && src1_type == GGML_TYPE_F16) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_F16 && src1_type == GGML_TYPE_F32) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_Q8_0 && src1_type == GGML_TYPE_F32) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_Q4_0 && src1_type == GGML_TYPE_F32) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_Q4_1 && src1_type == GGML_TYPE_F32) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q5_0) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_Q5_0 && src1_type == GGML_TYPE_F32) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q5_1) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_Q5_1 && src1_type == GGML_TYPE_F32) {
-                    return true;
-                }
-                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_IQ4_NL) {
-                    return true;
-                }
-                if(src0_type == GGML_TYPE_Q8_0 && src1_type == GGML_TYPE_Q8_0) {
-                    return true;
-                }
-                if(src0_type == GGML_TYPE_Q5_0 && src1_type == GGML_TYPE_Q5_0) {
-                    return true;
-                }
-                if(src0_type == GGML_TYPE_Q5_1 && src1_type == GGML_TYPE_Q5_1) {
-                    return true;
-                }
-                if(src0_type == GGML_TYPE_Q4_0 && src1_type == GGML_TYPE_Q4_0) {
-                    return true;
-                }
-                if(src0_type == GGML_TYPE_Q4_1 && src1_type == GGML_TYPE_Q4_1) {
-                    return true;
-                }
-                return false;
+
+                return true;
            }
        case GGML_OP_REPEAT_BACK:
            {
@@ -5806,7 +5855,7 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
        case GGML_OP_SCALE:
            return true;
        case GGML_OP_CONT:
-            return op->src[0]->type != GGML_TYPE_BF16;
+            return true;
        case GGML_OP_TRI:
            {
                const ggml_tensor * src0 = op->src[0];
@@ -5829,6 +5878,14 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
        case GGML_OP_IM2COL_3D:
        case GGML_OP_UPSCALE:
            return true;
+        case GGML_OP_COL2IM_1D:
+            return ggml_is_contiguous(op->src[0]) &&
+                   (op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16
+#ifdef GGML_SYCL_HAS_BF16
+                    || op->type == GGML_TYPE_BF16
+#endif
+                   ) &&
+                   op->src[0]->type == op->type;
        case GGML_OP_CONV_3D:
            return op->type == GGML_TYPE_F32 &&
                   (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16) &&
@@ -5852,9 +5909,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
        }
        case GGML_OP_POOL_2D:
        case GGML_OP_POOL_1D:
-            return true;
        case GGML_OP_ACC:
-            return ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]);
+            return true;
        case GGML_OP_PAD:
            if (ggml_get_op_params_i32(op, 8) != 0) {
                return false;
@@ -5887,6 +5943,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
        case GGML_OP_FILL:
        case GGML_OP_CUMSUM:
        case GGML_OP_DIAG:
+        case GGML_OP_CROSS_ENTROPY_LOSS:
+        case GGML_OP_CROSS_ENTROPY_LOSS_BACK:
            return true;
        case GGML_OP_SOLVE_TRI:
            return op->src[0]->ne[0] <= SYCL_SOLVE_TRI_MAX_N && op->src[1]->ne[0] <= SYCL_SOLVE_TRI_MAX_K;
@@ -19,6 +19,7 @@
 #define WARP_SIZE GGML_SYCL_WARP_SIZE
 #define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses

+#define SYCL_COL2IM_1D_BLOCK_SIZE 256
 #define SYCL_GELU_BLOCK_SIZE 256
 #define SYCL_SILU_BLOCK_SIZE 256
 #define SYCL_TANH_BLOCK_SIZE 256
@@ -62,7 +63,7 @@
 #endif

 #ifndef K_QUANTS_PER_ITERATION
-#define K_QUANTS_PER_ITERATION 2
+#define K_QUANTS_PER_ITERATION 1
 #else
 static_assert(K_QUANTS_PER_ITERATION == 1 || K_QUANTS_PER_ITERATION == 2, "K_QUANTS_PER_ITERATION must be 1 or 2");
 #endif
Author	SHA1	Message	Date
Neo Zhang	108f186d17	[SYCL] fix unsupported UT cases of CONT & CPY (#25231 ) * fix unsupported UT cases of CONT & CPY * update ops.md * rm unused head file	2026-07-07 12:20:52 +03:00
Neo Zhang	47e1de77aa	[SYCL] support op col2im_1d (#25264 ) * support op col2im_1d * update ops.md * rm unused words * update for bf16 * optimize 1%-11% as the review comments * fix the format issue * update as the review comments	2026-07-07 11:07:46 +03:00
Neo Zhang	55edb2de44	[SYCL] support OP cross_entropy_loss, cross_entropy_loss_back (#25236 ) * support OP cross_entropy_loss, cross_entropy_loss_back * correct format issue	2026-07-07 10:48:50 +03:00
Todd Malsbary	d209086157	sycl : set K_QUANTS_PER_ITERATION to 1 on DMMV path (#25063 ) * sycl: add supported types to ggml_sycl_supports_reorder_dmmv The reordered feature is implemented in ggml_sycl_op_dequantize_mul_mat_vec, but gated by ggml_sycl_supports_reorder_dmmv. This commit fixes the gate. Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: set K_QUANTS_PER_ITERATION=1 to improve utilization When combined with opening the reorder gate, this improves GPU utilization on B70, giving a significant boost to tg t/s. Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: replace QK_WARP_SIZE with WARP_SIZE for QK_5 Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * sycl: add missing types to ggml_backend_sycl_buffer_init_tensor Without this, the extra field is not allocated and the reorder path will not take effect. Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> --------- Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>	2026-07-07 10:43:41 +03:00
Neo Zhang	95e5254c0a	[SYCL] fix unsupport ACC UT cases for noncontiguous (#25124 ) * fix unsupport ACC UT cases for noncontiguous * update ops.md	2026-07-07 10:40:38 +03:00