sycl: s2_expert_grouped_parity passes - no FMA contraction in the transposed grouped kernels

The transposed grouped kernels (gu/down_grouped_t_kernel) write `acc += dw * dx * v` in one expression, which
-fp-model=precise contracts to an FMA; the previous kernels compute the product in chunk_dot and add it in the
caller, which it does not. nvcc contracts both, so they are bitwise equal on CUDA; on the B70 the last bit
differed and s2_expert_grouped_parity reported 4 failures (moe_grouped_s2 / moe_group_resident, both within 1e-7
of the host reference). `#pragma clang fp contract(off)` in the two kernels, kept by fixups.py (9d) across
re-migrations; 9c (the snapshot test's sycl::free) is back in fixups.py for the same reason.

sycl/src/kernels/cuda/s2_expert_grouped.dp.cpp

@@ -1350,6 +1350,7 @@ __dpct_inline__ void gu_grouped_t_kernel(
     const int32_t *__restrict__ ent_tok, const uint8_t *__restrict__ x_q8_0,
     const float *__restrict__ x_scales, float *__restrict__ gate_up,
     int cap_entries) {
+    #pragma clang fp contract(off)   // SYCL port: round as the previous kernels do (see tools/fixups.py)
     auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
     constexpr int NC = H / 32;
     auto &xs_w = *sycl::ext::oneapi::group_local_memory_for_overwrite<
@@ -1431,6 +1432,7 @@ __dpct_inline__ void down_grouped_t_kernel(
     const int32_t *__restrict__ grp_start, const int32_t *__restrict__ n_groups,
     const int32_t *__restrict__ ent_dst, const uint8_t *__restrict__ h_q8_0,
     const float *__restrict__ h_scales, float *__restrict__ out) {
+    #pragma clang fp contract(off)   // SYCL port: round as the previous kernels do (see tools/fixups.py)
     auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
     constexpr int NC = FF / 32;
     auto &hs_w = *sycl::ext::oneapi::group_local_memory_for_overwrite<

sycl/tools/fixups.py

@@ -234,6 +234,26 @@ edit("src/program/generate.cpp", lambda s: s.replace(
     "gs ? gs->adapt_stream\n                           : adapt_stream->memcpy(",
     "(gs ? gs->adapt_stream : adapt_stream)->memcpy("))
 
+# 9c. conversation_snapshot_test: dpct migrated cudaFreeHost(p) on sycl::malloc_host memory to libc free(p) - the
+#     test aborted with "munmap_chunk(): invalid pointer". USM host memory goes back through sycl::free.
+edit("src/core/conversation_snapshot_test.cpp", lambda s: s.replace(
+    "for (void *p : host) DPCT_CHECK_ERROR(free(p));",
+    "for (void *p : host) DPCT_CHECK_ERROR(sycl::free(p, dpct::get_in_order_queue()));   // SYCL port: dpct wrote free(p) for cudaFreeHost"))
+
+# 9d. s2_expert_grouped: the new transposed kernels (gu/down_grouped_t_kernel) write `acc += dw * dx * v` in one
+#     expression, which -fp-model=precise contracts to an FMA; the previous kernels compute the product in
+#     chunk_dot and add it in the caller, which it does not. nvcc contracts both, so on CUDA they are bitwise equal
+#     (s2_expert_grouped_parity checks that); here they differed in the last bit. No contraction in the new ones.
+def s2_no_contract(s):
+    import re as _re
+    note = "    #pragma clang fp contract(off)   // SYCL port: round as the previous kernels do (see tools/fixups.py)\n"
+    for name in ("gu_grouped_t_kernel", "down_grouped_t_kernel"):
+        m = _re.search(r"__dpct_inline__ void " + name + r"\([^{]*\{\n", s)
+        if m and not s[m.end():].startswith(note):
+            s = s[:m.end()] + note + s[m.end():]
+    return s
+edit("src/kernels/cuda/s2_expert_grouped.dp.cpp", s2_no_contract)
+
 # 10. %globaltimer: there is no device-side wall clock in SPIR-V; the verify-window stage profiler reads zeros.
 edit("src/kernels/cuda/verify_kernels.dp.cpp", lambda s: s.replace(
     'asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(t));',