Skip to content

Commit 8872ab5

Browse files
arthwarthw
authored
sycl : enhance set_rows to support q1_0, mxfp4, nvfp4 (ggml-org#24564)
1 parent 987fbd8 commit 8872ab5

4 files changed

Lines changed: 334 additions & 40 deletions

File tree

‎docs/ops/SYCL.csv‎

Lines changed: 36 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -582,42 +582,42 @@
582582
"SYCL0","SET_ROWS","type=q8_0,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=1","support","1","yes","SYCL"
583583
"SYCL0","SET_ROWS","type=q8_0,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=1","support","1","yes","SYCL"
584584
"SYCL0","SET_ROWS","type=q8_0,type_idx=i64,ne=[96,3,7,1],nr23=[2,3],r=2,v=1","support","1","yes","SYCL"
585-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=0","support","0","no","SYCL"
586-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=0","support","0","no","SYCL"
587-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,1,1],nr23=[2,3],r=2,v=0","support","0","no","SYCL"
588-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=1","support","0","no","SYCL"
589-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=1","support","0","no","SYCL"
590-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,1,1],nr23=[2,3],r=2,v=1","support","0","no","SYCL"
591-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=0","support","0","no","SYCL"
592-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=0","support","0","no","SYCL"
593-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,7,1],nr23=[2,3],r=2,v=0","support","0","no","SYCL"
594-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=1","support","0","no","SYCL"
595-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=1","support","0","no","SYCL"
596-
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,7,1],nr23=[2,3],r=2,v=1","support","0","no","SYCL"
597-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=0","support","0","no","SYCL"
598-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=0","support","0","no","SYCL"
599-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,1,1],nr23=[2,3],r=2,v=0","support","0","no","SYCL"
600-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=1","support","0","no","SYCL"
601-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=1","support","0","no","SYCL"
602-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,1,1],nr23=[2,3],r=2,v=1","support","0","no","SYCL"
603-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=0","support","0","no","SYCL"
604-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=0","support","0","no","SYCL"
605-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,7,1],nr23=[2,3],r=2,v=0","support","0","no","SYCL"
606-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=1","support","0","no","SYCL"
607-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=1","support","0","no","SYCL"
608-
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,7,1],nr23=[2,3],r=2,v=1","support","0","no","SYCL"
609-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=0","support","0","no","SYCL"
610-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=0","support","0","no","SYCL"
611-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,1,1],nr23=[2,3],r=2,v=0","support","0","no","SYCL"
612-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=1","support","0","no","SYCL"
613-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=1","support","0","no","SYCL"
614-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,1,1],nr23=[2,3],r=2,v=1","support","0","no","SYCL"
615-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=0","support","0","no","SYCL"
616-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=0","support","0","no","SYCL"
617-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,7,1],nr23=[2,3],r=2,v=0","support","0","no","SYCL"
618-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=1","support","0","no","SYCL"
619-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=1","support","0","no","SYCL"
620-
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,7,1],nr23=[2,3],r=2,v=1","support","0","no","SYCL"
585+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=0","support","1","yes","SYCL"
586+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=0","support","1","yes","SYCL"
587+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,1,1],nr23=[2,3],r=2,v=0","support","1","yes","SYCL"
588+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=1","support","1","yes","SYCL"
589+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=1","support","1","yes","SYCL"
590+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,1,1],nr23=[2,3],r=2,v=1","support","1","yes","SYCL"
591+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=0","support","1","yes","SYCL"
592+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=0","support","1","yes","SYCL"
593+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,7,1],nr23=[2,3],r=2,v=0","support","1","yes","SYCL"
594+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=1","support","1","yes","SYCL"
595+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=1","support","1","yes","SYCL"
596+
"SYCL0","SET_ROWS","type=q1_0,type_idx=i64,ne=[384,3,7,1],nr23=[2,3],r=2,v=1","support","1","yes","SYCL"
597+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=0","support","1","yes","SYCL"
598+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=0","support","1","yes","SYCL"
599+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,1,1],nr23=[2,3],r=2,v=0","support","1","yes","SYCL"
600+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=1","support","1","yes","SYCL"
601+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=1","support","1","yes","SYCL"
602+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,1,1],nr23=[2,3],r=2,v=1","support","1","yes","SYCL"
603+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=0","support","1","yes","SYCL"
604+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=0","support","1","yes","SYCL"
605+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,7,1],nr23=[2,3],r=2,v=0","support","1","yes","SYCL"
606+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=1","support","1","yes","SYCL"
607+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=1","support","1","yes","SYCL"
608+
"SYCL0","SET_ROWS","type=mxfp4,type_idx=i64,ne=[96,3,7,1],nr23=[2,3],r=2,v=1","support","1","yes","SYCL"
609+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=0","support","1","yes","SYCL"
610+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=0","support","1","yes","SYCL"
611+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,1,1],nr23=[2,3],r=2,v=0","support","1","yes","SYCL"
612+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=1","support","1","yes","SYCL"
613+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=1","support","1","yes","SYCL"
614+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,1,1],nr23=[2,3],r=2,v=1","support","1","yes","SYCL"
615+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=0","support","1","yes","SYCL"
616+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=0","support","1","yes","SYCL"
617+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,7,1],nr23=[2,3],r=2,v=0","support","1","yes","SYCL"
618+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,5,7,3],nr23=[1,1],r=1,v=1","support","1","yes","SYCL"
619+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[256,11,1,7],nr23=[2,3],r=7,v=1","support","1","yes","SYCL"
620+
"SYCL0","SET_ROWS","type=nvfp4,type_idx=i64,ne=[192,3,7,1],nr23=[2,3],r=2,v=1","support","1","yes","SYCL"
621621
"SYCL0","SET_ROWS","type=q2_K,type_idx=i64,ne=[256,5,1,3],nr23=[1,1],r=1,v=0","support","0","no","SYCL"
622622
"SYCL0","SET_ROWS","type=q2_K,type_idx=i64,ne=[256,11,1,1],nr23=[2,3],r=7,v=0","support","0","no","SYCL"
623623
"SYCL0","SET_ROWS","type=q2_K,type_idx=i64,ne=[768,3,1,1],nr23=[2,3],r=2,v=0","support","0","no","SYCL"

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

