common : add missing <fstream> include in common.h (#25220)

Signed-off-by: zhangrunda <zhangrunda1234@outlook.com>

common/common.h

@@ -14,6 +14,7 @@
 #include <vector>
 #include <map>
 #include <algorithm>
+#include <fstream>
 
 #if defined(_WIN32) && !defined(_WIN32_WINNT)
 #define _WIN32_WINNT 0x0A00

docs/ops.md

@@ -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 | ❌ | ✅ | ✅ | ✅ | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ | ❌ |

docs/ops/SYCL.csv

@@ -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"

ggml/src/ggml-hip/CMakeLists.txt

@@ -156,4 +156,4 @@ endif()
 
 target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas)
 
-target_compile_options(ggml-hip PRIVATE "$<$<COMPILE_LANGUAGE:HIP>:-ffast-math>")
+target_compile_options(ggml-hip PRIVATE "$<$<COMPILE_LANGUAGE:HIP>:-ffast-math;-fno-finite-math-only>")

ggml/src/ggml-sycl/backend.hpp

@@ -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"

ggml/src/ggml-sycl/col2im-1d.cpp

@@ -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);
+    }
+}

ggml/src/ggml-sycl/col2im-1d.hpp

@@ -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

ggml/src/ggml-sycl/cpy.cpp

@@ -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));

ggml/src/ggml-sycl/cross_entropy_loss.cpp

@@ -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);
+            });
+    }
+}

ggml/src/ggml-sycl/cross_entropy_loss.hpp

@@ -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);

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

@@ -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
@@ -4770,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);
@@ -5078,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;
@@ -5090,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;
@@ -5639,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;
@@ -5707,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:
             {
@@ -5812,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];
@@ -5835,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) &&
@@ -5892,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;

ggml/src/ggml-sycl/presets.hpp

@@ -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

src/llama-graph.cpp

