From 06a23190e7ce7e1de30e463d132fe0617ba268d4 Mon Sep 17 00:00:00 2001 From: johnnynunez Date: Thu, 11 Jun 2026 17:23:19 +0200 Subject: [PATCH 1/2] ThreadFusion docs: make the opt-in semantics explicit (Oscar's guidance) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit thread_fusion was ALWAYS opt-in (default False, matching FKL's own TransformDPP<> = TF::DISABLED default), but the docstring/README sold it too enthusiastically. Now both state explicitly: disabled by default, matching upstream, because per Oscar it only improves performance in a small set of cases — benchmark your pipeline before enabling. --- README.md | 12 ++++++++---- fkl/jit.py | 7 +++++-- 2 files changed, 13 insertions(+), 6 deletions(-) diff --git a/README.md b/README.md index 0e06298..8daae25 100644 --- a/README.md +++ b/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` 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 diff --git a/fkl/jit.py b/fkl/jit.py index 2ae2711..59e9583 100644 --- a/fkl/jit.py +++ b/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) From bdd9ab3b0444ccca5f85b72045b2dd11d8d9228e Mon Sep 17 00:00:00 2001 From: johnnynunez Date: Thu, 11 Jun 2026 17:33:17 +0200 Subject: [PATCH 2/2] spikes: add CircularTensor direct-use OOB probe (issue #250 evidence) C++-only probe using fk::CircularTensor exactly as designed (float, COLOR_PLANES=3, NewestFirst, TensorSplit write). Values come out CORRECT (rotation 5/4/3) but compute-sanitizer reports 33 invalid reads/writes per run on unfixed headers: update's 1-plane read + copy's BATCH-plane read sum to BATCH+1 launched planes. Sanitizer-clean with FKL PR #252 applied. This is the evidence posted on FusedKernelLibrary#250. --- spikes/repro_circular_tensor_oob.cu | 44 +++++++++++++++++++++++++++++ 1 file changed, 44 insertions(+) create mode 100644 spikes/repro_circular_tensor_oob.cu diff --git a/spikes/repro_circular_tensor_oob.cu b/spikes/repro_circular_tensor_oob.cu new file mode 100644 index 0000000..c6bd47b --- /dev/null +++ b/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. +// Use the canonical config from the class design: T=float, COLOR_PLANES=3 +// (VectorType_t = float3 storage), packed Standard. +#include +#include +#include +#include +#include +#include +using namespace fk; + +int main() { + constexpr int W = 8, H = 4, B = 3; + using CT = CircularTensor; + CT ct(W, H, MemType::Device); + + Stream stream; + Ptr2D 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 h(W * H, make_(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::build(frame), + TensorSplit::build(ct)); + stream.sync(); + + // read back: Tensor base, planes layout = B batch x 3 color planes + std::vector 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 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; +}