Niko1221/Strata
1
0 Fork 0
mirror of https://github.com/Niko1221/Strata.git synced 2026-11-07 14:21:06 -10:00
Code Issues Packages Projects Releases 1 Wiki Activity

sycl: gr_parity skips the STRATA_QFUSE read check - the port's fused GR read does not write q8_1

Browse Source 
fused_gr_read_multi on SYCL never writes the q8_1 image of `mixed` (fused_gr.dp.cpp returns false: the caller
quantizes), so gr_parity's S26 check failed on every SYCL device with "QFUSE: T=1 rep 0: not written". The
engine's path is fine (verify.cpp quantizes when the read returns false); the test now says "skipped" when the
read writes nothing at T=1, and a read that writes different bytes still fails.
This commit is contained in:
maxfridbe committed 2026-10-19 12:09:29 -05:00
1 parent 5fbef34f98
commit 344e43e1d5
1 file changed
+6 -2
Download Patch File Download Diff File Expand all files Collapse all files
sycl/src/kernels/gr_parity.cpp
+6 -2
View File
@@ -615,8 +615,9 @@ int fused_multi_lds_parity(const float* d_norm, const uint16_t* d_down, const ui
.wait()),
"qfuse counters zero");
int qbad = 0;
bool qskip = false; // SYCL port: see below
std::vector<uint8_t> hq(qbytes), hr(qbytes);
for (int tt = 1; tt <= T && !qbad; ++tt) {
for (int tt = 1; tt <= T && !qbad && !qskip; ++tt) {
std::vector<FusedGrArgs> qa(args.begin(), args.begin() + tt);
for (int t = 0; t < tt; ++t) { qa[t].q8_mixed = d_q + (size_t) t * (N / 32) * 36; qa[t].q8_cnt = d_cnt; }
for (int rep = 0; rep < 3 && !qbad; ++rep) {
@@ -653,6 +654,9 @@ int fused_multi_lds_parity(const float* d_norm, const uint16_t* d_down, const ui
check(DPCT_CHECK_ERROR(stream->wait()), "qfuse sync");
if (qx) { delete (qx); delete (qg); }
if (v3) { std::printf(" QFUSE: the v3 read writes no q8_1 (%s)\n", wrote ? "WRONG: it says it did" : "ok"); qbad += wrote; break; }
// SYCL port: fused_gr_read_multi never writes the q8_1 image here (fused_gr.dp.cpp: the caller quantizes),
// so there is nothing to compare - the engine's STRATA_QFUSE path falls back to native_quantize_q8_1
if (!wrote && tt == 1 && rep == 0) { std::printf(" QFUSE: not built in the SYCL port (the caller quantizes): skipped\n"); qskip = true; break; }
check(DPCT_CHECK_ERROR((dpct::get_current_device().queues_wait_and_throw(), dpct::get_in_order_queue())
.memcpy(hq.data(), d_q,
(size_t)tt * (N / 32) * 36)
@@ -669,7 +673,7 @@ int fused_multi_lds_parity(const float* d_norm, const uint16_t* d_down, const ui
}
}
}
std::printf(" fused GR read + q8_1 (STRATA_QFUSE), T 1..8, direct and graph replays: %s\n", qbad ? "FAIL" : "pass");
std::printf(" fused GR read + q8_1 (STRATA_QFUSE), T 1..8, direct and graph replays: %s\n", qbad ? "FAIL" : qskip ? "skipped" : "pass");
bad += qbad;
sycl::free(d_q, dpct::get_in_order_queue());
sycl::free(d_ref, dpct::get_in_order_queue());