@@ -63,26 +63,6 @@ static bool can_reuse_kq_mask(
 
 // impl
 
-static ggml_tensor * ggml_mul_mat_aux(
-        ggml_context * ctx,
-        ggml_tensor * cur,
-        ggml_tensor * rot) {
-    const auto n = rot->ne[0];
-
-    ggml_tensor * res;
-
-    if (!ggml_is_contiguous(cur)) {
-        res = ggml_cont_2d   (ctx, cur, n, ggml_nelements(cur)/n);
-    } else {
-        res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
-    }
-    res = ggml_mul_mat   (ctx, rot, res);
-    ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
-    res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
-
-    return res;
-}
-
 void llm_graph_input_embd::set_input(const llama_ubatch * ubatch) {
     if (ubatch->token) {
         const int64_t n_tokens = ubatch->n_tokens;
@@ -881,6 +861,14 @@ void llm_graph_input_dsv4::set_input(const llama_ubatch * ubatch) {
     dsv4_set_comp_inputs(inp_hca, plan_hca, "hca", debug > 0, ubatch->n_tokens, n_stream);
     dsv4_set_comp_inputs(inp_lid, plan_lid, "lid", debug > 0, ubatch->n_tokens, n_stream);
 
+    if (inp_csa.k_rot && inp_csa.k_rot->buffer) {
+        mctx->get_csa()->set_input_k_rot(inp_csa.k_rot);
+    }
+
+    if (inp_hca.k_rot && inp_hca.k_rot->buffer) {
+        mctx->get_hca()->set_input_k_rot(inp_hca.k_rot);
+    }
+
     if (inp_lid.k_rot && inp_lid.k_rot->buffer) {
         mctx->get_lid()->set_input_k_rot(inp_lid.k_rot);
     }
@@ -2633,12 +2621,12 @@ ggml_tensor * llm_graph_context::build_attn(
     GGML_ASSERT(v_mla == nullptr);
 
     if (inp->self_k_rot) {
-        q_cur = ggml_mul_mat_aux(ctx0, q_cur, inp->self_k_rot);
-        k_cur = ggml_mul_mat_aux(ctx0, k_cur, inp->self_k_rot);
+        q_cur = llama_mul_mat_hadamard(ctx0, q_cur, inp->self_k_rot);
+        k_cur = llama_mul_mat_hadamard(ctx0, k_cur, inp->self_k_rot);
     }
 
     if (inp->self_v_rot) {
-        v_cur = ggml_mul_mat_aux(ctx0, v_cur, inp->self_v_rot);
+        v_cur = llama_mul_mat_hadamard(ctx0, v_cur, inp->self_v_rot);
     }
 
     // these nodes are added to the graph together so that they are not reordered
@@ -2669,7 +2657,7 @@ ggml_tensor * llm_graph_context::build_attn(
     cb(cur, "kqv_out", il);
 
     if (inp->self_v_rot) {
-        cur = ggml_mul_mat_aux(ctx0, cur, inp->self_v_rot);
+        cur = llama_mul_mat_hadamard(ctx0, cur, inp->self_v_rot);
     }
 
     if (wo) {
@@ -2874,14 +2862,14 @@ ggml_tensor * llm_graph_context::build_attn(
     auto * v_rot = is_swa ? inp->self_v_rot_swa : inp->self_v_rot;
 
     if (k_rot) {
-        q_cur = ggml_mul_mat_aux(ctx0, q_cur, k_rot);
+        q_cur = llama_mul_mat_hadamard(ctx0, q_cur, k_rot);
         if (k_cur) {
-            k_cur = ggml_mul_mat_aux(ctx0, k_cur, k_rot);
+            k_cur = llama_mul_mat_hadamard(ctx0, k_cur, k_rot);
         }
     }
     if (v_rot) {
         if (v_cur) {
-            v_cur = ggml_mul_mat_aux(ctx0, v_cur, v_rot);
+            v_cur = llama_mul_mat_hadamard(ctx0, v_cur, v_rot);
         }
     }
 
@@ -2924,7 +2912,7 @@ ggml_tensor * llm_graph_context::build_attn(
     cb(cur, "kqv_out", il);
 
     if (v_rot) {
-        cur = ggml_mul_mat_aux(ctx0, cur, v_rot);
+        cur = llama_mul_mat_hadamard(ctx0, cur, v_rot);
     }
 
     if (wo) {
@@ -3084,6 +3072,8 @@ llm_graph_input_dsv4 * llm_graph_context::build_inp_dsv4() const {
     dsv4_build_comp_inputs(ctx0, inp->inp_csa, mctx_cur->get_csa_plan(ubatch), "csa", n_stream);
     dsv4_build_comp_inputs(ctx0, inp->inp_hca, mctx_cur->get_hca_plan(ubatch), "hca", n_stream);
     dsv4_build_comp_inputs(ctx0, inp->inp_lid, mctx_cur->get_lid_plan(ubatch), "lid", n_stream);
+    inp->inp_csa.k_rot = mctx_cur->get_csa()->build_input_k_rot(ctx0);
+    inp->inp_hca.k_rot = mctx_cur->get_hca()->build_input_k_rot(ctx0);
     inp->inp_lid.k_rot = mctx_cur->get_lid()->build_input_k_rot(ctx0);
 
     return (llm_graph_input_dsv4 *) res->add_input(std::move(inp));

src/llama-impl.h

@@ -54,6 +54,26 @@ static inline dst_t llama_cast(src_t v) {
     }
 }
 
+static inline ggml_tensor * llama_mul_mat_hadamard(
+        ggml_context * ctx,
+        ggml_tensor * cur,
+        ggml_tensor * rot) {
+    const auto n = rot->ne[0];
+
+    ggml_tensor * res;
+
+    if (!ggml_is_contiguous(cur)) {
+        res = ggml_cont_2d(ctx, cur, n, ggml_nelements(cur)/n);
+    } else {
+        res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
+    }
+    res = ggml_mul_mat(ctx, rot, res);
+    ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
+    res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
+
+    return res;
+}
+
 struct time_meas {
     time_meas(int64_t & t_acc, bool disable = false);
     ~time_meas();

src/llama-kv-cache.cpp

@@ -57,22 +57,6 @@ static void ggml_gen_hadamard(ggml_tensor * tensor) {
     }
 }
 
-static ggml_tensor * ggml_mul_mat_aux(
-        ggml_context * ctx,
-        ggml_tensor * cur,
-        ggml_tensor * rot) {
-    const auto n = rot->ne[0];
-
-    ggml_tensor * res;
-
-    res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
-    res = ggml_mul_mat   (ctx, rot, res);
-    ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
-    res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
-
-    return res;
-}
-
 //
 // llama_kv_cache
 //
@@ -1875,14 +1859,14 @@ ggml_tensor * llama_kv_cache::build_rope_shift(
         tmp = ggml_cast(ctx, cur, GGML_TYPE_F32);
 
         // rotate back
-        tmp = ggml_mul_mat_aux(ctx, tmp, rot);
+        tmp = llama_mul_mat_hadamard(ctx, tmp, rot);
 
         tmp = ggml_rope_ext(ctx, tmp,
                 shift, factors, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
                 yarn_ext_factor, yarn_attn_factor, yarn_beta_fast, yarn_beta_slow);
 
         // rotate fwd
-        tmp = ggml_mul_mat_aux(ctx, tmp, rot);
+        tmp = llama_mul_mat_hadamard(ctx, tmp, rot);
 
         tmp = ggml_cpy(ctx, tmp, cur);
     } else {

src/models/deepseek4.cpp

@@ -557,7 +557,7 @@ ggml_tensor * llama_model_deepseek4::graph::build_lid_top_k(
     cb(indexer_q_pe, "lid_q_pe", il);
 
     indexer_q = ggml_concat(ctx0, indexer_q_nope, indexer_q_pe, 0);
-    indexer_q = ggml_mul_mat(ctx0, inp_lid.k_rot, indexer_q);
+    indexer_q = llama_mul_mat_hadamard(ctx0, indexer_q, inp_lid.k_rot);
     cb(indexer_q, "lid_q_rot", il);
 
     ggml_tensor * indexer_weights = build_lora_mm(layer.indexer_proj, cur);
@@ -652,10 +652,15 @@ ggml_tensor * llama_model_deepseek4::graph::build_csa_lid_attention(
         int il) const {
     const auto & inp_csa = inp_dsv4->get_csa();
     GGML_ASSERT(inp_csa.kq_mask);
-    GGML_ASSERT(inp_attn->self_k_rot == nullptr);
 
     ggml_tensor * top_k = build_lid_top_k(model, inp_dsv4, qr, cur, inp_pos, il);
 
+    ggml_tensor * k_rot = inp_attn->self_k_rot;
+    if (k_rot) {
+        q  = llama_mul_mat_hadamard(ctx0, q, k_rot);
+        kv = llama_mul_mat_hadamard(ctx0, kv, k_rot);
+    }
+
     ggml_build_forward_expand(gf, q);
     ggml_build_forward_expand(gf, kv);
 
@@ -696,6 +701,9 @@ ggml_tensor * llama_model_deepseek4::graph::build_csa_lid_attention(
 
     ggml_tensor * kq_b = dsv4_build_kq_zero_bias(ctx0, cparams, kq_mask, q->ne[1]);
     ggml_tensor * out = build_attn_mha(q, k_all, k_all, kq_b, kq_mask, sinks, nullptr, kq_scale, il);
+    if (k_rot) {
+        out = llama_mul_mat_hadamard(ctx0, out, k_rot);
+    }
     cb(out, "attn_csa_lid", il);
 
     return out;
@@ -711,7 +719,12 @@ ggml_tensor * llama_model_deepseek4::graph::build_hca_attention(
         int il) const {
     const auto & inp_hca = inp_dsv4->get_hca();
     GGML_ASSERT(inp_hca.kq_mask);
-    GGML_ASSERT(inp_attn->self_k_rot == nullptr);
+
+    ggml_tensor * k_rot = inp_attn->self_k_rot;
+    if (k_rot) {
+        q  = llama_mul_mat_hadamard(ctx0, q, k_rot);
+        kv = llama_mul_mat_hadamard(ctx0, kv, k_rot);
+    }
 
     ggml_build_forward_expand(gf, q);
     ggml_build_forward_expand(gf, kv);
@@ -753,6 +766,9 @@ ggml_tensor * llama_model_deepseek4::graph::build_hca_attention(
 
     ggml_tensor * kq_b = dsv4_build_kq_zero_bias(ctx0, cparams, kq_mask, q->ne[1]);
     ggml_tensor * out = build_attn_mha(q, k_all, k_all, kq_b, kq_mask, sinks, nullptr, kq_scale, il);
+    if (k_rot) {
+        out = llama_mul_mat_hadamard(ctx0, out, k_rot);
+    }
     cb(out, "attn_hca", il);
 
     return out;
@@ -770,8 +786,8 @@ ggml_tensor * llama_model_deepseek4::graph::build_raw_attention(
     ggml_tensor * k_rot = inp_attn->self_k_rot;
 
     if (k_rot) {
-        q  = ggml_mul_mat(ctx0, k_rot, q);
-        kv = ggml_mul_mat(ctx0, k_rot, kv);
+        q  = llama_mul_mat_hadamard(ctx0, q, k_rot);
+        kv = llama_mul_mat_hadamard(ctx0, kv, k_rot);
     }
 
     ggml_build_forward_expand(gf, q);
@@ -788,6 +804,9 @@ ggml_tensor * llama_model_deepseek4::graph::build_raw_attention(
 
     ggml_tensor * kq_b = dsv4_build_kq_zero_bias(ctx0, cparams, kq_mask, q->ne[1]);
     ggml_tensor * out = build_attn_mha(q, k, k, kq_b, kq_mask, sinks, nullptr, kq_scale, il);
+    if (k_rot) {
+        out = llama_mul_mat_hadamard(ctx0, out, k_rot);
+    }
     cb(out, "attn_raw", il);
 
     return out;
@@ -917,6 +936,11 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
                 "csa_state_compress",
                 il);
 
+        if (inp_dsv4->get_csa().k_rot) {
+            kv_comp_csa_state = llama_mul_mat_hadamard(ctx0, kv_comp_csa_state, inp_dsv4->get_csa().k_rot);
+            cb(kv_comp_csa_state, "csa_state_compress_rot", il);
+        }
+
         ggml_build_forward_expand(gf, inp_dsv4->mctx->get_csa()->cpy_k(ctx0,
                     kv_comp_csa_state, inp_dsv4->get_csa().state_write_idxs, il));
 
@@ -965,7 +989,7 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
                 il);
 
         if (inp_dsv4->get_lid().k_rot) {
-            kv_comp_lid_state = ggml_mul_mat(ctx0, inp_dsv4->get_lid().k_rot, kv_comp_lid_state);
+            kv_comp_lid_state = llama_mul_mat_hadamard(ctx0, kv_comp_lid_state, inp_dsv4->get_lid().k_rot);
             cb(kv_comp_lid_state, "lid_state_compress_rot", il);
         }
 
@@ -1007,6 +1031,11 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
                 "hca_state_compress",
                 il);
 
+        if (inp_dsv4->get_hca().k_rot) {
+            kv_comp_hca = llama_mul_mat_hadamard(ctx0, kv_comp_hca, inp_dsv4->get_hca().k_rot);
+            cb(kv_comp_hca, "hca_state_compress_rot", il);
+        }
+
         ggml_build_forward_expand(gf, inp_dsv4->mctx->get_hca()->cpy_k(ctx0,
                     kv_comp_hca, inp_dsv4->get_hca().state_write_idxs, il));
         hca_state_dep = kv_comp_hca;
@@ -1035,13 +1064,11 @@ ggml_tensor * llama_model_deepseek4::graph::build_attention(
     if (ratio == DSV4_CSA_RATIO &&
             inp_dsv4->get_csa().kq_mask &&
             inp_dsv4->get_lid().kq_mask &&
-            inp_dsv4->get_lid().k_rot &&
-            inp_attn->self_k_rot == nullptr) {
+            inp_dsv4->get_lid().k_rot) {
         out = build_csa_lid_attention(model, inp_dsv4, inp_attn, q, kv, qr, cur, inp_pos, layer.attn_sinks,
                 1.0f/sqrtf(float(n_embd_head)), il);
     } else if (ratio == DSV4_HCA_RATIO &&
-            inp_dsv4->get_hca().kq_mask &&
-            inp_attn->self_k_rot == nullptr) {
+            inp_dsv4->get_hca().kq_mask) {
         out = build_hca_attention(inp_dsv4, inp_attn, q, kv, layer.attn_sinks,
                 1.0f/sqrtf(float(n_embd_head)), il);
     } else {