feat(chameleon): Chameleon-7B Thor FP8 and Orin SM87 INT8/QuaRot frontends

CMakeLists.txt

@@ -156,6 +156,13 @@ if((FLASHRT_ENABLE_COSMOS3_EDGE OR FLASHRT_ENABLE_COSMOS3_REASONER) AND
     "for non-Thor builds.")
 endif()
 
+# Chameleon-7B model kernels: QK Norm/RoPE, SM87 INT8/INT4 GEMM + FHT/QuaRot,
+# FA2 FP16 causal instances, and SM100/110 causal FMHA libraries. Off by
+# default so unrelated builds pay no compile/link/symbol cost. Model-neutral
+# fp16 norm/quant/activation helpers stay in the common layer.
+option(FLASHRT_ENABLE_CHAMELEON
+       "Build Chameleon-7B model kernels and bindings" OFF)
+
 # Motus beta integration. Motus-specific kernels are additive and must keep
 # their symbols prefixed with ``motus_``. Keeping a build tag lets the public
 # package compile without Motus kernels when debugging unrelated model paths.
@@ -483,6 +490,71 @@ if(ENABLE_SM100_CUTLASS)
   )
   target_link_libraries(fmha_fp16_strided PRIVATE CUDA::cudart)
   message(STATUS "libfmha_fp16_strided.so: building for sm_${GPU_ARCH} (Thor FMHA for SigLIP)")
+
+  if(FLASHRT_ENABLE_CHAMELEON)
+  # ── libfmha_fp16_causal.so — CUTLASS SM100 FP16 causal FMHA for Chameleon ──
+  # Same as fmha_fp16_strided but with CausalMask<true>/<false> and
+  # CausalIndividualTileScheduler. Used by Chameleon-7B LLM self-attention
+  # (is_causal=True). Loaded at runtime via ctypes in
+  # hardware/thor/attn_backend_chameleon.py.
+  add_library(fmha_fp16_causal SHARED csrc/attention/fmha_fp16_causal.cu)
+  set_target_properties(fmha_fp16_causal PROPERTIES
+    LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/flash_rt
+    RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/flash_rt
+    CUDA_STANDARD 17
+    POSITION_INDEPENDENT_CODE ON
+    CUDA_ARCHITECTURES "${GPU_ARCH}a"
+    CUDA_RESOLVE_DEVICE_SYMBOLS ON
+    PREFIX "lib"
+    OUTPUT_NAME "fmha_fp16_causal"
+  )
+  target_include_directories(fmha_fp16_causal PRIVATE
+    ${CMAKE_CURRENT_SOURCE_DIR}/csrc/attention
+    ${CUTLASS_DIR}/examples/77_blackwell_fmha
+    ${CUTLASS_INCLUDE}
+    ${CUTLASS_DIR}/tools/util/include
+  )
+  target_compile_options(fmha_fp16_causal PRIVATE
+    $<$<COMPILE_LANGUAGE:CUDA>:
+      --expt-relaxed-constexpr --expt-extended-lambda -O3
+      ${GPU_GENCODE}
+    >
+  )
+  target_link_libraries(fmha_fp16_causal PRIVATE CUDA::cudart)
+  message(STATUS "libfmha_fp16_causal.so: building for sm_${GPU_ARCH} (Thor causal FMHA for Chameleon)")
+
+  # ── libfmha_fp8_causal.so — CUTLASS SM100 FP8 causal FMHA ──
+  # FP8 (E4M3) input variant of fmha_fp16_causal for Chameleon-7B. Inputs
+  # Q/K/V are FP8, outputs O FP16, accumulators FP32. CUTLASS Sm100 FMHA
+  # mainloop has FP8-aware kPRescale logic that triggers when
+  # ``Element == cutlass::float_e4m3_t``. Loaded via dlopen at runtime
+  # alongside libfmha_fp16_causal.so as a drop-in alternative.
+  add_library(fmha_fp8_causal SHARED csrc/attention/fmha_fp8_causal.cu)
+  set_target_properties(fmha_fp8_causal PROPERTIES
+    LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/flash_rt
+    RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/flash_rt
+    CUDA_STANDARD 17
+    POSITION_INDEPENDENT_CODE ON
+    CUDA_ARCHITECTURES "${GPU_ARCH}a"
+    CUDA_RESOLVE_DEVICE_SYMBOLS ON
+    PREFIX "lib"
+    OUTPUT_NAME "fmha_fp8_causal"
+  )
+  target_include_directories(fmha_fp8_causal PRIVATE
+    ${CMAKE_CURRENT_SOURCE_DIR}/csrc/attention
+    ${CUTLASS_DIR}/examples/77_blackwell_fmha
+    ${CUTLASS_INCLUDE}
+    ${CUTLASS_DIR}/tools/util/include
+  )
+  target_compile_options(fmha_fp8_causal PRIVATE
+    $<$<COMPILE_LANGUAGE:CUDA>:
+      --expt-relaxed-constexpr --expt-extended-lambda -O3
+      ${GPU_GENCODE}
+    >
+  )
+  target_link_libraries(fmha_fp8_causal PRIVATE CUDA::cudart)
+  message(STATUS "libfmha_fp8_causal.so: building for sm_${GPU_ARCH} (Thor FP8 causal FMHA for Chameleon)")
+  endif()  # FLASHRT_ENABLE_CHAMELEON
 endif()
 
 # ── CUTLASS SM120a block-128 FP8 GEMM (Path B for Qwen3.6) ──