Lines changed: 281 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -48,6 +48,287 @@ inline void cpy_blck_f32_q8_0(const char * cxi, char * cdsti) {
4848
}
4949
}
5050

51+
inline void cpy_blck_f32_q1_0(const char * cxi, char * cdsti) {
52+
const float * xi = (const float *) cxi;
53+
block_q1_0 * dsti = (block_q1_0 *) cdsti;
54+
55+
float sum_abs = 0.0f;
56+
for (int j = 0; j < QK1_0; ++j) {
57+
sum_abs += sycl::fabs((float) xi[j]);
58+
}
59+
60+
dsti->d = sum_abs / QK1_0;
61+
62+
for (int j = 0; j < QK1_0 / 8; ++j) {
63+
dsti->qs[j] = 0;
64+
}
65+
66+
for (int j = 0; j < QK1_0; ++j) {
67+
if (xi[j] >= 0.0f) {
68+
dsti->qs[j / 8] |= (1u << (j % 8));
69+
}
70+
}
71+
}
72+
73+
inline int best_index_mxfp4(const float x, const float e) {
74+
int best_index = 0;
75+
float best_err = sycl::fabs((float) (kvalues_mxfp4[0] * e - x));
76+
for (int i = 1; i < 16; ++i) {
77+
const float err = sycl::fabs((float) (kvalues_mxfp4[i] * e - x));
78+
if (err < best_err) {
79+
best_index = i;
80+
best_err = err;
81+
}
82+
}
83+
return best_index;
84+
}
85+
86+
inline int nearest_int_sycl(float x) {
87+
const float val = x + 12582912.0f;
88+
int i;
89+
memcpy(&i, &val, sizeof(int));
90+
return (i & 0x007fffff) - 0x00400000;
91+
}
92+
93+
inline int nearest_int_ggml_sycl(float x) {
94+
return (int) sycl::round((float) x);
95+
}
96+
97+
inline uint8_t clamp_u8(const int x, const int lo, const int hi) {
98+
return (uint8_t) dpct::max(lo, dpct::min(hi, x));
99+
}
100+
101+
inline int8_t clamp_i8(const int x, const int lo, const int hi) {
102+
return (int8_t) dpct::max(lo, dpct::min(hi, x));
103+
}
104+
105+
constexpr float GROUP_MAX_EPS_SYCL = 1e-15f;
106+
107+
inline float make_qx_quants_sycl(int n, int nmax, const float * x, int8_t * L, int rmse_type, const float * qw) {
108+
float max = 0.0f;
109+
float amax = 0.0f;
110+
for (int i = 0; i < n; ++i) {
111+
const float ax = sycl::fabs(x[i]);
112+
if (ax > amax) {
113+
amax = ax;
114+
max = x[i];
115+
}
116+
}
117+
if (amax < GROUP_MAX_EPS_SYCL) {
118+
for (int i = 0; i < n; ++i) {
119+
L[i] = 0;
120+
}
121+
return 0.0f;
122+
}
123+
124+
float iscale = -nmax / max;
125+
if (rmse_type == 0) {
126+
for (int i = 0; i < n; ++i) {
127+
int l = nearest_int_ggml_sycl(iscale * x[i]);
128+
L[i] = (int8_t) (nmax + dpct::max(-nmax, dpct::min(nmax - 1, l)));
129+
}
130+
return 1.0f / iscale;
131+
}
132+
133+
bool return_early = false;
134+
if (rmse_type < 0) {
135+
rmse_type = -rmse_type;
136+
return_early = true;
137+
}
138+
139+
float sumlx = 0.0f;
140+
float suml2 = 0.0f;
141+
for (int i = 0; i < n; ++i) {
142+
int l = nearest_int_ggml_sycl(iscale * x[i]);
143+
l = dpct::max(-nmax, dpct::min(nmax - 1, l));
144+
L[i] = (int8_t) (l + nmax);
145+
146+
const float w = qw ? qw[i] : (rmse_type == 1 ? x[i] * x[i] :
147+
rmse_type == 2 ? 1.0f : rmse_type == 3 ? sycl::fabs(x[i]) : sycl::sqrt(sycl::fabs(x[i])));
148+
149+
sumlx += w * x[i] * l;
150+
suml2 += w * l * l;
151+
}
152+
153+
float scale = suml2 ? sumlx / suml2 : 0.0f;
154+
if (return_early) {
155+
return suml2 > 0.0f ? 0.5f * (scale + 1.0f / iscale) : 1.0f / iscale;
156+
}
157+
158+
float best = scale * sumlx;
159+
for (int is = -9; is <= 9; ++is) {
160+
if (is == 0) {
161+
continue;
162+
}
163+
iscale = -(nmax + 0.1f * is) / max;
164+
sumlx = 0.0f;
165+
suml2 = 0.0f;
166+
for (int i = 0; i < n; ++i) {
167+
int l = nearest_int_ggml_sycl(iscale * x[i]);
168+
l = dpct::max(-nmax, dpct::min(nmax - 1, l));
169+
const float w = qw ? qw[i] : (rmse_type == 1 ? x[i] * x[i] :
170+
rmse_type == 2 ? 1.0f : rmse_type == 3 ? sycl::fabs(x[i]) : sycl::sqrt(sycl::fabs(x[i])));
171+
sumlx += w * x[i] * l;
172+
suml2 += w * l * l;
173+
}
174+
175+
if (suml2 > 0.0f && sumlx * sumlx > best * suml2) {
176+
for (int i = 0; i < n; ++i) {
177+
int l = nearest_int_ggml_sycl(iscale * x[i]);
178+
L[i] = (int8_t) (nmax + dpct::max(-nmax, dpct::min(nmax - 1, l)));
179+
}
180+
scale = sumlx / suml2;
181+
best = scale * sumlx;
182+
}
183+
}
184+
185+
return scale;
186+
}
187+
188+
inline float make_q3_quants_sycl(int n, int nmax, const float * x, int8_t * L, bool do_rmse) {
189+
float max = 0.0f;
190+
float amax = 0.0f;
191+
for (int i = 0; i < n; ++i) {
192+
const float ax = sycl::fabs(x[i]);
193+
if (ax > amax) {
194+
amax = ax;
195+
max = x[i];
196+
}
197+
}
198+
199+
if (amax < GROUP_MAX_EPS_SYCL) {
200+
for (int i = 0; i < n; ++i) {
201+
L[i] = 0;
202+
}
203+
return 0.0f;
204+
}
205+
206+
const float iscale = -nmax / max;
207+
if (do_rmse) {
208+
float sumlx = 0.0f;
209+
float suml2 = 0.0f;
210+
for (int i = 0; i < n; ++i) {
211+
int l = nearest_int_ggml_sycl(iscale * x[i]);
212+
l = dpct::max(-nmax, dpct::min(nmax - 1, l));
213+
L[i] = (int8_t) l;
214+
const float w = x[i] * x[i];
215+
sumlx += w * x[i] * l;
216+
suml2 += w * l * l;
217+
}
218+
219+
for (int itry = 0; itry < 5; ++itry) {
220+
int n_changed = 0;
221+
for (int i = 0; i < n; ++i) {
222+
const float w = x[i] * x[i];
223+
float slx = sumlx - w * x[i] * L[i];
224+
if (slx > 0.0f) {
225+
float sl2 = suml2 - w * L[i] * L[i];
226+
int new_l = nearest_int_ggml_sycl(x[i] * sl2 / slx);
227+
new_l = dpct::max(-nmax, dpct::min(nmax - 1, new_l));
228+
if (new_l != L[i]) {
229+
slx += w * x[i] * new_l;
230+
sl2 += w * new_l * new_l;
231+
if (sl2 > 0.0f && slx * slx * suml2 > sumlx * sumlx * sl2) {
232+
L[i] = (int8_t) new_l;
233+
sumlx = slx;
234+
suml2 = sl2;
235+
++n_changed;
236+
}
237+
}
238+
}
239+
}
240+
if (!n_changed) {
241+
break;
242+
}
243+
}
244+
245+
for (int i = 0; i < n; ++i) {
246+
L[i] += nmax;
247+
}
248+
return suml2 > 0.0f ? sumlx / suml2 : 0.0f;
249+
}
250+
251+
for (int i = 0; i < n; ++i) {
252+
int l = nearest_int_ggml_sycl(iscale * x[i]);
253+
l = dpct::max(-nmax, dpct::min(nmax - 1, l));
254+
L[i] = (int8_t) (l + nmax);
255+
}
256+
257+
return 1.0f / iscale;
258+
}
259+
260+
inline void set_scale_min_k4(int j, uint8_t * q, uint8_t d, uint8_t m) {
261+
if (j < 4) {
262+
q[j] = (q[j] & 0xC0) | (d & 0x3F);
263+
q[j + 4] = (q[j + 4] & 0xC0) | (m & 0x3F);
264+
} else {
265+
q[j + 4] = (d & 0x0F) | ((m & 0x0F) << 4);
266+
q[j - 4] = (q[j - 4] & 0x3F) | ((d >> 4) << 6);
267+
q[j - 0] = (q[j - 0] & 0x3F) | ((m >> 4) << 6);
268+
}
269+
}
270+
271+
inline void get_scale_min_k4_local(int j, const uint8_t * q, uint8_t & d, uint8_t & m) {
272+
if (j < 4) {
273+
d = q[j] & 63;
274+
m = q[j + 4] & 63;
275+
} else {
276+
d = (q[j + 4] & 0xF) | ((q[j - 4] >> 6) << 4);
277+
m = (q[j + 4] >> 4) | ((q[j - 0] >> 6) << 4);
278+
}
279+
}
280+
281+
inline void cpy_blck_f32_mxfp4(const char * cxi, char * cdsti) {
282+
const float * xi = (const float *) cxi;
283+
block_mxfp4 * dsti = (block_mxfp4 *) cdsti;
284+
285+
float amax = 0.0f;
286+
for (int j = 0; j < QK_MXFP4; ++j) {
287+
amax = sycl::fmax(amax, sycl::fabs((float) xi[j]));
288+
}
289+
290+
const uint8_t e = amax > 0.0f ? (uint8_t) (sycl::floor(sycl::log2(amax)) - 2 + 127) : 0;
291+
const float d = GGML_E8M0_TO_FP32_HALF(e);
292+
293+
dsti->e = e;
294+
295+
for (int j = 0; j < QK_MXFP4 / 2; ++j) {
296+
const uint8_t x0 = best_index_mxfp4(xi[0 + j], d);
297+
const uint8_t x1 = best_index_mxfp4(xi[QK_MXFP4 / 2 + j], d);
298+
299+
dsti->qs[j] = x0;
300+
dsti->qs[j] |= x1 << 4;
301+
}
302+
}
303+
304+
inline void cpy_blck_f32_nvfp4(const char * cxi, char * cdsti) {
305+
const float * xi = (const float *) cxi;
306+
block_nvfp4 * dsti = (block_nvfp4 *) cdsti;
307+
308+
constexpr int n_sub = QK_NVFP4 / QK_NVFP4_SUB;
309+
310+
for (int s = 0; s < n_sub; ++s) {
311+
const float * xb = xi + s * QK_NVFP4_SUB;
312+
313+
float amax = 0.0f;
314+
for (int j = 0; j < QK_NVFP4_SUB; ++j) {
315+
amax = sycl::fmax(amax, sycl::fabs((float) xb[j]));
316+
}
317+
318+
const uint8_t ue = ggml_fp32_to_ue4m3(amax / 6.0f);
319+
dsti->d[s] = ue;
320+
const float d = ggml_ue4m3_to_fp32(ue);
321+
322+
for (int j = 0; j < QK_NVFP4_SUB / 2; ++j) {
323+
const uint8_t x0 = best_index_mxfp4(xb[0 + j], d);
324+
const uint8_t x1 = best_index_mxfp4(xb[QK_NVFP4_SUB / 2 + j], d);
325+
326+
dsti->qs[s * (QK_NVFP4_SUB / 2) + j] = x0 | (x1 << 4);
327+
}
328+
}
329+
}
330+
331+
51332
inline void cpy_blck_f32_q4_0(const char * cxi, char * cdsti) {
52333
const float * xi = (const float *) cxi;
53334
block_q4_0 * dsti = (block_q4_0 *) cdsti;

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

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5242,10 +5242,15 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
52425242

52435243
case GGML_OP_SET_ROWS:
52445244
{
5245-
return ((op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 ||
5245+
5246+
auto res = ((op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 ||
52465247
op->type == GGML_TYPE_Q8_0 || op->type == GGML_TYPE_Q5_1 || op->type == GGML_TYPE_Q5_0 ||
5247-
op->type == GGML_TYPE_Q4_1 || op->type == GGML_TYPE_Q4_0 || op->type == GGML_TYPE_IQ4_NL) &&
5248+
op->type == GGML_TYPE_Q1_0 ||
5249+
op->type == GGML_TYPE_Q4_1 || op->type == GGML_TYPE_Q4_0 || op->type == GGML_TYPE_IQ4_NL ||
5250+
op->type == GGML_TYPE_MXFP4 || op->type == GGML_TYPE_NVFP4) &&
5251+
op->src[0]->type == GGML_TYPE_F32 &&
52485252
(op->src[1]->type == GGML_TYPE_I64 || op->src[1]->type == GGML_TYPE_I32));
5253+
return res;
52495254
}
52505255
break;
52515256
case GGML_OP_CPY:

0 commit comments

Comments
 (0)