For each kernel, add a standalone path to build the cpp source.

README.md

@@ -14,7 +14,7 @@ Paper: [SyclKittens: A Tile Programming Model for Programmers and Coding Agents
 - Intel oneAPI DPC++ compiler with `icpx` and SYCL 2020 support
 - Intel Data Center GPU Max (PVC) for the supplied AOT target
 - C++20
-- XPU-enabled PyTorch and pybind11 for Python extension kernels
+- XPU-enabled PyTorch and pybind11 for Python extension kernels (`CONFIG=pytorch`)
 - oneMKL for the decode GEMV fallback
 - Level Zero development headers for the direct-copy all-gather and distributed
   GEMM targets; set
@@ -78,10 +78,36 @@ make -C kernels/collective all_reduce
 make -C kernels/context_attention ring_overlap
 ```
 
-Python extension targets use `PYTHON` to select an XPU-enabled environment:
+### Build modes
+
+Each kernel in `attention`, `gemm`, `norm`, `rotary`, `activation`, and `decode`
+builds in one of two modes, selected with `CONFIG`:
+
+| `CONFIG` | Output | Needs |
+|---|---|---|
+| `standalone` (default) | executable; `main()` comes from `<kernel>_harness.impl` | SYCL only |
+| `pytorch` | Python extension (`sk_<kernel>*.so`) | XPU PyTorch + pybind11 |
+
+A standalone run times the kernel and checks it against a CPU fp32 reference;
+arguments are `[iters] [warmup]`, and the exit code is non-zero on failure:
+
+```bash
+make -C kernels/attention forward
+./kernels/attention/build/gqa_forward 100 10
+```
+
+To build and run every standalone harness at once (summary at the end, non-zero
+exit on any failure):
+
+```bash
+./kernels/run_standalone.sh [iters] [warmup]
+```
+
+Python extension targets use `CONFIG=pytorch` and `PYTHON` to select an
+XPU-enabled environment:
 
 ```bash
-make -C kernels/norm all \
+make -C kernels/norm all CONFIG=pytorch \
   PYTHON=/path/to/xpu/python3 RMSNORM_D=4096 LN_D=4096
 ```
 

kernels/activation/Makefile

@@ -4,7 +4,13 @@ include ../common.mk
 
 all: silu_mul
 
+ifeq ($(CONFIG),standalone)
+silu_mul: check-sycl
+	@mkdir -p $(BUILD_DIR)
+	$(CXX) $(COMMON_SYCL_FLAGS) silu_mul.dp.cpp -o $(BUILD_DIR)/silu_mul
+else
 silu_mul: check-python
 	@mkdir -p $(BUILD_DIR)
 	$(CXX) $(EXTENSION_FLAGS) -DTORCH_EXTENSION_NAME=sk_silu_mul silu_mul.dp.cpp \
 	  -o $(BUILD_DIR)/sk_silu_mul$(EXT_SUFFIX) $(TORCH_LIBS)
+endif

kernels/activation/silu_mul.dp.cpp

@@ -28,10 +28,6 @@
 
 #include <sycl/sycl.hpp>
 
-#include "pyutils/torch_helpers.dp.hpp"
-#include <torch/extension.h>
-#include <c10/xpu/XPUStream.h>
-
 #ifndef SILU_VEC
 #define SILU_VEC 8
 #endif
@@ -111,10 +107,15 @@ void launch_silu_mul(
     // NO queue.wait() — caller synchronizes via torch.xpu.synchronize()
 }
 
+#ifdef TORCH_COMPILE
 // ============================================================
 // PyTorch dispatch
 // ============================================================
 
+#include "pyutils/torch_helpers.dp.hpp"
+#include <torch/extension.h>
+#include <c10/xpu/XPUStream.h>
+
 torch::Tensor dispatch_silu_mul(
     torch::Tensor gate,   // (B, N, inter), bf16
     torch::Tensor up      // (B, N, inter), bf16
@@ -153,3 +154,8 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
           "Fused SiLU*mul: gate,up (bf16 [B,N,inter]) -> out (bf16 [B,N,inter])",
           py::arg("gate"), py::arg("up"));
 }
+
+#else
+// Standalone executable (CONFIG=standalone): main() lives in the harness.
+#include "silu_mul_harness.impl"
+#endif  // TORCH_COMPILE

kernels/activation/silu_mul_harness.impl

@@ -0,0 +1,78 @@
+// Standalone harness for silu_mul.dp.cpp (CONFIG=standalone, no PyTorch).
+// Included at the end of silu_mul.dp.cpp when TORCH_COMPILE is not defined,
+// so bf16_t and launch_silu_mul() are already in scope.
+//
+//   usage: silu_mul [iters=100] [warmup=10]
+//   pick the GPU with ONEAPI_DEVICE_SELECTOR / ZE_AFFINITY_MASK.
+
+#include <chrono>
+#include <cmath>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <vector>
+
+// Deterministic values in [-4, 4) so the run is reproducible.
+static float lcg_uniform(uint32_t &state) {
+    state = state * 1664525u + 1013904223u;
+    return (static_cast<float>(state >> 8) / static_cast<float>(1u << 23) - 1.0f) * 4.0f;
+}
+
+static uint16_t to_bits(float f) { return sycl::bit_cast<uint16_t>(bf16_t(f)); }
+static float from_bits(uint16_t b) { return static_cast<float>(sycl::bit_cast<bf16_t>(b)); }
+
+int main(int argc, char **argv) {
+    const int iters  = argc > 1 ? std::atoi(argv[1]) : 100;
+    const int warmup = argc > 2 ? std::atoi(argv[2]) : 10;
+    constexpr int INTER = 14336;
+
+    sycl::queue queue{sycl::gpu_selector_v, sycl::property::queue::in_order{}};
+    std::printf("Device: %s\n",
+                queue.get_device().get_info<sycl::info::device::name>().c_str());
+
+    bool all_pass = true;
+    for (int T : {1, 2048, 8192}) {
+        const size_t total = size_t(T) * INTER;
+        std::vector<uint16_t> h_gate(total), h_up(total), h_out(total);
+        uint32_t seed = 12345u + T;
+        for (auto &x : h_gate) x = to_bits(lcg_uniform(seed));
+        for (auto &x : h_up)   x = to_bits(lcg_uniform(seed));
+
+        uint16_t *d_gate = sycl::malloc_device<uint16_t>(total, queue);
+        uint16_t *d_up   = sycl::malloc_device<uint16_t>(total, queue);
+        uint16_t *d_out  = sycl::malloc_device<uint16_t>(total, queue);
+        queue.memcpy(d_gate, h_gate.data(), total * sizeof(uint16_t));
+        queue.memcpy(d_up, h_up.data(), total * sizeof(uint16_t));
+        queue.wait();
+
+        for (int i = 0; i < warmup; i++) launch_silu_mul(queue, d_gate, d_up, d_out, total);
+        queue.wait();
+        const auto t0 = std::chrono::high_resolution_clock::now();
+        for (int i = 0; i < iters; i++) launch_silu_mul(queue, d_gate, d_up, d_out, total);
+        queue.wait();
+        const auto t1 = std::chrono::high_resolution_clock::now();
+        const double ms = std::chrono::duration<double, std::milli>(t1 - t0).count() / iters;
+        const double gbps = 3.0 * total * sizeof(uint16_t) / (ms * 1e-3) / 1e9;   // r gate, r up, w out
+
+        // Full CPU reference (fp32 math on the bf16 inputs).
+        queue.memcpy(h_out.data(), d_out, total * sizeof(uint16_t)).wait();
+        double err2 = 0.0, ref2 = 0.0;
+        for (size_t i = 0; i < total; i++) {
+            const float g = from_bits(h_gate[i]), u = from_bits(h_up[i]);
+            const double ref = double(g / (1.0f + std::exp(-g)) * u);
+            const double d = double(from_bits(h_out[i])) - ref;
+            err2 += d * d;
+            ref2 += ref * ref;
+        }
+        const double rel = std::sqrt(err2 / (ref2 + 1e-30));
+        const bool pass = rel < 1e-2;
+        all_pass &= pass;
+        std::printf("silu_mul T=%d inter=%d: %.4f ms  %.1f GB/s  (iters=%d warmup=%d)  rel_l2=%.2e  %s\n",
+                    T, INTER, ms, gbps, iters, warmup, rel, pass ? "PASS" : "FAIL");
+
+        sycl::free(d_gate, queue);
+        sycl::free(d_up, queue);
+        sycl::free(d_out, queue);
+    }
+    return all_pass ? 0 : 1;
+}

kernels/attention/Makefile

@@ -7,22 +7,33 @@ BWD_H ?= 16
 BWD_H_KV ?= 16
 BWD_D ?= 128
 
+FWD_DEFS := -DATTN_H_VALUE=$(FWD_H) -DATTN_H_KV_VALUE=$(FWD_H_KV) -DATTN_D_VALUE=$(FWD_D)
+BWD_DEFS := -DATTN_H_VALUE=$(BWD_H) -DATTN_H_KV_VALUE=$(BWD_H_KV) -DATTN_D_VALUE=$(BWD_D) \
+            -ftarget-register-alloc-mode=pvc:large
+
 .PHONY: all forward backward
 
 all: forward backward
 
+ifeq ($(CONFIG),standalone)
+forward: check-sycl
+	@mkdir -p $(BUILD_DIR)
+	$(CXX) $(COMMON_SYCL_FLAGS) $(FWD_DEFS) gqa_forward.dp.cpp -o $(BUILD_DIR)/gqa_forward
+
+backward: check-sycl
+	@mkdir -p $(BUILD_DIR)
+	$(CXX) $(COMMON_SYCL_FLAGS) $(BWD_DEFS) gqa_backward.dp.cpp -o $(BUILD_DIR)/gqa_backward
+else
 forward: check-python
 	@mkdir -p $(BUILD_DIR)
-	$(CXX) $(EXTENSION_FLAGS) \
-	  -DATTN_H_VALUE=$(FWD_H) -DATTN_H_KV_VALUE=$(FWD_H_KV) -DATTN_D_VALUE=$(FWD_D) \
+	$(CXX) $(EXTENSION_FLAGS) $(FWD_DEFS) \
 	  -DTORCH_EXTENSION_NAME=sk_gqa_forward gqa_forward.dp.cpp \
 	  -o $(BUILD_DIR)/sk_gqa_forward$(EXT_SUFFIX) $(TORCH_LIBS)
 
 backward: check-python
 	@mkdir -p $(BUILD_DIR)
-	$(CXX) $(EXTENSION_FLAGS) \
-	  -DATTN_H_VALUE=$(BWD_H) -DATTN_H_KV_VALUE=$(BWD_H_KV) -DATTN_D_VALUE=$(BWD_D) \
-	  -ftarget-register-alloc-mode=pvc:large \
+	$(CXX) $(EXTENSION_FLAGS) $(BWD_DEFS) \
 	  -DTORCH_EXTENSION_NAME=sk_gqa_backward gqa_backward.dp.cpp \
 	  -o $(BUILD_DIR)/sk_gqa_backward$(EXT_SUFFIX) \
 	  $(TORCH_LIBS)
+endif

kernels/attention/gqa_backward.dp.cpp

@@ -11,10 +11,6 @@
 
 #include "kittens.dp.hpp"
 #include "ops/warp/register/tile/conversions.dp.hpp"
-#include "pyutils/torch_helpers.dp.hpp"
-
-#include <torch/extension.h>
-#include <c10/xpu/XPUStream.h>
 using namespace kittens;
 
 // ============================================================
@@ -370,21 +366,13 @@ SYCL_EXTERNAL void attend_bwd_fused_ker(
 }
 
 // ============================================================
-// Dispatch
+// Shared raw-pointer entries (both builds). All tensors BNHD except
+// L/delta [B,H,N] f32. Q/K/V/O/dO bf16; dQ/dK/dV f32.
 // ============================================================
 
-void dispatch_prep(torch::Tensor O, torch::Tensor dO, torch::Tensor delta) {
-    CHECK_INPUT(O); CHECK_INPUT(dO); CHECK_INPUT(delta);
-    const int B = O.size(0), N = O.size(1), H = O.size(2);
-
-    auto stream = c10::xpu::getCurrentXPUStream(O.device().index());
-    auto &queue = stream.queue();
-    queue.wait();
-
-    bf16  *d_O_bf  = reinterpret_cast<bf16*>(O.data_ptr<c10::BFloat16>());
-    bf16  *d_dO_bf = reinterpret_cast<bf16*>(dO.data_ptr<c10::BFloat16>());
-    float *d_delta = delta.data_ptr<float>();
-
+// delta[b,h,q] = rowsum(dO * O). Blocks until done.
+static void run_prep(sycl::queue &queue, bf16 *d_O_bf, bf16 *d_dO_bf,
+                     float *d_delta, int B, int N, int H) {
     bwd_gl<ATTN_D> Og(d_O_bf,  B, N, H, nullptr);
     bwd_gl<ATTN_D> dOg(d_dO_bf, B, N, H, nullptr);
 
@@ -407,43 +395,13 @@ void dispatch_prep(torch::Tensor O, torch::Tensor dO, torch::Tensor delta) {
     queue.wait();
 }
 
-// Allocate the P^T / dS^T reorder scratch: [B, kv_blocks*128, H_KV, 128] bf16.
-static torch::Tensor make_reorder_scratch(int B, int N, int H_KV, torch::Tensor like) {
-    auto opts = torch::TensorOptions().dtype(torch::kBFloat16).device(like.device());
-    return torch::empty({(long)B, (long)N, (long)H_KV, (long)ATTN_D}, opts);
-}
-
-void dispatch_bwd(torch::Tensor Q, torch::Tensor K, torch::Tensor V,
-                  torch::Tensor dO, torch::Tensor L, torch::Tensor delta,
-                  torch::Tensor dQ, torch::Tensor dK, torch::Tensor dV) {
-    CHECK_INPUT(Q); CHECK_INPUT(K); CHECK_INPUT(V); CHECK_INPUT(dO);
-    CHECK_INPUT(L); CHECK_INPUT(delta);
-    CHECK_INPUT(dQ); CHECK_INPUT(dK); CHECK_INPUT(dV);
-
-    const int B = Q.size(0), N = Q.size(1), H = Q.size(2), H_KV = K.size(2);
-
-    auto stream = c10::xpu::getCurrentXPUStream(Q.device().index());
-    auto &queue = stream.queue();
-    queue.wait();
-
-    bf16  *d_Q  = reinterpret_cast<bf16*>(Q.data_ptr<c10::BFloat16>());
-    bf16  *d_K  = reinterpret_cast<bf16*>(K.data_ptr<c10::BFloat16>());
-    bf16  *d_V  = reinterpret_cast<bf16*>(V.data_ptr<c10::BFloat16>());
-    bf16  *d_dO = reinterpret_cast<bf16*>(dO.data_ptr<c10::BFloat16>());
-    float *d_L     = L.data_ptr<float>();
-    float *d_delta = delta.data_ptr<float>();
-    float *d_dQ    = dQ.data_ptr<float>();
-    float *d_dK    = dK.data_ptr<float>();
-    float *d_dV    = dV.data_ptr<float>();
-
-    // Reorder scratch (bf16 P^T and dS^T), shaped [B, N(kv), H_KV, D(q=128)].
-    torch::Tensor Pt  = make_reorder_scratch(B, N, H_KV, Q);
-    torch::Tensor dSt = make_reorder_scratch(B, N, H_KV, Q);
-    torch::Tensor dSq = make_reorder_scratch(B, N, H_KV, Q);  // dQ q-major scratch
-    bf16 *d_Pt  = reinterpret_cast<bf16*>(Pt.data_ptr<c10::BFloat16>());
-    bf16 *d_dSt = reinterpret_cast<bf16*>(dSt.data_ptr<c10::BFloat16>());
-    bf16 *d_dSq = reinterpret_cast<bf16*>(dSq.data_ptr<c10::BFloat16>());
-
+// Fused causal dV/dK/dQ. d_Pt/d_dSt/d_dSq are caller-owned bf16 reorder scratch
+// buffers of B*N*H_KV*ATTN_D elements each. Blocks until done.
+static void run_bwd(sycl::queue &queue, bf16 *d_Q, bf16 *d_K, bf16 *d_V,
+                    bf16 *d_dO, float *d_L, float *d_delta,
+                    float *d_dQ, float *d_dK, float *d_dV,
+                    bf16 *d_Pt, bf16 *d_dSt, bf16 *d_dSq,
+                    int B, int N, int H, int H_KV) {
     bwd_gl<ATTN_D> Qg(d_Q, B, N, H, nullptr);
     bwd_gl<ATTN_D> Kg(d_K, B, N, H_KV, nullptr);
     bwd_gl<ATTN_D> Vg(d_V, B, N, H_KV, nullptr);
@@ -476,7 +434,70 @@ void dispatch_bwd(torch::Tensor Q, torch::Tensor K, torch::Tensor V,
     }
 }
 
+#ifdef TORCH_COMPILE
+// ============================================================
+// PyTorch dispatch
+// ============================================================
+#include "pyutils/torch_helpers.dp.hpp"
+#include <torch/extension.h>
+#include <c10/xpu/XPUStream.h>
+
+void dispatch_prep(torch::Tensor O, torch::Tensor dO, torch::Tensor delta) {
+    CHECK_INPUT(O); CHECK_INPUT(dO); CHECK_INPUT(delta);
+    const int B = O.size(0), N = O.size(1), H = O.size(2);
+
+    auto stream = c10::xpu::getCurrentXPUStream(O.device().index());
+    auto &queue = stream.queue();
+    queue.wait();
+
+    run_prep(queue, reinterpret_cast<bf16*>(O.data_ptr<c10::BFloat16>()),
+             reinterpret_cast<bf16*>(dO.data_ptr<c10::BFloat16>()),
+             delta.data_ptr<float>(), B, N, H);
+}
+
+// Allocate the P^T / dS^T reorder scratch: [B, kv_blocks*128, H_KV, 128] bf16.
+static torch::Tensor make_reorder_scratch(int B, int N, int H_KV, torch::Tensor like) {
+    auto opts = torch::TensorOptions().dtype(torch::kBFloat16).device(like.device());
+    return torch::empty({(long)B, (long)N, (long)H_KV, (long)ATTN_D}, opts);
+}
+
+void dispatch_bwd(torch::Tensor Q, torch::Tensor K, torch::Tensor V,
+                  torch::Tensor dO, torch::Tensor L, torch::Tensor delta,
+                  torch::Tensor dQ, torch::Tensor dK, torch::Tensor dV) {
+    CHECK_INPUT(Q); CHECK_INPUT(K); CHECK_INPUT(V); CHECK_INPUT(dO);
+    CHECK_INPUT(L); CHECK_INPUT(delta);
+    CHECK_INPUT(dQ); CHECK_INPUT(dK); CHECK_INPUT(dV);
+
+    const int B = Q.size(0), N = Q.size(1), H = Q.size(2), H_KV = K.size(2);
+
+    auto stream = c10::xpu::getCurrentXPUStream(Q.device().index());
+    auto &queue = stream.queue();
+    queue.wait();
+
+    // Reorder scratch (bf16 P^T and dS^T), shaped [B, N(kv), H_KV, D(q=128)].
+    torch::Tensor Pt  = make_reorder_scratch(B, N, H_KV, Q);
+    torch::Tensor dSt = make_reorder_scratch(B, N, H_KV, Q);
+    torch::Tensor dSq = make_reorder_scratch(B, N, H_KV, Q);  // dQ q-major scratch
+
+    run_bwd(queue,
+            reinterpret_cast<bf16*>(Q.data_ptr<c10::BFloat16>()),
+            reinterpret_cast<bf16*>(K.data_ptr<c10::BFloat16>()),
+            reinterpret_cast<bf16*>(V.data_ptr<c10::BFloat16>()),
+            reinterpret_cast<bf16*>(dO.data_ptr<c10::BFloat16>()),
+            L.data_ptr<float>(), delta.data_ptr<float>(),
+            dQ.data_ptr<float>(), dK.data_ptr<float>(), dV.data_ptr<float>(),
+            reinterpret_cast<bf16*>(Pt.data_ptr<c10::BFloat16>()),
+            reinterpret_cast<bf16*>(dSt.data_ptr<c10::BFloat16>()),
+            reinterpret_cast<bf16*>(dSq.data_ptr<c10::BFloat16>()),
+            B, N, H, H_KV);
+}
+
 PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
     m.def("dispatch_prep", &dispatch_prep);
     m.def("dispatch_bwd", &dispatch_bwd);
 }
+
+#else
+// Standalone executable (CONFIG=standalone): main() lives in the harness.
+#include "gqa_backward_harness.impl"
+#endif  // TORCH_COMPILE