@@ -869,11 +941,18 @@ if(ENABLE_FA2 AND
       csrc/attention/fa2_causal_inst/flash_fwd_split_hdim256_bf16_sm80_causal.cu
     )
   endif()
+  # fp16 hdim=128 serves Chameleon-7B causal prefill/decode on Orin SM87.
+  if(FLASHRT_ENABLE_CHAMELEON AND "128" IN_LIST FA2_HDIMS AND "fp16" IN_LIST FA2_DTYPES)
+    list(APPEND FA2_SRCS
+      csrc/attention/fa2_causal_inst/flash_fwd_hdim128_fp16_sm80_causal.cu
+      csrc/attention/fa2_causal_inst/flash_fwd_split_hdim128_fp16_sm80_causal.cu
+    )
+  endif()
   if(FLASHRT_ENABLE_NATIVE_CPP)
     # The native C boundary has a stable five-symbol surface. Its causal
     # wrapper contains native-only fail-fast dispatch for a slim matrix.
     list(APPEND FA2_SRCS csrc/attention/fa2_wrapper_causal.cu)
-  elseif("bf16" IN_LIST FA2_DTYPES AND
+  elseif(("bf16" IN_LIST FA2_DTYPES OR "fp16" IN_LIST FA2_DTYPES) AND
          ("128" IN_LIST FA2_HDIMS OR "256" IN_LIST FA2_HDIMS))
     # Preserve the existing Python-only source matrix exactly by default.
     list(APPEND FA2_SRCS csrc/attention/fa2_wrapper_causal.cu)
@@ -1582,6 +1661,26 @@ if(ENABLE_SM80_INT8_CUTLASS)
   target_compile_definitions(flash_rt_kernels PRIVATE ENABLE_SM80_INT8_CUTLASS=1)
 endif()
 
+# ── Chameleon-7B model kernels (opt-in, see FLASHRT_ENABLE_CHAMELEON) ──
+# QK Norm/RoPE fused, AWQ FP16 quant, SM80 INT8/INT4 rowwise GEMM fp16-out +
+# FHT/QuaRot rotation. Everything here is Chameleon-specific; model-neutral
+# fp16 norm/quant/activation helpers live in the common layer.
+if(FLASHRT_ENABLE_CHAMELEON)
+  target_sources(flash_rt_kernels PRIVATE
+    csrc/kernels/qk_norm_rope_fused.cu
+    csrc/quantize/awq_quant_fp8_static_fp16.cu)
+  if(ENABLE_SM80_INT8_CUTLASS)
+    target_sources(flash_rt_kernels PRIVATE
+      csrc/gemm/cutlass_sm80_int8_rowwise_fp16out.cu
+      csrc/gemm/cutlass_sm80_int8_rowwise_fp16out_t64x128.cu
+      csrc/gemm/cutlass_sm80_int8_rowwise_fp16out_t256x128.cu
+      csrc/gemm/cutlass_sm80_int4_rowwise.cu
+      csrc/kernels/fht_int4.cu)
+  endif()
+  target_compile_definitions(flash_rt_kernels PRIVATE FLASHRT_ENABLE_CHAMELEON=1)
+  message(STATUS "Chameleon-7B model kernels: ENABLED")
+endif()
+
 # SM120a CUTLASS block-128 FP8 GEMM (Path B for Qwen3.6).
 if(GPU_ARCH STREQUAL "120")
   target_sources(flash_rt_kernels PRIVATE

README.md

@@ -1072,6 +1072,7 @@ examples/
 - **Cosmos3-Nano text-to-video** (`config="cosmos3_video"`) — RTX 5090 BF16/FP8 denoise and complete benchmark workflow; [usage and performance](docs/cosmos3_video_usage.md)
 - **Cosmos3-Edge AV inverse dynamics and Reasoner** (`config="cosmos3_edge"`) — Jetson AGX Thor official baseline, 6.60x no-cache AV denoise, and NVFP4 multimodal chat decode; [complete usage and performance](docs/cosmos3_edge_thor.md)
 - **Qwen3-VL-8B** — RTX 5090 NVFP4/FP8 multimodal path, RTX 4090 official-FP8 path, and Jetson BF16 paths for Thor and Orin; [RTX 5090 usage](docs/qwen3_vl_nvfp4.md), [RTX 4090 usage](docs/qwen3_vl_fp8_sm89.md), [Jetson Thor usage](docs/qwen3_vl_thor.md), [Jetson Orin usage](docs/qwen3_vl_rtx_bf16.md)
+- **Chameleon-7B** — Jetson Thor dynamic-FP8 prefill (~120 ms E2E, ~30 tok/s decode) and Jetson Orin INT8/QuaRot-INT4 path; [usage](docs/chameleon_usage.md), [Thor notes](docs/chameleon_thor_sm110.md), [Orin SM87 notes](docs/chameleon7b_rtx_sm87.md)
 - **MiniMax-Remover** — FP8 transformer + NVFP4 VAE video inpainting; [usage and performance](docs/minimax_remover_usage.md)
 - **MelBandRoformer** — kernelized FP8 audio source separation; [usage and performance](docs/melband_roformer_usage.md)
 - **OmniVoice TTS** — BF16/FP4 acceleration and HTTP serving; [serving quickstart](serving/omnivoice_agent/README.md)

USAGE.md

@@ -681,6 +681,39 @@ once and reused, so per-chunk cost is just the infer row.
 This path does not change CMake targets, C++ bindings, or existing
 Pi0/Pi0.5/GROOT N1.6 runtime dispatch.
 
+### Chameleon-7B (Thor)
+
+Standalone Chameleon-7B (text + image) is a direct-instantiation Thor
+frontend — it is registered in `_PIPELINE_MAP` but is **not** dispatched by
+`flash_rt.load_model` (same pattern as Qwen3-VL). The VQGAN image
+tokenizer defaults to the generic eager Chameleon path; if compatible
+TensorRT engines exist in the deployment, it is recommended to opt in
+explicitly (`use_trt_vqgan=True`). A Jetson Orin (SM87) INT8/QuaRot
+frontend is also available; see [`docs/chameleon_usage.md`](docs/chameleon_usage.md),
+[`docs/chameleon_thor_sm110.md`](docs/chameleon_thor_sm110.md) and
+[`docs/chameleon7b_rtx_sm87.md`](docs/chameleon7b_rtx_sm87.md).
+
+```python
+from flash_rt.frontends.torch.chameleon_thor import ChameleonTorchFrontendThor
+
+fe = ChameleonTorchFrontendThor(
+    "/path/to/Chameleon_7B_mGPT",
+    use_fp8=True,            # dynamic per-tensor FP8 (recommended default)
+    use_cuda_graph=True,
+    use_trt_vqgan=False,     # generic default = eager VQGAN; set True when engines exist
+    target_size=512,
+)
+
+out = fe.prefill("Describe the image.", [pil_image])
+# out["logits"]: (65536,) fp32 with mask_image_logits applied
+```
+
+FA4 attention is an explicit opt-in fast path
+(`use_fa4_attn=True` / `FLASHRT_CHAMELEON_FA4_ATTN=1`, needs the
+`thor-fa4` pip extra); measured transformer-prefill-only ≈ **104 ms** at
+Se≈1056 vs ≈111 ms with CUTLASS FMHA. E2E (TRT VQGAN + FA4) ≈ **120 ms**
+vs HF BF16 ≈ 403 ms transformer-only (~3.4×).
+
 ### Wan2.2 TI2V-5B
 
 Wan2.2 TI2V-5B is exposed as an RTX SM120 official-pipeline baseline:

benchmarks/chameleon_thor_latency.py

@@ -0,0 +1,181 @@
+#!/usr/bin/env python3
+"""Chameleon-7B (Thor sm_110) latency benchmark.
+
+Measures standalone Chameleon-7B prefill latency on real images with clean
+stage separation:
+
+- ``--reuse-input-ids``: build real-image input ids once, then time only
+  embed + backbone + lm_head (transformer-prefill-only, HF-comparable).
+- Default: full ``prefill()`` E2E including VQGAN tokenization.
+- ``--use-trt-vqgan``: explicit TensorRT VQGAN opt-in (recommended when
+  compatible engines exist; the generic default stays eager).
+- FA4 attention: enable with ``FLASHRT_CHAMELEON_FA4_ATTN=1`` (needs the
+  ``thor-fa4`` pip extra; prints whether it is active).
+
+Latency is wall-clock P50 (per CONTRIBUTING.md: quickstart --benchmark
+style; CUDA-graph replayed latency is what the pipeline measures inside the
+graph). Every result row records device, VQGAN backend, FA4 state, Se,
+fp8/fp16 and graph settings for reproducible reporting.
+
+Usage:
+
+    python benchmarks/chameleon_thor_latency.py \
+        --checkpoint /path/to/Chameleon_7B_mGPT \
+        --image-dir /path/to/images \
+        --iters 20 --warmup 5
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import pathlib
+import time
+from typing import Dict, List
+
+import torch
+
+
+def _stats(xs: List[float]) -> Dict[str, float]:
+    a = sorted(xs)
+    n = len(a)
+    return {"mean": sum(a) / n, "p50": a[n // 2], "min": a[0], "max": a[-1]}
+
+
+def _load_images(image_dir: pathlib.Path, max_images: int):
+    from PIL import Image
+
+    paths = sorted(
+        p for p in image_dir.iterdir()
+        if p.suffix.lower() in (".jpg", ".jpeg", ".png", ".bmp"))
+    if not paths:
+        raise FileNotFoundError(f"No real images under {image_dir}")
+    paths = paths[:max_images]
+    return [Image.open(p).convert("RGB") for p in paths], [str(p) for p in paths]
+
+
+def _pad_ids(input_ids: List[int], pad_id: int = 1):
+    real_len = len(input_ids)
+    padded = list(input_ids)
+    rem = len(padded) % 16
+    if rem:
+        padded.extend([pad_id] * (16 - rem))
+    return padded, real_len
+
+
+def _prefill_once(fe, prompt, images, cached_ids, use_graph: bool) -> Dict[str, float]:
+    """One timed prefill; returns stage latencies in ms."""
+    times: Dict[str, float] = {}
+    t0 = time.perf_counter()
+    ids = fe.encode_prompt(prompt, images) if cached_ids is None else cached_ids
+    torch.cuda.synchronize()
+    times["encode_ms"] = (time.perf_counter() - t0) * 1000.0
+
+    padded, real_len = _pad_ids(ids)
+    fe._real_len = real_len
+    fe.Se = len(padded)
+    fe._last_input_ids = padded
+    if fe._use_autotune:
+        fe._autotune_gemms(fe.Se)
+    torch.cuda.synchronize()
+    t1 = time.perf_counter()
+
+    fe._embed_ids(padded)
+    torch.cuda.synchronize()
+    t2 = time.perf_counter()
+
+    if use_graph:
+        fe._capture_graph(fe.Se)
+        fe._infer_graph.replay()
+    else:
+        fe._run_backbone(fe.Se)
+    torch.cuda.synchronize()
+    t3 = time.perf_counter()
+
+    fe._project_last()
+    torch.cuda.synchronize()
+    t4 = time.perf_counter()
+
+    times["prepare_ms"] = (t1 - t0) * 1000.0 - times["encode_ms"]
+    times["embed_ms"] = (t2 - t1) * 1000.0
+    times["backbone_ms"] = (t3 - t2) * 1000.0
+    times["lm_head_ms"] = (t4 - t3) * 1000.0
+    times["transformer_ms"] = times["embed_ms"] + times["backbone_ms"] + times["lm_head_ms"]
+    times["total_ms"] = (t4 - t0) * 1000.0
+    return times
+
+
+def main() -> None:
+    ap = argparse.ArgumentParser(description=__doc__)
+    ap.add_argument("--checkpoint", required=True)
+    ap.add_argument("--image-dir", required=True,
+                    help="Directory of real input images")
+    ap.add_argument("--prompt", default="Describe the image.")
+    ap.add_argument("--max-images", type=int, default=1)
+    ap.add_argument("--target-size", type=int, default=512)
+    ap.add_argument("--use-trt-vqgan", action="store_true",
+                    help="Use TensorRT VQGAN if compatible engines exist "
+                         "(recommended when engines are available; default is eager VQGAN)")
+    ap.add_argument("--reuse-input-ids", action="store_true",
+                    help="time transformer prefill only (VQGAN excluded)")
+    ap.add_argument("--no-graph", action="store_true")
+    ap.add_argument("--use-fp16", action="store_true",
+                    help="FP16 reference path instead of dynamic FP8")
+    ap.add_argument("--iters", type=int, default=20)
+    ap.add_argument("--warmup", type=int, default=5)
+    ap.add_argument("--output", default=None, help="JSON output path")
+    args = ap.parse_args()
+
+    from flash_rt.frontends.torch.chameleon_thor import ChameleonTorchFrontendThor
+    from flash_rt.hardware.thor import fa4_backend
+
+    images, image_paths = _load_images(pathlib.Path(args.image_dir), args.max_images)
+    use_graph = not args.no_graph
+    fp8 = not args.use_fp16
+
+    fe = ChameleonTorchFrontendThor(
+        args.checkpoint, use_fp8=fp8, use_cuda_graph=use_graph,
+        target_size=args.target_size, use_trt_vqgan=args.use_trt_vqgan)
+
+    cached_ids = None
+    if args.reuse_input_ids:
+        cached_ids = fe.encode_prompt(args.prompt, images)
+        torch.cuda.synchronize()
+
+    for _ in range(args.warmup):
+        _prefill_once(fe, args.prompt, images, cached_ids, use_graph)
+
+    rows: Dict[str, List[float]] = {}
+    for _ in range(args.iters):
+        t = _prefill_once(fe, args.prompt, images, cached_ids, use_graph)
+        for k, v in t.items():
+            rows.setdefault(k, []).append(v)
+
+    result = {
+        "model": "chameleon-7b",
+        "device": torch.cuda.get_device_name(0),
+        "sm_count": torch.cuda.get_device_properties(0).multi_processor_count,
+        "checkpoint": args.checkpoint,
+        "image_paths": image_paths,
+        "prompt": args.prompt,
+        "target_size": args.target_size,
+        "fp8": fp8,
+        "cuda_graph": use_graph,
+        "vqgan_backend": fe.vqgan_backend,
+        "fa4_attn": fe.fa4_attn_active,
+        "fa4_status": fa4_backend.status(),
+        "reuse_input_ids": bool(args.reuse_input_ids),
+        "Se": int(fe.Se),
+        "latency_ms": {k: _stats(v) for k, v in rows.items()},
+    }
+    for k, v in result["latency_ms"].items():
+        print(f"[chameleon] {k:14s} p50={v['p50']:8.1f} ms  mean={v['mean']:8.1f}")
+    print(f"[chameleon] device={result['device']} vqgan={fe.vqgan_backend} "
+          f"fa4={fe.fa4_attn_active} fp8={fp8} graph={use_graph} Se={fe.Se}")
+
+    if args.output:
+        pathlib.Path(args.output).write_text(json.dumps(result, indent=2))
+        print(f"[chameleon] wrote {args.output}")
+
+
+if __name__ == "__main__":
+    main()