Docs: make ThreadFusion's opt-in semantics explicit (matches FKL default, per Oscar's guidance)

README.md

@@ -106,11 +106,15 @@ the GPU backend.
 ## ThreadFusion
 
 ```python
-pipe = fkl.compose(..., thread_fusion=True)   # TF::ENABLED, GPU only
+pipe = fkl.compose(..., thread_fusion=True)   # opt-in, GPU only
 ```
-Vectorized multi-element threads. Auto-falls back to scalar for shapes
-whose row bytes aren't 16-aligned (external tight-pitch pointers would
-fault on float4 loads).
+**Disabled by default**, matching FKL's own default (`TransformDPP<>` =
+`TF::DISABLED`). Per Oscar: ThreadFusion only improves performance in a
+small set of cases (wide images, trivial per-pixel chains, bandwidth-bound)
+— benchmark YOUR pipeline before enabling it; it is not a general speedup.
+When enabled it emits `TransformDPP<GPU_NVIDIA, TF::ENABLED>` and
+auto-falls back to scalar for shapes whose row bytes aren't 16-aligned
+(external tight-pitch pointers would fault on vectorized loads).
 
 ## Multi-GPU
 

fkl/jit.py

@@ -226,9 +226,12 @@ def source_for(self, dtype_spec, shape):
 
 def compose(*ops: Op, target: str = "gpu",
             thread_fusion: bool = False) -> FusedKernel:
-    """thread_fusion=True enables FKL's ThreadFusion (TF::ENABLED):
+    """thread_fusion=True opts into FKL's ThreadFusion (TF::ENABLED):
     each thread processes multiple elements with vectorized accesses.
-    GPU-only; best for wide images with simple per-pixel chains."""
+    DEFAULT IS FALSE, matching FKL's own TransformDPP<> default
+    (TF::DISABLED): per upstream guidance it only pays off in a small set
+    of cases (wide images, trivial chains, bandwidth-bound) — benchmark
+    your pipeline before enabling. GPU-only."""
     return FusedKernel(list(ops), target=target, thread_fusion=thread_fusion)
 
 

spikes/repro_circular_tensor_oob.cu

@@ -0,0 +1,44 @@
+// GROUND TRUTH: does fk::CircularTensor (used directly, as designed) work?
+// float3, COLOR_PLANES=1? No: aggregate path needs is_aggregate_v<T>.
+// Use the canonical config from the class design: T=float, COLOR_PLANES=3
+// (VectorType_t<float,3> = float3 storage), packed Standard.
+#include <fused_kernel/fused_kernel.h>
+#include <fused_kernel/core/execution_model/execution_model.h>
+#include <fused_kernel/core/data/circular_tensor.h>
+#include <fused_kernel/algorithms/algorithms.h>
+#include <cstdio>
+#include <vector>
+using namespace fk;
+
+int main() {
+    constexpr int W = 8, H = 4, B = 3;
+    using CT = CircularTensor<float, 3, B, CircularTensorOrder::NewestFirst, ColorPlanes::Standard>;
+    CT ct(W, H, MemType::Device);
+
+    Stream stream;
+    Ptr2D<float3> frame(W, H);
+
+    // push frames k=1..5, frame k = constant (k, k+0.1, k+0.2)
+    for (int k = 1; k <= 5; ++k) {
+        std::vector<float3> h(W * H, make_<float3>(float(k), k + 0.1f, k + 0.2f));
+        cudaMemcpy2D(frame.ptr().data, frame.ptr().dims.pitch, h.data(),
+                     W * sizeof(float3), W * sizeof(float3), H, cudaMemcpyHostToDevice);
+        ct.update(stream,
+                  PerThreadRead<ND::_2D, float3>::build(frame),
+                  TensorSplit<float3>::build(ct));
+        stream.sync();
+
+        // read back: Tensor base, planes layout = B batch x 3 color planes
+        std::vector<float> out(W * H * B * 3);
+        cudaMemcpy(out.data(), ct.ptr().data, out.size() * sizeof(float), cudaMemcpyDeviceToHost);
+        printf("after push %d: plane0.x=%5.1f plane1.x=%5.1f plane2.x=%5.1f\n",
+               k, out[0], out[W*H*3], out[2*W*H*3]);
+    }
+    // expectation (NewestFirst): planes = [5, 4, 3]
+    std::vector<float> out(W * H * B * 3);
+    cudaMemcpy(out.data(), ct.ptr().data, out.size() * sizeof(float), cudaMemcpyDeviceToHost);
+    const bool ok = out[0] == 5.f && out[W*H*3] == 4.f && out[2*W*H*3] == 3.f;
+    printf(ok ? "fk::CircularTensor DIRECT USE: CORRECT\n"
+              : "fk::CircularTensor DIRECT USE: WRONG ORDER/VALUES\n");
+    return ok ? 0 : 1;
+}