From a29b8960b0f887c20e74dafa16a24c37d6508b4e Mon Sep 17 00:00:00 2001 From: Daniel Maddern Date: Wed, 26 Aug 2026 00:48:44 +0700 Subject: [PATCH] Integrate guarded FC2 cuBLASLt schedule --- CURRENT_STATE.md | 14 +- Dockerfile.spark | 2 +- PERFORMANCE_ROADMAP.md | 12 +- ...-nvfp4-production-block-gate-20260825.json | 2025 ++++++++++++++ ...production-trajectory-12step-20260825.json | 1776 ++++++++++++ ...-production-trajectory-2step-20260825.json | 1776 ++++++++++++ .../gb10-fc2-nvfp4-shape-gate-20260825.json | 2404 +++++++++++++++++ compose.spark-stack.yml | 1 + compose.spark.yml | 1 + pyproject.toml | 3 + research/artifact_manifest.json | 94 +- research/experiment_registry.json | 22 +- research/fc2_nvfp4_scheduling/README.md | 15 +- research/fc2_nvfp4_scheduling/RESULTS.md | 58 +- .../csrc}/fc2_nvfp4_lt.cpp | 42 +- src/h3_blackwell_runtime/fc2_lt.py | 197 ++ src/h3_blackwell_runtime/nvfp4.py | 11 +- src/h3_blackwell_runtime/nvfp4_quant.py | 9 +- src/h3_blackwell_runtime/runtime.py | 6 + tests/test_fc2_lt.py | 287 ++ tools/benchmark_fc2_nvfp4_algorithms.py | 184 +- 21 files changed, 8879 insertions(+), 60 deletions(-) create mode 100644 benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json create mode 100644 benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json create mode 100644 benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json create mode 100644 benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json rename {research/fc2_nvfp4_scheduling => src/h3_blackwell_runtime/csrc}/fc2_nvfp4_lt.cpp (93%) create mode 100644 src/h3_blackwell_runtime/fc2_lt.py create mode 100644 tests/test_fc2_lt.py diff --git a/CURRENT_STATE.md b/CURRENT_STATE.md index 77c9678..a15c8f5 100644 --- a/CURRENT_STATE.md +++ b/CURRENT_STATE.md @@ -160,7 +160,8 @@ video and audio tensors. Spark enables the path with `SAGE2_BLACKWELL_DESIGN.md`, and `benchmarks/gb10-post-optimization-profile-summary.json`. -The isolated FC2 cuBLASLt scheduling study is complete. The production +The FC2 cuBLASLt scheduling study and guarded canonical-shape integration are +complete. The production heuristic's `_stream_k` kernel requests the same `25.664 GB` of operands as the retained public split-K-1 schedule, but its L2 hit rate is only `53.32%` versus `91.10%`; it incurs `9.853 GB` more L2 read misses and spends heavily in @@ -168,8 +169,15 @@ synchronization polling. Algorithm 70, tile 20, stages 37, split-K 1 is byte-exact with zero workspace. It improves complete blocks 0, 24, and 49 by `8.16-8.88%`, the two-step trajectory by `7.50%`, and the canonical 12-step trajectory from `278.201 s` to `255.371 s` (`8.21%`) with exact video and audio -latents. This remains a research-retained integration candidate: production -dispatch and configuration are unchanged. See +latents. The final hardened production method improves 20-round blocks by +`7.86-9.36%` and the canonical 12-step trajectory from `286.431 s` to +`262.979 s` (`8.19%`), with 600 successful dispatches, zero fallback, and exact +latents. `H3_NVFP4_FC2_LT_SPLITK1=1` enables only the validated `M=37,810` +descriptor; nearby row counts can differ by two BF16 elements and therefore +retain the existing Comfy fallback. The extension and measured runtime ABI are +prepared during H3 model loading rather than on the first canonical request. +The Spark hot service is intentionally +stopped until performance work is complete. See `research/fc2_nvfp4_scheduling/RESULTS.md`. The Spark hot runtime was rebuilt and recreated with image diff --git a/Dockerfile.spark b/Dockerfile.spark index 904191c..685efe5 100644 --- a/Dockerfile.spark +++ b/Dockerfile.spark @@ -33,7 +33,7 @@ RUN python -m pip install --no-cache-dir --no-deps -e . \ ENV TORCH_EXTENSIONS_DIR=/opt/h3-blackwell-runtime/.torch_extensions RUN TORCH_CUDA_ARCH_LIST=12.1 python -c \ - "from h3_blackwell_runtime.nvfp4_quant import _vortex_scale_extension; print(_vortex_scale_extension().__file__)" + "from h3_blackwell_runtime.nvfp4_quant import _vortex_scale_extension; from h3_blackwell_runtime.fc2_lt import _fc2_lt_extension; print(_vortex_scale_extension().__file__, _fc2_lt_extension().__file__)" ENV H3_MODEL_PATH=/models/minimax_h3_ref2va_pruned_nvfp4.safetensors ENV PYTHONPATH=/opt/ComfyUI-sol-attn diff --git a/PERFORMANCE_ROADMAP.md b/PERFORMANCE_ROADMAP.md index f29de8c..5afd8d4 100644 --- a/PERFORMANCE_ROADMAP.md +++ b/PERFORMANCE_ROADMAP.md @@ -394,12 +394,18 @@ tensor-pipe activity from `25.01%` to `82.42%`. There is no material global partial-accumulator or output-reduction traffic; the baseline instead spends heavily in Stream-K synchronization polling and loses traversal locality. -The candidate is byte-exact and improves FC2 p50 from `52.521 ms` to +The candidate is byte-exact and improves FC2 p50 from `53.618 ms` to `15.636 ms`. Complete blocks 0, 24, and 49 improve by `8.16-8.88%`; two-step and canonical 12-step trajectories improve by `7.50%` and `8.21%`, with exact video and audio latents. The custom persistent-kernel branch is therefore -closed. Production integration remains separate work, so current dispatch is -unchanged. See `research/fc2_nvfp4_scheduling/RESULTS.md` and +closed. Guarded production integration is complete behind +`H3_NVFP4_FC2_LT_SPLITK1=1`. Its direct 20-round block gate improves blocks +0/24/49 by `9.36%`, `8.16%`, and `7.86%`; the exact canonical 12-step gate +improves by `8.19%` +with 600 successful dispatches and zero fallback. Shape testing found two-BF16 +differences at several nearby row counts, so only `M=37,810` is eligible and +all variable shapes retain Comfy. The hot service remains stopped until the +performance program is complete. See `research/fc2_nvfp4_scheduling/RESULTS.md` and `benchmarks/gb10-fc2-nvfp4-library-sweep-20260825.json`. The follow-up real block-24 Sage2 decomposition now selects the next exact diff --git a/benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json b/benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json new file mode 100644 index 0000000..44c49f9 --- /dev/null +++ b/benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json @@ -0,0 +1,2025 @@ +{ + "mode": "production-block-gate", + "environment": { + "platform": "Linux-6.17.0-1026-nvidia-aarch64-with-glibc2.39", + "python": "3.12.3 (main, Mar 23 2026, 19:04:32) [GCC 13.3.0]", + "torch": "2.9.1+cu130", + "torch_cuda": "13.0", + "device": "NVIDIA GB10", + "device_capability": [ + 12, + 1 + ], + "driver": null, + "git_commit": null, + "checkpoint_path": "/models/minimax_h3_fl2va_pruned_nvfp4.safetensors", + "checkpoint_sha256": null, + "checkpoint_hash_note": "not calculated", + "environment_switches": { + "CUDA_DEVICE_MAX_CONNECTIONS": "1", + "CUDA_DEVICE_MAX_COPY_CONNECTIONS": "4", + "CUDA_HOME": "/usr/local/cuda", + "CUDA_INC_PATH": "/usr/local/cuda/include", + "CUDA_MANAGED_FORCE_DEVICE_ALLOC": "1", + "CUDA_MODULE_LOADING": "EAGER", + "CUDA_VERSION": "13.0.2", + "H3_FUSED_ELEMENTWISE": "1", + "H3_MODEL_PATH": "/models/minimax_h3_ref2va_pruned_nvfp4.safetensors", + "H3_NVFP4_MODULATE_FUSION": "1", + "H3_NVFP4_SCALE_BACKEND": "vortex", + "H3_NVFP4_SCALE_VERSION": "1", + "H3_NVFP4_SWIGLU_FUSION": "1", + "H3_SAGE_QKV_LAYOUT": "strided_nhd", + "TORCH_COMPILE_DISABLE": "0", + "TORCH_CUDA_ARCH_LIST": "12.1a", + "TORCH_EXTENSIONS_DIR": "/opt/h3-blackwell-runtime/.torch_extensions" + }, + "extension": { + "cuda_version": 13000, + "cublas_version": 130100, + "cuda_runtime_version": 13000, + "cublaslt_runtime_version": 130000, + "stream_k_public_control": false, + "stream_k_note": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + }, + "comfy_kitchen": "0.2.31 package without __version__" + }, + "workload": { + "resolution": [ + 1344, + 768 + ], + "frames": 124, + "steps": 12, + "sampler_step": 1, + "seed": 440420, + "text_tokens": 100, + "tokens": 37810, + "hidden_shape": [ + 37810, + 5376 + ], + "segments": [ + [ + 0, + 100, + 1 + ], + [ + 100, + 514, + 2 + ], + [ + 514, + 37810, + 0 + ] + ] + }, + "retained_blocks": [ + 0, + 24, + 49 + ], + "immutable_cloned_block_inputs": { + "0": [ + 37810, + 5376 + ], + "24": [ + 37810, + 5376 + ], + "49": [ + 37810, + 5376 + ] + }, + "fc2_boundary": { + "block": 24, + "gate_up_shape": [ + 37810, + 28672 + ], + "activation_qdata_shape": [ + 37824, + 7168 + ], + "weight_qdata_shape": [ + 5376, + 7168 + ], + "logical_mnk": [ + 37810, + 5376, + 14336 + ], + "descriptor_mnk_after_padding": [ + 37824, + 5376, + 14336 + ], + "producer": "vortex_native_quantize_swiglu_nvfp4", + "no_bias": true + }, + "baseline_kernel_metadata": { + "path": "fc2.forward_swiglu -> accepted producer -> Comfy Kitchen 0.2.31 scaled_mm_nvfp4", + "descriptors": { + "packed_input_output": "row-major [M,K] @ [N,K].T -> BF16 [M,N]", + "block_scale_mode": "VEC16_UE4M3", + "compute_and_scale": "FP32", + "scalar_pointer_mode": "device", + "bias": null, + "beta": 0.0, + "comfy_kitchen_version": "0.2.31" + }, + "profiler_cuda_events_available": false, + "profiler_note": "Torch profiler returned no CUDA kernel events on this build; use --mode profile with NCU for kernel metadata.", + "top_cuda_events": [], + "output_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f" + }, + "heuristics": [ + { + "max_workspace_bytes": 0, + "requested_count": 32, + "returned_count": 5, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 4194304, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 8388608, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 16777216, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 33554432, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 67108864, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + } + ], + "explicit_split_k_checks": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 1, + "reduction_scheme": 0 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 4 + } + } + ], + "selected": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0 + }, + "production_block_gate": [ + { + "block": 0, + "production_method": "block.mlp.fc2.forward_swiglu", + "candidate_vs_baseline": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147", + "expected_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147" + }, + "baseline_vs_traversal": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147", + "expected_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147" + }, + "candidate_vs_traversal": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147", + "expected_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 462.0343322753906, + 519.793701171875, + 463.6451721191406, + 461.61883544921875, + 461.3515930175781, + 464.995361328125, + 463.2869873046875, + 462.8765563964844, + 461.7970886230469, + 465.5448913574219, + 461.8421936035156, + 465.0537109375, + 464.8384094238281, + 465.4224548339844, + 463.8840637207031, + 463.4743347167969, + 465.61236572265625, + 464.4220275878906, + 464.6580810546875, + 466.87091064453125 + ], + "p50_ms": 464.1530456542969, + "p95_ms": 469.51705017089853, + "mean_ms": 466.6511535644531 + }, + "candidate": { + "samples_ms": [ + 571.5564575195312, + 422.9067077636719, + 416.4256591796875, + 419.8518981933594, + 419.189453125, + 420.7510070800781, + 420.09149169921875, + 419.6890869140625, + 419.6440734863281, + 420.3362731933594, + 421.8302917480469, + 417.80389404296875, + 420.0198669433594, + 420.9107971191406, + 421.89117431640625, + 420.67724609375, + 421.85693359375, + 422.25421142578125, + 422.38531494140625, + 423.3605041503906 + ], + "p50_ms": 420.71412658691406, + "p95_ms": 430.77030181884777, + "mean_ms": 428.17161712646487 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147", + "expected_sha256": "6bdb03e2bdf7733fde92b2cf5c78a373e9818cb5ec9d213d873a632fad0af147" + } + }, + "dispatch_delta": { + "attempts": 23, + "successes": 23, + "fallbacks": 0 + } + }, + { + "block": 24, + "production_method": "block.mlp.fc2.forward_swiglu", + "candidate_vs_baseline": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75", + "expected_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75" + }, + "baseline_vs_traversal": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75", + "expected_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75" + }, + "candidate_vs_traversal": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75", + "expected_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 462.13031005859375, + 1074.8060302734375, + 499.2489318847656, + 460.1708984375, + 459.4143371582031, + 459.58917236328125, + 458.5694274902344, + 459.37249755859375, + 457.447509765625, + 462.1033020019531, + 460.47076416015625, + 458.94989013671875, + 461.4758605957031, + 461.214111328125, + 458.8883972167969, + 459.40771484375, + 461.8358154296875, + 465.3328857421875, + 457.39892578125, + 460.1976623535156 + ], + "p50_ms": 460.1842803955078, + "p95_ms": 528.0267868041996, + "mean_ms": 492.9012222290039 + }, + "candidate": { + "samples_ms": [ + 427.8656921386719, + 427.14617919921875, + 525.4254150390625, + 420.90576171875, + 421.28857421875, + 422.8717041015625, + 423.4268493652344, + 422.6097106933594, + 422.38580322265625, + 422.58343505859375, + 423.83770751953125, + 422.4466552734375, + 422.70172119140625, + 422.8758544921875, + 421.4642333984375, + 422.4757995605469, + 424.78662109375, + 422.1180725097656, + 421.7055358886719, + 424.3546447753906 + ], + "p50_ms": 422.6557159423828, + "p95_ms": 432.7436782836915, + "mean_ms": 428.26379852294923 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75", + "expected_sha256": "d8f5c6335e823dc4e3ed8cdd34d90dda0d69297ea47de9b328962dfae00c9f75" + } + }, + "dispatch_delta": { + "attempts": 23, + "successes": 23, + "fallbacks": 0 + } + }, + { + "block": 49, + "production_method": "block.mlp.fc2.forward_swiglu", + "candidate_vs_baseline": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6", + "expected_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6" + }, + "baseline_vs_traversal": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6", + "expected_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6" + }, + "candidate_vs_traversal": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6", + "expected_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 457.59906005859375, + 455.51739501953125, + 454.50775146484375, + 455.77972412109375, + 454.094482421875, + 453.72442626953125, + 456.2005615234375, + 457.22589111328125, + 455.93505859375, + 454.9208679199219, + 454.7513732910156, + 456.42205810546875, + 456.1644592285156, + 455.8465270996094, + 459.75469970703125, + 456.6866149902344, + 456.8373107910156, + 458.20733642578125, + 456.2602233886719, + 456.6114807128906 + ], + "p50_ms": 456.18251037597656, + "p95_ms": 458.2847045898437, + "mean_ms": 456.15236511230466 + }, + "candidate": { + "samples_ms": [ + 421.10394287109375, + 427.0021057128906, + 418.5732116699219, + 420.71209716796875, + 419.9263916015625, + 420.857177734375, + 417.8297424316406, + 421.5742492675781, + 421.7043762207031, + 419.58203125, + 419.05865478515625, + 421.8150634765625, + 421.2725524902344, + 419.3066711425781, + 419.87371826171875, + 420.12255859375, + 419.9741516113281, + 420.2537841796875, + 420.41552734375, + 422.7670593261719 + ], + "p50_ms": 420.33465576171875, + "p95_ms": 422.9788116455078, + "mean_ms": 420.6862533569336 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6", + "expected_sha256": "3723ebf82eb1e279cc004832b68d071422a45dcc74adac93010e5fde33b715d6" + } + }, + "dispatch_delta": { + "attempts": 23, + "successes": 23, + "fallbacks": 0 + } + } + ], + "errors_and_unsupported": [ + { + "feature": "Stream-K", + "supported_public_control": false, + "reason": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + } + ] +} diff --git a/benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json b/benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json new file mode 100644 index 0000000..673cbc2 --- /dev/null +++ b/benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json @@ -0,0 +1,1776 @@ +{ + "mode": "production-trajectory", + "environment": { + "platform": "Linux-6.17.0-1026-nvidia-aarch64-with-glibc2.39", + "python": "3.12.3 (main, Mar 23 2026, 19:04:32) [GCC 13.3.0]", + "torch": "2.9.1+cu130", + "torch_cuda": "13.0", + "device": "NVIDIA GB10", + "device_capability": [ + 12, + 1 + ], + "driver": null, + "git_commit": null, + "checkpoint_path": "/models/minimax_h3_fl2va_pruned_nvfp4.safetensors", + "checkpoint_sha256": null, + "checkpoint_hash_note": "not calculated", + "environment_switches": { + "CUDA_DEVICE_MAX_CONNECTIONS": "1", + "CUDA_DEVICE_MAX_COPY_CONNECTIONS": "4", + "CUDA_HOME": "/usr/local/cuda", + "CUDA_INC_PATH": "/usr/local/cuda/include", + "CUDA_MANAGED_FORCE_DEVICE_ALLOC": "1", + "CUDA_MODULE_LOADING": "EAGER", + "CUDA_VERSION": "13.0.2", + "H3_FUSED_ELEMENTWISE": "1", + "H3_MODEL_PATH": "/models/minimax_h3_ref2va_pruned_nvfp4.safetensors", + "H3_NVFP4_MODULATE_FUSION": "1", + "H3_NVFP4_SCALE_BACKEND": "vortex", + "H3_NVFP4_SCALE_VERSION": "1", + "H3_NVFP4_SWIGLU_FUSION": "1", + "TORCH_COMPILE_DISABLE": "0", + "TORCH_CUDA_ARCH_LIST": "12.1a", + "TORCH_EXTENSIONS_DIR": "/opt/h3-blackwell-runtime/.torch_extensions" + }, + "extension": { + "cuda_version": 13000, + "cublas_version": 130100, + "cuda_runtime_version": 13000, + "cublaslt_runtime_version": 130000, + "stream_k_public_control": false, + "stream_k_note": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + }, + "comfy_kitchen": "0.2.31 package without __version__" + }, + "workload": { + "resolution": [ + 1344, + 768 + ], + "frames": 124, + "steps": 12, + "sampler_step": 1, + "seed": 440420, + "text_tokens": 100, + "tokens": 37810, + "hidden_shape": [ + 37810, + 5376 + ], + "segments": [ + [ + 0, + 100, + 1 + ], + [ + 100, + 514, + 2 + ], + [ + 514, + 37810, + 0 + ] + ] + }, + "retained_blocks": [ + 0, + 24, + 49 + ], + "immutable_cloned_block_inputs": { + "0": [ + 37810, + 5376 + ], + "24": [ + 37810, + 5376 + ], + "49": [ + 37810, + 5376 + ] + }, + "fc2_boundary": { + "block": 24, + "gate_up_shape": [ + 37810, + 28672 + ], + "activation_qdata_shape": [ + 37824, + 7168 + ], + "weight_qdata_shape": [ + 5376, + 7168 + ], + "logical_mnk": [ + 37810, + 5376, + 14336 + ], + "descriptor_mnk_after_padding": [ + 37824, + 5376, + 14336 + ], + "producer": "vortex_native_quantize_swiglu_nvfp4", + "no_bias": true + }, + "baseline_kernel_metadata": { + "path": "fc2.forward_swiglu -> accepted producer -> Comfy Kitchen 0.2.31 scaled_mm_nvfp4", + "descriptors": { + "packed_input_output": "row-major [M,K] @ [N,K].T -> BF16 [M,N]", + "block_scale_mode": "VEC16_UE4M3", + "compute_and_scale": "FP32", + "scalar_pointer_mode": "device", + "bias": null, + "beta": 0.0, + "comfy_kitchen_version": "0.2.31" + }, + "profiler_cuda_events_available": false, + "profiler_note": "Torch profiler returned no CUDA kernel events on this build; use --mode profile with NCU for kernel metadata.", + "top_cuda_events": [], + "output_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f" + }, + "heuristics": [ + { + "max_workspace_bytes": 0, + "requested_count": 32, + "returned_count": 5, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 4194304, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 8388608, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 16777216, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 33554432, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 67108864, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + } + ], + "explicit_split_k_checks": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 1, + "reduction_scheme": 0 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 4 + } + } + ], + "selected": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0 + }, + "trajectory": { + "steps": 12, + "resolution": [ + 1344, + 768 + ], + "frames": 124, + "seed": 440420, + "baseline_seconds": 286.43065266899976, + "candidate_seconds": 262.97871506399997, + "improvement_percent": 8.187649396624085, + "video_parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "c62d23a42972eab907ba42f93c50247ff17a9c454b4a53fe93d2e34f9fefe578", + "expected_sha256": "c62d23a42972eab907ba42f93c50247ff17a9c454b4a53fe93d2e34f9fefe578" + }, + "audio_parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "852005383770480a6503504e1ffec86dd1fb63a69c6400f92da18e39e0986de2", + "expected_sha256": "852005383770480a6503504e1ffec86dd1fb63a69c6400f92da18e39e0986de2" + }, + "bf16_exact": true, + "selected_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0 + }, + "supplied_workspace_bytes": 0, + "all_50_fc2_calls_replaced": true, + "accepted_swiglu_producer_preserved": true, + "accepted_gate_and_residual_path_preserved": true, + "production_dispatch": true, + "dispatch_delta": { + "attempts": 600, + "successes": 600, + "fallbacks": 0 + } + }, + "errors_and_unsupported": [ + { + "feature": "Stream-K", + "supported_public_control": false, + "reason": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + } + ] +} diff --git a/benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json b/benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json new file mode 100644 index 0000000..906fe4c --- /dev/null +++ b/benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json @@ -0,0 +1,1776 @@ +{ + "mode": "production-trajectory", + "environment": { + "platform": "Linux-6.17.0-1026-nvidia-aarch64-with-glibc2.39", + "python": "3.12.3 (main, Mar 23 2026, 19:04:32) [GCC 13.3.0]", + "torch": "2.9.1+cu130", + "torch_cuda": "13.0", + "device": "NVIDIA GB10", + "device_capability": [ + 12, + 1 + ], + "driver": null, + "git_commit": null, + "checkpoint_path": "/models/minimax_h3_fl2va_pruned_nvfp4.safetensors", + "checkpoint_sha256": null, + "checkpoint_hash_note": "not calculated", + "environment_switches": { + "CUDA_DEVICE_MAX_CONNECTIONS": "1", + "CUDA_DEVICE_MAX_COPY_CONNECTIONS": "4", + "CUDA_HOME": "/usr/local/cuda", + "CUDA_INC_PATH": "/usr/local/cuda/include", + "CUDA_MANAGED_FORCE_DEVICE_ALLOC": "1", + "CUDA_MODULE_LOADING": "EAGER", + "CUDA_VERSION": "13.0.2", + "H3_FUSED_ELEMENTWISE": "1", + "H3_MODEL_PATH": "/models/minimax_h3_ref2va_pruned_nvfp4.safetensors", + "H3_NVFP4_MODULATE_FUSION": "1", + "H3_NVFP4_SCALE_BACKEND": "vortex", + "H3_NVFP4_SCALE_VERSION": "1", + "H3_NVFP4_SWIGLU_FUSION": "1", + "TORCH_COMPILE_DISABLE": "0", + "TORCH_CUDA_ARCH_LIST": "12.1a", + "TORCH_EXTENSIONS_DIR": "/opt/h3-blackwell-runtime/.torch_extensions" + }, + "extension": { + "cuda_version": 13000, + "cublas_version": 130100, + "cuda_runtime_version": 13000, + "cublaslt_runtime_version": 130000, + "stream_k_public_control": false, + "stream_k_note": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + }, + "comfy_kitchen": "0.2.31 package without __version__" + }, + "workload": { + "resolution": [ + 1344, + 768 + ], + "frames": 124, + "steps": 2, + "sampler_step": 1, + "seed": 440420, + "text_tokens": 100, + "tokens": 37810, + "hidden_shape": [ + 37810, + 5376 + ], + "segments": [ + [ + 0, + 100, + 1 + ], + [ + 100, + 514, + 2 + ], + [ + 514, + 37810, + 0 + ] + ] + }, + "retained_blocks": [ + 0, + 24, + 49 + ], + "immutable_cloned_block_inputs": { + "0": [ + 37810, + 5376 + ], + "24": [ + 37810, + 5376 + ], + "49": [ + 37810, + 5376 + ] + }, + "fc2_boundary": { + "block": 24, + "gate_up_shape": [ + 37810, + 28672 + ], + "activation_qdata_shape": [ + 37824, + 7168 + ], + "weight_qdata_shape": [ + 5376, + 7168 + ], + "logical_mnk": [ + 37810, + 5376, + 14336 + ], + "descriptor_mnk_after_padding": [ + 37824, + 5376, + 14336 + ], + "producer": "vortex_native_quantize_swiglu_nvfp4", + "no_bias": true + }, + "baseline_kernel_metadata": { + "path": "fc2.forward_swiglu -> accepted producer -> Comfy Kitchen 0.2.31 scaled_mm_nvfp4", + "descriptors": { + "packed_input_output": "row-major [M,K] @ [N,K].T -> BF16 [M,N]", + "block_scale_mode": "VEC16_UE4M3", + "compute_and_scale": "FP32", + "scalar_pointer_mode": "device", + "bias": null, + "beta": 0.0, + "comfy_kitchen_version": "0.2.31" + }, + "profiler_cuda_events_available": false, + "profiler_note": "Torch profiler returned no CUDA kernel events on this build; use --mode profile with NCU for kernel metadata.", + "top_cuda_events": [], + "output_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f" + }, + "heuristics": [ + { + "max_workspace_bytes": 0, + "requested_count": 32, + "returned_count": 5, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 4194304, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 8388608, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 16777216, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 33554432, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 67108864, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + } + ], + "explicit_split_k_checks": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 1, + "reduction_scheme": 0 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 4 + } + } + ], + "selected": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0 + }, + "trajectory": { + "steps": 2, + "resolution": [ + 1344, + 768 + ], + "frames": 124, + "seed": 440420, + "baseline_seconds": 47.30056222799976, + "candidate_seconds": 43.727866895999796, + "improvement_percent": 7.553177306389591, + "video_parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "7d5e50d8f0ee9a639feede8a0e6aea1bd160087634ecc7ae80d0ef5b96fd8e78", + "expected_sha256": "7d5e50d8f0ee9a639feede8a0e6aea1bd160087634ecc7ae80d0ef5b96fd8e78" + }, + "audio_parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "df518aa8c645ec45e7cc46be58a941790f9766b420adf2cdcbaf14ef8b0e5605", + "expected_sha256": "df518aa8c645ec45e7cc46be58a941790f9766b420adf2cdcbaf14ef8b0e5605" + }, + "bf16_exact": true, + "selected_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0 + }, + "supplied_workspace_bytes": 0, + "all_50_fc2_calls_replaced": true, + "accepted_swiglu_producer_preserved": true, + "accepted_gate_and_residual_path_preserved": true, + "production_dispatch": true, + "dispatch_delta": { + "attempts": 100, + "successes": 100, + "fallbacks": 0 + } + }, + "errors_and_unsupported": [ + { + "feature": "Stream-K", + "supported_public_control": false, + "reason": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + } + ] +} diff --git a/benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json b/benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json new file mode 100644 index 0000000..e874f0a --- /dev/null +++ b/benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json @@ -0,0 +1,2404 @@ +{ + "mode": "shape-gate", + "environment": { + "platform": "Linux-6.17.0-1026-nvidia-aarch64-with-glibc2.39", + "python": "3.12.3 (main, Mar 23 2026, 19:04:32) [GCC 13.3.0]", + "torch": "2.9.1+cu130", + "torch_cuda": "13.0", + "device": "NVIDIA GB10", + "device_capability": [ + 12, + 1 + ], + "driver": null, + "git_commit": null, + "checkpoint_path": "/models/minimax_h3_fl2va_pruned_nvfp4.safetensors", + "checkpoint_sha256": null, + "checkpoint_hash_note": "not calculated", + "environment_switches": { + "CUDA_DEVICE_MAX_CONNECTIONS": "1", + "CUDA_DEVICE_MAX_COPY_CONNECTIONS": "4", + "CUDA_HOME": "/usr/local/cuda", + "CUDA_INC_PATH": "/usr/local/cuda/include", + "CUDA_MANAGED_FORCE_DEVICE_ALLOC": "1", + "CUDA_MODULE_LOADING": "EAGER", + "CUDA_VERSION": "13.0.2", + "H3_FUSED_ELEMENTWISE": "1", + "H3_MODEL_PATH": "/models/minimax_h3_ref2va_pruned_nvfp4.safetensors", + "H3_NVFP4_MODULATE_FUSION": "1", + "H3_NVFP4_SCALE_BACKEND": "vortex", + "H3_NVFP4_SCALE_VERSION": "1", + "H3_NVFP4_SWIGLU_FUSION": "1", + "H3_SAGE_QKV_LAYOUT": "strided_nhd", + "TORCH_COMPILE_DISABLE": "0", + "TORCH_CUDA_ARCH_LIST": "12.1a", + "TORCH_EXTENSIONS_DIR": "/opt/h3-blackwell-runtime/.torch_extensions" + }, + "extension": { + "cuda_version": 13000, + "cublas_version": 130100, + "stream_k_public_control": false, + "stream_k_note": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + }, + "comfy_kitchen": "0.2.31 package without __version__" + }, + "workload": { + "resolution": [ + 1344, + 768 + ], + "frames": 124, + "steps": 12, + "sampler_step": 1, + "seed": 440420, + "text_tokens": 100, + "tokens": 37810, + "hidden_shape": [ + 37810, + 5376 + ], + "segments": [ + [ + 0, + 100, + 1 + ], + [ + 100, + 514, + 2 + ], + [ + 514, + 37810, + 0 + ] + ] + }, + "retained_blocks": [ + 0, + 24, + 49 + ], + "immutable_cloned_block_inputs": { + "0": [ + 37810, + 5376 + ], + "24": [ + 37810, + 5376 + ], + "49": [ + 37810, + 5376 + ] + }, + "fc2_boundary": { + "block": 24, + "gate_up_shape": [ + 37810, + 28672 + ], + "activation_qdata_shape": [ + 37824, + 7168 + ], + "weight_qdata_shape": [ + 5376, + 7168 + ], + "logical_mnk": [ + 37810, + 5376, + 14336 + ], + "descriptor_mnk_after_padding": [ + 37824, + 5376, + 14336 + ], + "producer": "vortex_native_quantize_swiglu_nvfp4", + "no_bias": true + }, + "baseline_kernel_metadata": { + "path": "fc2.forward_swiglu -> accepted producer -> Comfy Kitchen 0.2.31 scaled_mm_nvfp4", + "descriptors": { + "packed_input_output": "row-major [M,K] @ [N,K].T -> BF16 [M,N]", + "block_scale_mode": "VEC16_UE4M3", + "compute_and_scale": "FP32", + "scalar_pointer_mode": "device", + "bias": null, + "beta": 0.0, + "comfy_kitchen_version": "0.2.31" + }, + "profiler_cuda_events_available": false, + "profiler_note": "Torch profiler returned no CUDA kernel events on this build; use --mode profile with NCU for kernel metadata.", + "top_cuda_events": [], + "output_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f" + }, + "heuristics": [ + { + "max_workspace_bytes": 0, + "requested_count": 32, + "returned_count": 5, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 4194304, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 8388608, + "requested_count": 32, + "returned_count": 7, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 16777216, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 33554432, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + }, + { + "max_workspace_bytes": 67108864, + "requested_count": 32, + "returned_count": 6, + "algorithms": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": -2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + } + ] + } + ], + "explicit_split_k_checks": [ + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 1, + "reduction_scheme": 0 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 2, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 2, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 4, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 4, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 8, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 8, + "reduction_scheme": 4 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 2, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 2 + } + }, + { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 16, + "reduction_scheme": 4, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 0.0, + "state": 15, + "api_status": 15, + "valid": false, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + }, + "requested_config": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "split_k": 16, + "reduction_scheme": 4 + } + } + ], + "selected": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0 + }, + "shape_gate": [ + { + "text_tokens": 1, + "logical_rows": 37711, + "packed_rows": 37712, + "parity": { + "bf16_exact": false, + "different_elements": 2, + "max_abs": 0.03125, + "mean_abs": 1.553468603754382e-10, + "actual_sha256": "c1ee129ef17b93c7814dc0dfb4dfc9fc0218a3d4252cddb02c80eb98e0471ab5", + "expected_sha256": "e017519cd7e89a93770f7648e81c637ccac517fe2715dd213f46627dfdfed711" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 57.960289001464844 + ], + "p50_ms": 57.960289001464844, + "p95_ms": 57.960289001464844, + "mean_ms": 57.960289001464844, + "dense_tflop_s_p50": 100.28933571475277 + }, + "candidate": { + "samples_ms": [ + 18.854400634765625 + ], + "p50_ms": 18.854400634765625, + "p95_ms": 18.854400634765625, + "mean_ms": 18.854400634765625, + "dense_tflop_s_p50": 308.2993193150771 + }, + "parity": { + "bf16_exact": false, + "different_elements": 2, + "max_abs": 0.03125, + "mean_abs": 1.553468603754382e-10, + "actual_sha256": "c1ee129ef17b93c7814dc0dfb4dfc9fc0218a3d4252cddb02c80eb98e0471ab5", + "expected_sha256": "e017519cd7e89a93770f7648e81c637ccac517fe2715dd213f46627dfdfed711" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": false + }, + { + "text_tokens": 15, + "logical_rows": 37725, + "packed_rows": 37728, + "parity": { + "bf16_exact": false, + "different_elements": 2, + "max_abs": 0.03125, + "mean_abs": 1.5528919816709674e-10, + "actual_sha256": "c47d817c29526fcf16d800bba07d8ac659831fe16076c07b5b5ba5347ed63c7e", + "expected_sha256": "310f6ecf71f385122512a2289282a9a8723e231aa252223201e23e9857219eda" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 57.9911994934082 + ], + "p50_ms": 57.9911994934082, + "p95_ms": 57.9911994934082, + "mean_ms": 57.9911994934082, + "dense_tflop_s_p50": 100.27309146900781 + }, + "candidate": { + "samples_ms": [ + 18.6856632232666 + ], + "p50_ms": 18.6856632232666, + "p95_ms": 18.6856632232666, + "mean_ms": 18.6856632232666, + "dense_tflop_s_p50": 311.1988470368801 + }, + "parity": { + "bf16_exact": false, + "different_elements": 2, + "max_abs": 0.03125, + "mean_abs": 1.5528919816709674e-10, + "actual_sha256": "c47d817c29526fcf16d800bba07d8ac659831fe16076c07b5b5ba5347ed63c7e", + "expected_sha256": "310f6ecf71f385122512a2289282a9a8723e231aa252223201e23e9857219eda" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": false + }, + { + "text_tokens": 32, + "logical_rows": 37742, + "packed_rows": 37744, + "parity": { + "bf16_exact": false, + "different_elements": 2, + "max_abs": 0.03125, + "mean_abs": 1.5521925411654536e-10, + "actual_sha256": "c50b1f08167249b90c80632f4bf1c379b252446cd766ac1ba331087c2de2744d", + "expected_sha256": "42816a3c87083bd5ed98a975cd6c7b31f35746761089fa8f2d19c92f47fed3df" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 57.83795166015625 + ], + "p50_ms": 57.83795166015625, + "p95_ms": 57.83795166015625, + "mean_ms": 57.83795166015625, + "dense_tflop_s_p50": 100.58408148350189 + }, + "candidate": { + "samples_ms": [ + 18.732128143310547 + ], + "p50_ms": 18.732128143310547, + "p95_ms": 18.732128143310547, + "mean_ms": 18.732128143310547, + "dense_tflop_s_p50": 310.56680789905454 + }, + "parity": { + "bf16_exact": false, + "different_elements": 2, + "max_abs": 0.03125, + "mean_abs": 1.5521925411654536e-10, + "actual_sha256": "c50b1f08167249b90c80632f4bf1c379b252446cd766ac1ba331087c2de2744d", + "expected_sha256": "42816a3c87083bd5ed98a975cd6c7b31f35746761089fa8f2d19c92f47fed3df" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": false + }, + { + "text_tokens": 64, + "logical_rows": 37774, + "packed_rows": 37776, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "9b272d355084277500dc89e4d633d44c7daa81481b092d3b96f577e2c207ddac", + "expected_sha256": "9b272d355084277500dc89e4d633d44c7daa81481b092d3b96f577e2c207ddac" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 58.0667839050293 + ], + "p50_ms": 58.0667839050293, + "p95_ms": 58.0667839050293, + "mean_ms": 58.0667839050293, + "dense_tflop_s_p50": 100.27264044192569 + }, + "candidate": { + "samples_ms": [ + 19.07366371154785 + ], + "p50_ms": 19.07366371154785, + "p95_ms": 19.07366371154785, + "mean_ms": 19.07366371154785, + "dense_tflop_s_p50": 305.26435991439087 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "9b272d355084277500dc89e4d633d44c7daa81481b092d3b96f577e2c207ddac", + "expected_sha256": "9b272d355084277500dc89e4d633d44c7daa81481b092d3b96f577e2c207ddac" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": true + }, + { + "text_tokens": 99, + "logical_rows": 37809, + "packed_rows": 37824, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "dcc4ac6d680a76bf563145e569ea114364a0ff3e76f3f2f656d1dd70adfba812", + "expected_sha256": "dcc4ac6d680a76bf563145e569ea114364a0ff3e76f3f2f656d1dd70adfba812" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 57.88643264770508 + ], + "p50_ms": 57.88643264770508, + "p95_ms": 57.88643264770508, + "mean_ms": 57.88643264770508, + "dense_tflop_s_p50": 100.6782487895987 + }, + "candidate": { + "samples_ms": [ + 18.773759841918945 + ], + "p50_ms": 18.773759841918945, + "p95_ms": 18.773759841918945, + "mean_ms": 18.773759841918945, + "dense_tflop_s_p50": 310.4282102637308 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "dcc4ac6d680a76bf563145e569ea114364a0ff3e76f3f2f656d1dd70adfba812", + "expected_sha256": "dcc4ac6d680a76bf563145e569ea114364a0ff3e76f3f2f656d1dd70adfba812" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": true + }, + { + "text_tokens": 100, + "logical_rows": 37810, + "packed_rows": 37824, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f", + "expected_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 57.96451187133789 + ], + "p50_ms": 57.96451187133789, + "p95_ms": 57.96451187133789, + "mean_ms": 57.96451187133789, + "dense_tflop_s_p50": 100.54529263105621 + }, + "candidate": { + "samples_ms": [ + 18.34467124938965 + ], + "p50_ms": 18.34467124938965, + "p95_ms": 18.34467124938965, + "mean_ms": 18.34467124938965, + "dense_tflop_s_p50": 317.6976424973496 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f", + "expected_sha256": "ad26c744afbb0249bff9afea503984e51217f85a59434e44fde347ae7d9b852f" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": true + }, + { + "text_tokens": 128, + "logical_rows": 37838, + "packed_rows": 37840, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "adb567fdd1568c406d514983f0dbf799b3d9121656788a853fe8620bca7cb67b", + "expected_sha256": "adb567fdd1568c406d514983f0dbf799b3d9121656788a853fe8620bca7cb67b" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 57.80854415893555 + ], + "p50_ms": 57.80854415893555, + "p95_ms": 57.80854415893555, + "mean_ms": 57.80854415893555, + "dense_tflop_s_p50": 100.89122346864156 + }, + "candidate": { + "samples_ms": [ + 19.78188705444336 + ], + "p50_ms": 19.78188705444336, + "p95_ms": 19.78188705444336, + "mean_ms": 19.78188705444336, + "dense_tflop_s_p50": 294.8340939913488 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "adb567fdd1568c406d514983f0dbf799b3d9121656788a853fe8620bca7cb67b", + "expected_sha256": "adb567fdd1568c406d514983f0dbf799b3d9121656788a853fe8620bca7cb67b" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": true + }, + { + "text_tokens": 256, + "logical_rows": 37966, + "packed_rows": 37968, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "e264df4d576726d22bf75e6e33a7bc3880543af28a6a528a8332b6fe25aa7676", + "expected_sha256": "e264df4d576726d22bf75e6e33a7bc3880543af28a6a528a8332b6fe25aa7676" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 58.237281799316406 + ], + "p50_ms": 58.237281799316406, + "p95_ms": 58.237281799316406, + "mean_ms": 58.237281799316406, + "dense_tflop_s_p50": 100.4872578586024 + }, + "candidate": { + "samples_ms": [ + 18.483808517456055 + ], + "p50_ms": 18.483808517456055, + "p95_ms": 18.483808517456055, + "mean_ms": 18.483808517456055, + "dense_tflop_s_p50": 316.60708601397215 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "e264df4d576726d22bf75e6e33a7bc3880543af28a6a528a8332b6fe25aa7676", + "expected_sha256": "e264df4d576726d22bf75e6e33a7bc3880543af28a6a528a8332b6fe25aa7676" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": true + }, + { + "text_tokens": 512, + "logical_rows": 38222, + "packed_rows": 38224, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "b6276cda42be6a858bab99ae69a211c71400104f39bc28b2d6723ad37be015b9", + "expected_sha256": "b6276cda42be6a858bab99ae69a211c71400104f39bc28b2d6723ad37be015b9" + }, + "timing": { + "order": "AB/BA alternates by round", + "baseline": { + "samples_ms": [ + 58.96633529663086 + ], + "p50_ms": 58.96633529663086, + "p95_ms": 58.96633529663086, + "mean_ms": 58.96633529663086, + "dense_tflop_s_p50": 99.91403968970451 + }, + "candidate": { + "samples_ms": [ + 20.196256637573242 + ], + "p50_ms": 20.196256637573242, + "p95_ms": 20.196256637573242, + "mean_ms": 20.196256637573242, + "dense_tflop_s_p50": 291.71568132201764 + }, + "parity": { + "bf16_exact": true, + "different_elements": 0, + "max_abs": 0.0, + "mean_abs": 0.0, + "actual_sha256": "b6276cda42be6a858bab99ae69a211c71400104f39bc28b2d6723ad37be015b9", + "expected_sha256": "b6276cda42be6a858bab99ae69a211c71400104f39bc28b2d6723ad37be015b9" + } + }, + "checked": { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, + "inner_shape": null, + "cluster_shape": null, + "required_workspace_bytes": 0, + "waves": 1.0, + "state": 0, + "api_status": 0, + "valid": true, + "capabilities": { + "split_k_support": 1, + "reduction_scheme_mask": 6, + "cta_swizzle_support": 0, + "custom_option_max": 0, + "strided_batch_support": 1, + "out_of_place_result_support": 1, + "tile_ids": [ + 20 + ], + "stages_ids": [ + 37 + ], + "inner_cluster_shape_capability_note": "This CUDA 13 cublasLt.h exposes config IDs but no public capability attributes that enumerate inner/cluster shape IDs." + } + }, + "supplied_workspace_bytes": 0, + "passed": true + } + ], + "errors_and_unsupported": [ + { + "feature": "Stream-K", + "supported_public_control": false, + "reason": "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute that directly selects Stream-K. Negative SPLITK_NUM values returned by heuristics are preserved as undocumented library sentinels, not claimed as public Stream-K control." + } + ] +} diff --git a/compose.spark-stack.yml b/compose.spark-stack.yml index a427862..b34fe06 100644 --- a/compose.spark-stack.yml +++ b/compose.spark-stack.yml @@ -20,6 +20,7 @@ services: H3_FUSED_ELEMENTWISE: "1" H3_NVFP4_MODULATE_FUSION: "1" H3_NVFP4_SWIGLU_FUSION: "1" + H3_NVFP4_FC2_LT_SPLITK1: "1" H3_SAGE_QKV_LAYOUT: "strided_nhd" H3_SOL_QKV_LAYOUT: "native" command: ["python", "/opt/h3-blackwell-runtime/tools/serve_hot_runtime.py", "--host", "0.0.0.0", "--port", "8000", "--attention", "sage2", "--warmup"] diff --git a/compose.spark.yml b/compose.spark.yml index 2f9c0ce..5d82bcd 100644 --- a/compose.spark.yml +++ b/compose.spark.yml @@ -37,6 +37,7 @@ services: H3_FUSED_ELEMENTWISE: "1" H3_NVFP4_MODULATE_FUSION: "1" H3_NVFP4_SWIGLU_FUSION: "1" + H3_NVFP4_FC2_LT_SPLITK1: "1" H3_SAGE_QKV_LAYOUT: "strided_nhd" H3_SOL_QKV_LAYOUT: "native" command: ["python", "/opt/h3-blackwell-runtime/tools/serve_hot_runtime.py", "--host", "0.0.0.0", "--port", "8000", "--attention", "sage2", "--turbo-lora", "4step=/turbo/minimax_h3_fl2v_turbo_4step_v1.1_768p_comfyui_bf16.safetensors", "--turbo-lora", "8step=/turbo/minimax_h3_fl2v_turbo_8step_v1.0_comfyui_bf16.safetensors", "--latent-upscaler", "/latent-upscaler/minimax_h3_latent_upscaler_3d_fp16.safetensors", "--warmup"] diff --git a/pyproject.toml b/pyproject.toml index 4ff4f44..d158804 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -17,3 +17,6 @@ flash4 = ["flash-attn-4[cu13]==4.0.0b27"] [tool.setuptools.packages.find] where = ["src"] + +[tool.setuptools.package-data] +h3_blackwell_runtime = ["csrc/*.cpp", "csrc/*.cu"] diff --git a/research/artifact_manifest.json b/research/artifact_manifest.json index 2dfd7f2..044ef26 100644 --- a/research/artifact_manifest.json +++ b/research/artifact_manifest.json @@ -2,36 +2,36 @@ "metadata": { "algorithm": "SHA-256", "generated_date": "2026-08-25", - "scope_note": "Local selected scope is the established 167-file set plus 10 retained FC2 scheduling artifacts. Spark records are the complete reproducible current top-level benchmark output set; the audit-retained 279-file/665950155-byte aggregate cannot be reconstructed because its path list was not retained.", + "scope_note": "Local selected scope is the established 167-file set plus 14 retained FC2 scheduling and integration artifacts. Spark records are the complete reproducible current top-level benchmark output set; the audit-retained 279-file/665950155-byte aggregate cannot be reconstructed because its path list was not retained.", "summary": { "local": { - "record_count": 177, - "size_bytes": 230157548, - "expected_record_count": 177, - "expected_size_bytes": 230157548, + "record_count": 181, + "size_bytes": 230398680, + "expected_record_count": 181, + "expected_size_bytes": 230398680, "reconciled": true }, "spark": { - "record_count": 290, - "size_bytes": 567477443, + "record_count": 294, + "size_bytes": 567718575, "expected_record_count": 279, "expected_size_bytes": 665950155, "reconciled": false, - "record_count_delta": 11, - "size_bytes_delta": -98472712 + "record_count_delta": 15, + "size_bytes_delta": -98231580 }, "total": { - "record_count": 467, - "size_bytes": 797634991 + "record_count": 475, + "size_bytes": 798117255 } }, "json_reconciliation": { - "identical": 112, - "mismatches": 2, + "identical": 118, + "mismatches": 0, "local_only": 40, "spark_only": 107, - "local_total": 154, - "spark_total": 221 + "local_total": 158, + "spark_total": 225 } }, "artifacts": [ @@ -195,6 +195,70 @@ "artifact_class": "benchmark_json", "location": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-trajectory-2step-20260825.json" }, + { + "scope": "local", + "path": "benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json", + "size_bytes": 61575, + "sha256": "a49664031af6c21a341cfb0787000e54801c4067cf8ad231cc614c11de251067", + "artifact_class": "benchmark_json", + "location": "benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json" + }, + { + "scope": "local", + "path": "benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json", + "size_bytes": 53189, + "sha256": "3e0e42addf5a748462f1bc600d5bb08875b31da5177d2af2b2abddfb1078ebd3", + "artifact_class": "benchmark_json", + "location": "benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json" + }, + { + "scope": "local", + "path": "benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json", + "size_bytes": 53192, + "sha256": "94a915acee1bf3d3c23a934f93bafc583ac2ae0f3e91ff1277ee57f628e2d254", + "artifact_class": "benchmark_json", + "location": "benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json" + }, + { + "scope": "local", + "path": "benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json", + "size_bytes": 73176, + "sha256": "cbde2aab33ceb9df32e65c8649838dcdf47e0e135aaaae4165b2f9b6e8d2734b", + "artifact_class": "benchmark_json", + "location": "benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json" + }, + { + "scope": "spark", + "path": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json", + "size_bytes": 61575, + "sha256": "a49664031af6c21a341cfb0787000e54801c4067cf8ad231cc614c11de251067", + "artifact_class": "benchmark_json", + "location": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json" + }, + { + "scope": "spark", + "path": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json", + "size_bytes": 53189, + "sha256": "3e0e42addf5a748462f1bc600d5bb08875b31da5177d2af2b2abddfb1078ebd3", + "artifact_class": "benchmark_json", + "location": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json" + }, + { + "scope": "spark", + "path": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json", + "size_bytes": 53192, + "sha256": "94a915acee1bf3d3c23a934f93bafc583ac2ae0f3e91ff1277ee57f628e2d254", + "artifact_class": "benchmark_json", + "location": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json" + }, + { + "scope": "spark", + "path": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json", + "size_bytes": 73176, + "sha256": "cbde2aab33ceb9df32e65c8649838dcdf47e0e135aaaae4165b2f9b6e8d2734b", + "artifact_class": "benchmark_json", + "location": "/home/daniel/StoryStudioAssets/H3-output/h3-blackwell-runtime/benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json" + }, { "scope": "local", "path": "benchmarks/audio-dialogue-format-sweep-sage2-seeds440420-440429.json", diff --git a/research/experiment_registry.json b/research/experiment_registry.json index ff061b6..b3eab07 100644 --- a/research/experiment_registry.json +++ b/research/experiment_registry.json @@ -504,22 +504,22 @@ "id": "fc2-cublaslt-splitk1-schedule", "name": "FC2 cuBLASLt public split-K-1 schedule", "family": "nvfp4-library-scheduling", - "status": "research_retained", + "status": "production_accepted", "hypothesis": "A documented cuBLASLt schedule can preserve the exact FC2 reduction result while avoiding the production heuristic's traffic and synchronization regression.", "implementation_strategy": "Reproduce the exact Comfy Kitchen descriptors in an isolated extension, enumerate checked cuBLASLt algorithms, and compare one selected public split-K-1 schedule against the accepted FC2 path.", - "source_locations": ["research/fc2_nvfp4_scheduling/README.md", "research/fc2_nvfp4_scheduling/RESULTS.md", "research/fc2_nvfp4_scheduling/fc2_nvfp4_lt.cpp", "tools/benchmark_fc2_nvfp4_algorithms.py"], - "active_source_location": "research/fc2_nvfp4_scheduling/fc2_nvfp4_lt.cpp", + "source_locations": ["research/fc2_nvfp4_scheduling/README.md", "research/fc2_nvfp4_scheduling/RESULTS.md", "src/h3_blackwell_runtime/csrc/fc2_nvfp4_lt.cpp", "tools/benchmark_fc2_nvfp4_algorithms.py"], + "active_source_location": "src/h3_blackwell_runtime/csrc/fc2_nvfp4_lt.cpp", "commit_hash": null, - "benchmark_artifacts": [{"path": "benchmarks/gb10-fc2-nvfp4-library-sweep-20260825.json", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-block-gate-20260825.json", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-trajectory-2step-20260825.json", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-trajectory-12step-20260825.json", "exists": true}], + "benchmark_artifacts": [{"path": "benchmarks/gb10-fc2-nvfp4-library-sweep-20260825.json", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json", "exists": true}], "profiler_artifacts": [{"path": "benchmarks/gb10-fc2-nvfp4-baseline-20260825.ncu-rep", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-splitk1-20260825.ncu-rep", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-baseline-20260825.csv", "exists": true}, {"path": "benchmarks/gb10-fc2-nvfp4-splitk1-20260825.csv", "exists": true}], - "environment": {"gpu": "NVIDIA GB10", "cuda": "13", "driver": null, "pytorch": null, "triton": null, "container_image": "sha256:1d340e14cb6fc45ccfdbe63dde8db2a2b3aea94b493702c8a08e1f8d5b4f7b83", "commit_context": "isolated extension; production source and dispatch unchanged"}, - "metrics": {"fc2_p50_ms_baseline": 53.618, "fc2_p50_ms_candidate": 15.636, "block_improvement_pct": [8.16, 8.63, 8.88], "two_step_improvement_pct": 7.50, "canonical_12_step_improvement_pct": 8.21, "trajectory_measurement_note": "single baseline-then-candidate paired runs", "l2_hit_rate_pct_baseline": 53.32, "l2_hit_rate_pct_candidate": 91.10, "l2_read_miss_reduction_bytes": 9853094784}, - "correctness_evidence": ["Selected FC2 output is byte-exact.", "Blocks 0, 24, and 49 are byte-exact against paired baseline and retained traversal.", "Two-step and canonical 12-step video and audio latents are byte-exact."], - "decision_rationale": "The library schedule passed all exactness and performance gates, explaining the prior traffic amplification without requiring a custom kernel. It is retained pending explicit production integration and deployment validation.", - "reproducer_commands": ["python tools/benchmark_fc2_nvfp4_algorithms.py --mode sweep --rounds 20 --workspace-bytes 67108864", "python tools/benchmark_fc2_nvfp4_algorithms.py --mode block-gate --candidate research/fc2_nvfp4_scheduling/candidate_splitk1.json --rounds 20 --workspace-bytes 0"], + "environment": {"gpu": "NVIDIA GB10", "cuda": "13", "driver": null, "pytorch": "2.9.1+cu130", "triton": null, "container_image": "sha256:d61c8d29e8b59306d140bfd279b0de31d827ad13bd92447ea72ea0a04182494d", "commit_context": "guarded canonical production integration"}, + "metrics": {"fc2_p50_ms_baseline": 53.618, "fc2_p50_ms_candidate": 15.636, "production_block_improvement_pct": [9.36, 8.16, 7.86], "production_two_step_seconds": {"baseline": 47.30056222799976, "candidate": 43.727866895999796, "improvement_percent": 7.553177306389591}, "production_canonical_12_step_seconds": {"baseline": 286.43065266899976, "candidate": 262.97871506399997, "improvement_percent": 8.187649396624085}, "production_12_step_dispatches": 600, "production_12_step_fallbacks": 0, "l2_hit_rate_pct_baseline": 53.32, "l2_hit_rate_pct_candidate": 91.10, "l2_read_miss_reduction_bytes": 9853094784}, + "correctness_evidence": ["Production blocks 0, 24, and 49 are byte-exact against paired baseline and retained traversal.", "Production two-step and canonical 12-step video and audio latents are byte-exact.", "Nearby M=37711/37725/37742 shape probes differ in two BF16 elements and are excluded by the canonical guard."], + "decision_rationale": "The shape-guarded library schedule passed production-method block and trajectory gates without requiring a custom kernel. Noncanonical shapes retain the accepted fallback because AlgoCheck alone does not guarantee byte exactness.", + "reproducer_commands": ["python tools/benchmark_fc2_nvfp4_algorithms.py --mode production-block-gate --candidate research/fc2_nvfp4_scheduling/candidate_splitk1.json --rounds 20 --workspace-bytes 0", "python tools/benchmark_fc2_nvfp4_algorithms.py --mode production-trajectory --candidate research/fc2_nvfp4_scheduling/candidate_splitk1.json --steps 12 --workspace-bytes 0", "python tools/benchmark_fc2_nvfp4_algorithms.py --mode shape-gate --candidate research/fc2_nvfp4_scheduling/candidate_splitk1.json --workspace-bytes 0"], "timestamp": "2026-08-25", - "evidence_missing": ["Production Nvfp4Linear integration", "Production deployment smoke", "Portable validation outside GB10/SM121"], - "production_behavior": "Not dispatched. Production still uses Comfy Kitchen's heuristic-selected FC2 path.", + "evidence_missing": ["Resident-service deployment validation deferred until performance work completes", "Trajectory timings are single baseline-then-candidate pairs rather than repeated AB/BA trials", "Shape-gate artifact predates runtime-version fields in build_info", "Portable validation outside GB10/SM121"], + "production_behavior": "H3_NVFP4_FC2_LT_SPLITK1 selects algorithm 70 only for the validated canonical FC2 descriptor; every other shape, device, library, LoRA, and failure case falls back to Comfy Kitchen.", "source_recovery": "The direct cuBLASLt extension, benchmark harness, selected candidate, raw NCU reports, and full gate artifacts are retained in this checkout." }, { diff --git a/research/fc2_nvfp4_scheduling/README.md b/research/fc2_nvfp4_scheduling/README.md index 800a43c..e379330 100644 --- a/research/fc2_nvfp4_scheduling/README.md +++ b/research/fc2_nvfp4_scheduling/README.md @@ -1,14 +1,15 @@ # FC2 NVFP4 Library Scheduling This directory contains the first isolated library-scheduling study for H3 FC2. -It does not register a PyTorch operator, alter `Nvfp4Linear`, or participate in -production dispatch. The extension is loaded only by -`tools/benchmark_fc2_nvfp4_algorithms.py`. +The validated extension source moved to +`src/h3_blackwell_runtime/csrc/fc2_nvfp4_lt.cpp` when the guarded production +integration was added. The research harness and production loader share that +source. The completed measurements and decision are in `RESULTS.md`. Algorithm 70 with public split-K 1 is byte-exact and passes the block, two-step, and canonical -12-step gates, but remains research-only until production integration and -deployment validation are performed. +12-step gates. Canonical-only production integration is complete behind an +opt-in flag; resident-service deployment validation remains deferred. ## Exact Operation @@ -175,6 +176,10 @@ latents; see `RESULTS.md` and the linked benchmark artifacts. ## Prohibited Experiments +These restrictions governed the isolated study. Production integration was +authorized only after the library candidate passed the complete validation +ladder; the canonical-only guard and automatic fallback preserve that boundary. + - Do not alter production source, configuration, dispatch, or existing research files. - Do not expose this extension through `Nvfp4Linear` or any production operator. diff --git a/research/fc2_nvfp4_scheduling/RESULTS.md b/research/fc2_nvfp4_scheduling/RESULTS.md index 761bd1b..8909507 100644 --- a/research/fc2_nvfp4_scheduling/RESULTS.md +++ b/research/fc2_nvfp4_scheduling/RESULTS.md @@ -1,7 +1,7 @@ # FC2 NVFP4 Scheduling Results -Status: research retained; validated as a canonical-shape production-integration -candidate, but not wired into production dispatch. +Status: production accepted for the canonical descriptor shape behind +`H3_NVFP4_FC2_LT_SPLITK1`; all other shapes retain the existing fallback. ## Decision @@ -70,13 +70,52 @@ traversal. | Canonical 12-step | `278.201 s` | `255.371 s` | `8.21%` | video/audio BF16 exact | Trajectory values are single paired runs in baseline-then-candidate -order, not repeated medians. Their approximately `38 ms` per-FC2 savings agree +order, not repeated medians or AB/BA trials. Their approximately `38 ms` per-FC2 savings agree with the alternating isolated and complete-block gates, but the precise end-to-end percentages retain run-order uncertainty. The 12-step saving of `22.829 s` matches approximately 600 FC2 invocations multiplied by the isolated roughly `38 ms` saving. +## Production Integration + +The packaged path lives in `src/h3_blackwell_runtime/fc2_lt.py` and is called +from `Nvfp4Linear.forward_swiglu` only after the accepted native producer. It +requires the FC2 role, unsharded `K=14,336/N=5,376`, BF16 output, no bias or +active LoRA, GB10/SM121, CUDA 13.0, cuBLAS 13.1 headers, cuBLASLt 13.0 runtime, +logical `M=37,810`, and packed +`M=37,824`. A shape-specific `cublasLtMatmulAlgoCheck` must pass with zero +workspace. Load, check, or execution failure falls back through the already +produced NVFP4 bytes without quantizing twice; strict mode is available through +`H3_NVFP4_FC2_LT_STRICT=1`. Production startup prepares the extension and +caches its measured runtime compatibility result before inference. + +The 20-round production-method gate exercised `Nvfp4Linear.forward_swiglu` +directly and recorded 23 successful dispatches with zero fallbacks per block: + +| Block | Baseline p50 | Production p50 | Improvement | +| ---: | ---: | ---: | ---: | +| 0 | `464.153 ms` | `420.714 ms` | `9.36%` | +| 24 | `460.184 ms` | `422.656 ms` | `8.16%` | +| 49 | `456.183 ms` | `420.335 ms` | `7.86%` | + +Production-method trajectories remained exact: + +| Trajectory | Baseline | Production | Improvement | Dispatch | +| --- | ---: | ---: | ---: | --- | +| Two-step | `47.301 s` | `43.728 s` | `7.55%` | `100/100`, zero fallback | +| Canonical 12-step | `286.431 s` | `262.979 s` | `8.19%` | `600/600`, zero fallback | + +The row-shape gate prevents unsafe generalization. Algorithm 70 was exact for +the tested `M=37,774`, `37,809`, `37,810`, `37,838`, `37,966`, and `38,222` +inputs, but differed in two BF16 elements at each of `M=37,711`, `37,725`, and +`37,742` despite passing `AlgoCheck`. Production therefore intentionally keeps +the canonical-only guard rather than inferring exactness from nearby shapes. +The retained shape artifact records compile-time CUDA/cuBLAS versions but +predates runtime-version fields in `build_info()`. The final extension separately +measured CUDA runtime `13000` and cuBLASLt runtime `130000`; the shape artifact +alone is not runtime-ABI evidence. + ## Traffic Attribution The one-pass distinct-data footprint is: @@ -150,9 +189,12 @@ synchronization waiting, not occupancy or a reduction workspace. - `benchmarks/gb10-fc2-nvfp4-splitk1-20260825.csv` - `benchmarks/gb10-fc2-nvfp4-baseline-20260825.ncu-rep` - `benchmarks/gb10-fc2-nvfp4-splitk1-20260825.ncu-rep` +- `benchmarks/gb10-fc2-nvfp4-production-block-gate-20260825.json` +- `benchmarks/gb10-fc2-nvfp4-production-trajectory-2step-20260825.json` +- `benchmarks/gb10-fc2-nvfp4-production-trajectory-12step-20260825.json` +- `benchmarks/gb10-fc2-nvfp4-shape-gate-20260825.json` +- `src/h3_blackwell_runtime/csrc/fc2_nvfp4_lt.cpp` -Production source and configuration were not changed by this study. The -candidate is validated only for the canonical descriptor shape and current -GB10/CUDA-library combination. Promotion still requires shape-specific -`AlgoCheck` with a safe fallback, integration behind the existing -`Nvfp4Linear` boundary, broader shape tests, and deployment validation. +The Spark hot runtime remains stopped while performance work continues. +Resident-service validation is intentionally deferred until the performance +program is finished. diff --git a/research/fc2_nvfp4_scheduling/fc2_nvfp4_lt.cpp b/src/h3_blackwell_runtime/csrc/fc2_nvfp4_lt.cpp similarity index 93% rename from research/fc2_nvfp4_scheduling/fc2_nvfp4_lt.cpp rename to src/h3_blackwell_runtime/csrc/fc2_nvfp4_lt.cpp index 3cf79d8..5abb693 100644 --- a/research/fc2_nvfp4_scheduling/fc2_nvfp4_lt.cpp +++ b/src/h3_blackwell_runtime/csrc/fc2_nvfp4_lt.cpp @@ -23,13 +23,22 @@ namespace py = pybind11; namespace { -thread_local cublasLtHandle_t handle = nullptr; +struct LtHandle { + cublasLtHandle_t value = nullptr; + + LtHandle() { LT_CHECK(cublasLtCreate(&value)); } + ~LtHandle() { + if (value) cublasLtDestroy(value); + } +}; cublasLtHandle_t get_handle() { - if (!handle) LT_CHECK(cublasLtCreate(&handle)); - return handle; + static LtHandle handle; + return handle.value; } +void prepare() { get_handle(); } + void require_cuda_contiguous(const torch::Tensor& tensor, const char* name) { TORCH_CHECK(tensor.is_cuda(), name, " must be CUDA"); TORCH_CHECK(tensor.is_contiguous(), name, " must be contiguous"); @@ -50,6 +59,19 @@ struct Problem { const torch::Tensor& activation_block_scale, const torch::Tensor& weight_qdata, const torch::Tensor& weight_block_scale) { + try { + initialize(activation_qdata, activation_block_scale, weight_qdata, + weight_block_scale); + } catch (...) { + reset(); + throw; + } + } + + void initialize(const torch::Tensor& activation_qdata, + const torch::Tensor& activation_block_scale, + const torch::Tensor& weight_qdata, + const torch::Tensor& weight_block_scale) { require_cuda_contiguous(activation_qdata, "activation_qdata"); require_cuda_contiguous(activation_block_scale, "activation_block_scale"); require_cuda_contiguous(weight_qdata, "weight_qdata"); @@ -122,13 +144,17 @@ struct Problem { LT_CHECK(cublasLtMatrixLayoutCreate(&d, CUDA_R_16BF, m, n, m)); } - ~Problem() { + void reset() noexcept { if (d) cublasLtMatrixLayoutDestroy(d); if (c) cublasLtMatrixLayoutDestroy(c); if (b) cublasLtMatrixLayoutDestroy(b); if (a) cublasLtMatrixLayoutDestroy(a); if (operation) cublasLtMatmulDescDestroy(operation); + d = c = b = a = nullptr; + operation = nullptr; } + + ~Problem() { reset(); } }; template @@ -367,8 +393,15 @@ void run(torch::Tensor activation_qdata, py::dict build_info() { py::dict result; + int cuda_runtime_version = 0; + const cudaError_t cuda_status = cudaRuntimeGetVersion(&cuda_runtime_version); + TORCH_CHECK(cuda_status == cudaSuccess, "cudaRuntimeGetVersion failed: ", + cudaGetErrorString(cuda_status)); result["cuda_version"] = CUDA_VERSION; result["cublas_version"] = CUBLAS_VERSION; + result["cuda_runtime_version"] = cuda_runtime_version; + result["cublaslt_runtime_version"] = + static_cast(cublasLtGetVersion()); result["stream_k_public_control"] = false; result["stream_k_note"] = "CUDA 13 cuBLASLt exposes no documented MatmulAlgoConfig attribute " @@ -381,6 +414,7 @@ py::dict build_info() { } // namespace PYBIND11_MODULE(TORCH_EXTENSION_NAME, module) { + module.def("prepare", &prepare); module.def("enumerate", &enumerate); module.def("check", &check); module.def("run", &run); diff --git a/src/h3_blackwell_runtime/fc2_lt.py b/src/h3_blackwell_runtime/fc2_lt.py new file mode 100644 index 0000000..ad68fbb --- /dev/null +++ b/src/h3_blackwell_runtime/fc2_lt.py @@ -0,0 +1,197 @@ +"""Guarded production dispatch for the canonical GB10 FC2 cuBLASLt schedule.""" + +from __future__ import annotations + +import os +from functools import lru_cache +from pathlib import Path + +import torch + + +_TRUTHY = {"1", "true", "yes", "on"} +_CANONICAL_LOGICAL_SHAPE = (37_810, 14_336, 5_376) +_CANONICAL_PACKED_ROWS = 37_824 +_CONFIG = { + "algorithm_id": 70, + "tile_id": 20, + "stages_id": 37, + "split_k": 1, + "reduction_scheme": 0, + "custom_option": 0, + "cta_swizzle": 0, +} +_VALIDATED: set[tuple[int, tuple[int, ...], tuple[int, ...]]] = set() +_STATS = {"attempts": 0, "successes": 0, "fallbacks": 0} + + +class _Fc2LtStrictError(RuntimeError): + pass + + +def fc2_lt_enabled() -> bool: + return os.getenv("H3_NVFP4_FC2_LT_SPLITK1", "").lower() in _TRUTHY + + +def _fc2_lt_strict() -> bool: + return os.getenv("H3_NVFP4_FC2_LT_STRICT", "").lower() in _TRUTHY + + +def fc2_lt_status() -> dict: + return {"enabled": fc2_lt_enabled(), **_STATS} + + +@lru_cache(maxsize=1) +def _fc2_lt_extension(): + from torch.utils.cpp_extension import CUDA_HOME, load + + if CUDA_HOME is None: + raise RuntimeError("CUDA_HOME is unavailable") + source = Path(__file__).resolve().parent / "csrc" / "fc2_nvfp4_lt.cpp" + return load( + name="h3_fc2_nvfp4_lt", + sources=[str(source)], + extra_include_paths=[str(Path(CUDA_HOME) / "include")], + extra_cflags=["-O2", "-std=c++17"], + extra_ldflags=[ + "-L" + str(Path(CUDA_HOME) / "lib64"), + "-lcublasLt", + "-lcublas", + "-lcudart", + ], + verbose=os.getenv("H3_NVFP4_FC2_LT_VERBOSE", "").lower() in _TRUTHY, + ) + + +@lru_cache(maxsize=1) +def _fc2_lt_extension_result(): + try: + return _fc2_lt_extension(), None + except Exception as error: + return None, error + + +@lru_cache(maxsize=1) +def _fc2_lt_compatible_result(): + extension, load_error = _fc2_lt_extension_result() + if extension is None: + return None, load_error + try: + info = dict(extension.build_info()) + compatible = ( + info.get("cuda_version") == 13_000 + and info.get("cublas_version") == 130_100 + and info.get("cuda_runtime_version") == 13_000 + and info.get("cublaslt_runtime_version") == 130_000 + ) + if not compatible: + return None, RuntimeError(f"unsupported FC2 cuBLASLt environment: {info}") + extension.prepare() + return extension, None + except Exception as error: + return None, error + + +def _canonical_fc2_supported(linear, gate_up: torch.Tensor, qdata: torch.Tensor) -> bool: + logical_rows, in_features, out_features = _CANONICAL_LOGICAL_SHAPE + return ( + linear.role == "h3_mlp_fc2" + and linear.in_features == in_features + and linear.out_features == out_features + and linear.output_dtype == torch.bfloat16 + and linear.bias is None + and tuple(gate_up.shape) == (logical_rows, in_features * 2) + and tuple(qdata.shape) == (_CANONICAL_PACKED_ROWS, in_features // 2) + and torch.cuda.get_device_capability(gate_up.device) == (12, 1) + ) + + +def _fallback_or_raise(message: str, error: Exception | None = None): + _STATS["fallbacks"] += 1 + if _fc2_lt_strict(): + if error is None: + raise _Fc2LtStrictError(message) + raise _Fc2LtStrictError(message) from error + return None + + +def prepare_fc2_lt() -> bool: + """Load and validate the extension before latency-sensitive inference.""" + if not fc2_lt_enabled(): + return False + extension, error = _fc2_lt_compatible_result() + if extension is None: + _fallback_or_raise("FC2 cuBLASLt extension is incompatible", error) + return False + return True + + +def fc2_lt_linear( + linear, + gate_up: torch.Tensor, + tensor_scale: torch.Tensor, + qdata: torch.Tensor, + block_scale: torch.Tensor, +) -> torch.Tensor | None: + """Run the validated canonical schedule, or return ``None`` for fallback.""" + if not fc2_lt_enabled(): + return None + try: + supported = _canonical_fc2_supported(linear, gate_up, qdata) + except Exception as error: + return _fallback_or_raise("FC2 cuBLASLt eligibility check failed", error) + if not supported: + _STATS["fallbacks"] += 1 + return None + _STATS["attempts"] += 1 + + extension, load_error = _fc2_lt_compatible_result() + if extension is None: + return _fallback_or_raise( + "FC2 cuBLASLt extension failed to load or is incompatible", load_error + ) + + try: + device = gate_up.device.index if gate_up.device.index is not None else torch.cuda.current_device() + validation_key = (device, tuple(qdata.shape), tuple(linear.weight.shape)) + if validation_key not in _VALIDATED: + checked = dict( + extension.check( + qdata, + block_scale, + linear.weight, + linear.weight_scale, + _CONFIG, + ) + ) + if not checked.get("valid") or checked.get("required_workspace_bytes") != 0: + return _fallback_or_raise( + f"FC2 cuBLASLt algorithm check rejected the candidate: {checked}" + ) + _VALIDATED.add(validation_key) + + alpha = (tensor_scale.float() * linear.weight_scale_2.float()).reshape(1).contiguous() + beta = torch.zeros_like(alpha) + output = torch.empty( + (qdata.shape[0], linear.out_features), + device=qdata.device, + dtype=torch.bfloat16, + ) + workspace = torch.empty(0, device=qdata.device, dtype=torch.uint8) + extension.run( + qdata, + block_scale, + linear.weight, + linear.weight_scale, + alpha, + beta, + output, + workspace, + _CONFIG, + ) + _STATS["successes"] += 1 + return output[: gate_up.shape[0], : linear.out_features] + except _Fc2LtStrictError: + raise + except Exception as error: + return _fallback_or_raise("FC2 cuBLASLt candidate execution failed", error) diff --git a/src/h3_blackwell_runtime/nvfp4.py b/src/h3_blackwell_runtime/nvfp4.py index 2c2d0ef..a19750d 100644 --- a/src/h3_blackwell_runtime/nvfp4.py +++ b/src/h3_blackwell_runtime/nvfp4.py @@ -164,9 +164,16 @@ class Nvfp4Linear(DynamicLoraMixin, nn.Module): if gate_up.dtype != torch.bfloat16 or not gate_up.is_cuda or not gate_up.is_contiguous(): raise ValueError("SwiGLU NVFP4 dispatch requires contiguous CUDA BF16 input") - from .nvfp4_quant import vortex_quantize_swiglu_nvfp4 + from .nvfp4_quant import vortex_native_quantize_swiglu_nvfp4, wrap_native_swiglu_nvfp4 - packed_x = vortex_quantize_swiglu_nvfp4(gate_up) + packed = vortex_native_quantize_swiglu_nvfp4(gate_up) + if os.getenv("H3_NVFP4_FC2_LT_SPLITK1", "").lower() in {"1", "true", "yes", "on"}: + from .fc2_lt import fc2_lt_linear + + scheduled = fc2_lt_linear(self, gate_up, *packed) + if scheduled is not None: + return scheduled + packed_x = wrap_native_swiglu_nvfp4(gate_up, packed) bias = self.bias.to(gate_up) if self.bias is not None else None output = functional.linear(packed_x, self._packed_weight(), bias)[ : gate_up.shape[0], : self.out_features diff --git a/src/h3_blackwell_runtime/nvfp4_quant.py b/src/h3_blackwell_runtime/nvfp4_quant.py index b964944..2bb4f0b 100644 --- a/src/h3_blackwell_runtime/nvfp4_quant.py +++ b/src/h3_blackwell_runtime/nvfp4_quant.py @@ -328,9 +328,16 @@ def vortex_native_quantize_swiglu_nvfp4(tensor: torch.Tensor): def vortex_quantize_swiglu_nvfp4(tensor: torch.Tensor): """Wrap fused SwiGLU output as a Comfy TensorCoreNVFP4Layout tensor.""" + return wrap_native_swiglu_nvfp4( + tensor, vortex_native_quantize_swiglu_nvfp4(tensor), + ) + + +def wrap_native_swiglu_nvfp4(tensor: torch.Tensor, packed): + """Wrap already-produced fused SwiGLU bytes for the Comfy fallback.""" from comfy_kitchen.tensor import QuantizedTensor, TensorCoreNVFP4Layout - tensor_scale, qdata, block_scale = vortex_native_quantize_swiglu_nvfp4(tensor) + tensor_scale, qdata, block_scale = packed output_shape = (tensor.shape[0], tensor.shape[1] // 2) return QuantizedTensor( qdata, diff --git a/src/h3_blackwell_runtime/runtime.py b/src/h3_blackwell_runtime/runtime.py index c4a89aa..3b0e4e6 100644 --- a/src/h3_blackwell_runtime/runtime.py +++ b/src/h3_blackwell_runtime/runtime.py @@ -154,9 +154,14 @@ class H3HotRuntime: def _load_h3(self) -> H3PackedDenoiser: model = H3PackedDenoiser.from_checkpoint(self.checkpoint, attention_backend=self.config.attention).eval() configure_mlp_chunking(model, self.config.mlp_chunks, self.config.mlp_chunk_threshold) + from .fc2_lt import prepare_fc2_lt + + prepare_fc2_lt() return model def status(self) -> dict: + from .fc2_lt import fc2_lt_status + return { "ready": True, "initial_attention": self.config.attention, @@ -168,6 +173,7 @@ class H3HotRuntime: "vae_tile_size": self.config.vae_tile_size, "mlp_chunks": self.config.mlp_chunks, "mlp_chunk_threshold": self.config.mlp_chunk_threshold, + "fc2_lt": fc2_lt_status(), "loaded_at": self.loaded_at, "load_stages": self.load_stages, } diff --git a/tests/test_fc2_lt.py b/tests/test_fc2_lt.py new file mode 100644 index 0000000..35bc13a --- /dev/null +++ b/tests/test_fc2_lt.py @@ -0,0 +1,287 @@ +import os +import unittest +from types import SimpleNamespace +from unittest.mock import MagicMock, patch + +import torch + +from h3_blackwell_runtime.fc2_lt import ( + _VALIDATED, + _canonical_fc2_supported, + _fc2_lt_compatible_result, + _STATS, + fc2_lt_enabled, + fc2_lt_linear, + prepare_fc2_lt, +) +from h3_blackwell_runtime.nvfp4 import Nvfp4Linear +from h3_blackwell_runtime.runtime import H3HotRuntime + + +class Fc2LtContracts(unittest.TestCase): + def setUp(self): + _VALIDATED.clear() + _fc2_lt_compatible_result.cache_clear() + _STATS.update(attempts=0, successes=0, fallbacks=0) + + def test_flag_is_opt_in(self): + with patch.dict(os.environ, {}, clear=True): + self.assertFalse(fc2_lt_enabled()) + with patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "yes"}, clear=True): + self.assertTrue(fc2_lt_enabled()) + + def test_only_canonical_unsharded_fc2_is_eligible(self): + linear = SimpleNamespace( + role="h3_mlp_fc2", + in_features=14_336, + out_features=5_376, + output_dtype=torch.bfloat16, + bias=None, + ) + gate_up = SimpleNamespace( + shape=(37_810, 28_672), + device=torch.device("cuda", 0), + ) + qdata = SimpleNamespace(shape=(37_824, 7_168)) + with patch("torch.cuda.get_device_capability", return_value=(12, 1)): + self.assertTrue(_canonical_fc2_supported(linear, gate_up, qdata)) + linear.in_features //= 2 + self.assertFalse(_canonical_fc2_supported(linear, gate_up, qdata)) + + def test_prepare_loads_and_validates_extension(self): + extension = MagicMock() + extension.build_info.return_value = { + "cuda_version": 13_000, + "cublas_version": 130_100, + "cuda_runtime_version": 13_000, + "cublaslt_runtime_version": 130_000, + } + with ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch( + "h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", + return_value=(extension, None), + ), + ): + self.assertTrue(prepare_fc2_lt()) + extension.prepare.assert_called_once_with() + + def test_ineligible_shape_is_counted_as_fallback(self): + args = (MagicMock(), MagicMock(), MagicMock(), MagicMock(), MagicMock()) + with ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=False), + ): + self.assertIsNone(fc2_lt_linear(*args)) + self.assertEqual(_STATS["fallbacks"], 1) + + def test_h3_model_load_prepares_extension(self): + model = MagicMock() + model.eval.return_value = model + runtime = SimpleNamespace( + checkpoint=MagicMock(), + config=SimpleNamespace(attention="sage2", mlp_chunks=1, mlp_chunk_threshold=4096), + ) + with ( + patch( + "h3_blackwell_runtime.runtime.H3PackedDenoiser.from_checkpoint", + return_value=model, + ), + patch("h3_blackwell_runtime.runtime.configure_mlp_chunking") as configure, + patch("h3_blackwell_runtime.fc2_lt.prepare_fc2_lt") as prepare, + ): + self.assertIs(H3HotRuntime._load_h3(runtime), model) + configure.assert_called_once_with(model, 1, 4096) + prepare.assert_called_once_with() + + def test_extension_load_failure_falls_back_unless_strict(self): + args = (MagicMock(), MagicMock(), MagicMock(), MagicMock(), MagicMock()) + with ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch("h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", return_value=(None, RuntimeError("load"))), + ): + self.assertIsNone(fc2_lt_linear(*args)) + with ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1", "H3_NVFP4_FC2_LT_STRICT": "1"}, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch("h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", return_value=(None, RuntimeError("load"))), + ): + with self.assertRaisesRegex(RuntimeError, "failed to load"): + fc2_lt_linear(*args) + + def test_build_info_and_runtime_version_failures_fall_back(self): + extension = MagicMock() + extension.build_info.side_effect = RuntimeError("metadata") + args = (MagicMock(), MagicMock(), MagicMock(), MagicMock(), MagicMock()) + with ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch("h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", return_value=(extension, None)), + ): + self.assertIsNone(fc2_lt_linear(*args)) + + _fc2_lt_compatible_result.cache_clear() + extension.build_info.side_effect = None + extension.build_info.return_value = { + "cuda_version": 13_000, + "cublas_version": 130_100, + "cuda_runtime_version": 12_900, + "cublaslt_runtime_version": 130_000, + } + with ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch("h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", return_value=(extension, None)), + ): + self.assertIsNone(fc2_lt_linear(*args)) + + def test_algo_check_and_execution_failures_fall_back(self): + extension = MagicMock() + extension.build_info.return_value = { + "cuda_version": 13_000, + "cublas_version": 130_100, + "cuda_runtime_version": 13_000, + "cublaslt_runtime_version": 130_000, + } + linear = MagicMock() + linear.weight.shape = (5_376, 7_168) + gate_up = MagicMock() + gate_up.device.index = 0 + qdata = MagicMock() + qdata.shape = (37_824, 7_168) + args = (linear, gate_up, MagicMock(), qdata, MagicMock()) + common = ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch("h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", return_value=(extension, None)), + ) + extension.check.return_value = {"valid": False, "required_workspace_bytes": 0} + with common[0], common[1], common[2]: + self.assertIsNone(fc2_lt_linear(*args)) + + _VALIDATED.clear() + extension.check.return_value = {"valid": True, "required_workspace_bytes": 0} + extension.run.side_effect = RuntimeError("run") + with ( + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch("h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", return_value=(extension, None)), + patch("h3_blackwell_runtime.fc2_lt.torch.zeros_like", return_value=MagicMock()), + patch("h3_blackwell_runtime.fc2_lt.torch.empty", return_value=MagicMock()), + ): + self.assertIsNone(fc2_lt_linear(*args)) + + def test_strict_mode_raises_on_algo_check_and_execution_failures(self): + extension = MagicMock() + extension.build_info.return_value = { + "cuda_version": 13_000, + "cublas_version": 130_100, + "cuda_runtime_version": 13_000, + "cublaslt_runtime_version": 130_000, + } + linear = MagicMock() + linear.weight.shape = (5_376, 7_168) + gate_up = MagicMock() + gate_up.device.index = 0 + qdata = MagicMock() + qdata.shape = (37_824, 7_168) + args = (linear, gate_up, MagicMock(), qdata, MagicMock()) + environment = { + "H3_NVFP4_FC2_LT_SPLITK1": "1", + "H3_NVFP4_FC2_LT_STRICT": "1", + } + extension.check.return_value = {"valid": False, "required_workspace_bytes": 0} + with ( + patch.dict(os.environ, environment, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch( + "h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", + return_value=(extension, None), + ), + ): + with self.assertRaisesRegex(RuntimeError, "algorithm check rejected"): + fc2_lt_linear(*args) + + _VALIDATED.clear() + _fc2_lt_compatible_result.cache_clear() + extension.check.return_value = {"valid": True, "required_workspace_bytes": 0} + extension.run.side_effect = RuntimeError("run") + with ( + patch.dict(os.environ, environment, clear=True), + patch("h3_blackwell_runtime.fc2_lt._canonical_fc2_supported", return_value=True), + patch( + "h3_blackwell_runtime.fc2_lt._fc2_lt_extension_result", + return_value=(extension, None), + ), + patch("h3_blackwell_runtime.fc2_lt.torch.zeros_like", return_value=MagicMock()), + patch("h3_blackwell_runtime.fc2_lt.torch.empty", return_value=MagicMock()), + ): + with self.assertRaisesRegex(RuntimeError, "execution failed"): + fc2_lt_linear(*args) + + def test_forward_swiglu_uses_schedule_without_repacking(self): + linear = SimpleNamespace( + role="h3_mlp_fc2", + in_features=4, + out_features=3, + output_dtype=torch.bfloat16, + bias=None, + pre_quant_scale=None, + active_lora=None, + lora_strength=0.0, + full_precision_matrix_mult=False, + ) + gate_up = torch.randn(1, 8, dtype=torch.bfloat16) + packed = (MagicMock(), MagicMock(), MagicMock()) + expected = torch.randn(1, 3, dtype=torch.bfloat16) + with ( + torch.inference_mode(), + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch.object(torch.Tensor, "is_cuda", new_callable=unittest.mock.PropertyMock, return_value=True), + patch("h3_blackwell_runtime.nvfp4_quant.vortex_native_quantize_swiglu_nvfp4", return_value=packed) as quantize, + patch("h3_blackwell_runtime.nvfp4_quant.wrap_native_swiglu_nvfp4") as wrap, + patch("h3_blackwell_runtime.fc2_lt.fc2_lt_linear", return_value=expected) as scheduled, + ): + actual = Nvfp4Linear.forward_swiglu(linear, gate_up) + self.assertTrue(torch.equal(actual, expected)) + quantize.assert_called_once_with(gate_up) + scheduled.assert_called_once_with(linear, gate_up, *packed) + wrap.assert_not_called() + + def test_forward_swiglu_reuses_native_bytes_on_fallback(self): + packed_weight = MagicMock() + linear = SimpleNamespace( + role="h3_mlp_fc2", + in_features=4, + out_features=3, + output_dtype=torch.bfloat16, + bias=None, + pre_quant_scale=None, + active_lora=None, + lora_strength=0.0, + full_precision_matrix_mult=False, + _packed_weight=lambda: packed_weight, + ) + gate_up = torch.randn(1, 8, dtype=torch.bfloat16) + packed = (MagicMock(), MagicMock(), MagicMock()) + wrapped = MagicMock() + expected = torch.randn(1, 3, dtype=torch.bfloat16) + with ( + torch.inference_mode(), + patch.dict(os.environ, {"H3_NVFP4_FC2_LT_SPLITK1": "1"}, clear=True), + patch.object(torch.Tensor, "is_cuda", new_callable=unittest.mock.PropertyMock, return_value=True), + patch("h3_blackwell_runtime.nvfp4_quant.vortex_native_quantize_swiglu_nvfp4", return_value=packed) as quantize, + patch("h3_blackwell_runtime.nvfp4_quant.wrap_native_swiglu_nvfp4", return_value=wrapped) as wrap, + patch("h3_blackwell_runtime.fc2_lt.fc2_lt_linear", return_value=None), + patch("h3_blackwell_runtime.nvfp4.functional.linear", return_value=expected) as gemm, + ): + actual = Nvfp4Linear.forward_swiglu(linear, gate_up) + self.assertTrue(torch.equal(actual, expected)) + quantize.assert_called_once_with(gate_up) + wrap.assert_called_once_with(gate_up, packed) + gemm.assert_called_once_with(wrapped, packed_weight, None) + + +if __name__ == "__main__": + unittest.main() diff --git a/tools/benchmark_fc2_nvfp4_algorithms.py b/tools/benchmark_fc2_nvfp4_algorithms.py index 49de326..81b8cd8 100644 --- a/tools/benchmark_fc2_nvfp4_algorithms.py +++ b/tools/benchmark_fc2_nvfp4_algorithms.py @@ -27,7 +27,7 @@ from h3_blackwell_runtime.t2v import random_av_latents ROOT = Path(__file__).resolve().parents[1] -SOURCE = ROOT / "research" / "fc2_nvfp4_scheduling" / "fc2_nvfp4_lt.cpp" +SOURCE = ROOT / "src" / "h3_blackwell_runtime" / "csrc" / "fc2_nvfp4_lt.cpp" DEFAULT_BUDGETS = (0, 4 << 20, 8 << 20, 16 << 20, 32 << 20, 64 << 20) DEFAULT_BLOCKS = (0, 24, 49) ENV_PREFIXES = ("H3_", "COMFY_KITCHEN_", "CUDA_", "TORCH_") @@ -35,7 +35,15 @@ ENV_PREFIXES = ("H3_", "COMFY_KITCHEN_", "CUDA_", "TORCH_") def parse_args() -> argparse.Namespace: parser = argparse.ArgumentParser(description=__doc__) - parser.add_argument("--mode", choices=("compile", "characterize", "sweep", "selected", "profile", "block-gate", "trajectory"), default="characterize") + parser.add_argument( + "--mode", + choices=( + "compile", "characterize", "sweep", "selected", "profile", + "block-gate", "trajectory", "production-block-gate", + "production-trajectory", "shape-gate", + ), + default="characterize", + ) parser.add_argument("--model-path", default="/models/minimax_h3_fl2va_pruned_nvfp4.safetensors") parser.add_argument("--capture", type=Path, help="Directory containing input.pt, or an input .pt file") parser.add_argument("--output", type=Path, default=Path("/output/h3-blackwell-runtime/benchmarks/fc2-nvfp4-scheduling.json")) @@ -48,6 +56,10 @@ def parse_args() -> argparse.Namespace: parser.add_argument("--sampler-step", type=int, default=1) parser.add_argument("--seed", type=int, default=440420) parser.add_argument("--text-tokens", type=int, default=100) + parser.add_argument( + "--shape-text-tokens", type=int, nargs="+", + default=[1, 15, 32, 64, 99, 100, 128, 256, 512], + ) parser.add_argument("--blocks", type=int, nargs="+", default=list(DEFAULT_BLOCKS)) parser.add_argument("--probe-block", type=int, default=24) parser.add_argument("--workspace-budgets", type=int, nargs="+", default=list(DEFAULT_BUDGETS)) @@ -553,7 +565,109 @@ def block_gate(extension, model, block_inputs, block_outputs, gate_up, timesteps return rows -def trajectory_gate(extension, checkpoint, model, config, args): +def production_block_gate(model, block_inputs, block_outputs, timesteps, rotation, segments, args): + from h3_blackwell_runtime.fc2_lt import fc2_lt_status + + flag = "H3_NVFP4_FC2_LT_SPLITK1" + previous = os.environ.get(flag) + rows = [] + + def select(enabled: bool) -> None: + if enabled: + os.environ[flag] = "1" + else: + os.environ.pop(flag, None) + + try: + for index in DEFAULT_BLOCKS: + block = model.backbone.blocks[index] + adaln_values = tuple(value.detach().clone() for value in model.backbone.adaln[index](timesteps)) + + def baseline(): + select(False) + return block(block_inputs[index].detach().clone(), rotation, *adaln_values, segments) + + def candidate(): + select(True) + return block(block_inputs[index].detach().clone(), rotation, *adaln_values, segments) + + before = fc2_lt_status() + timing = benchmark_pair( + baseline, candidate, args.rounds, args.warmup, + block_outputs[index].shape[0], block_outputs[index].shape[1], 1, + report_dense=False, + ) + with torch.inference_mode(): + baseline_value = baseline().detach().clone() + candidate_value = candidate().detach().clone() + after = fc2_lt_status() + dispatch_delta = { + name: after[name] - before[name] + for name in ("attempts", "successes", "fallbacks") + } + row = { + "block": index, + "production_method": "block.mlp.fc2.forward_swiglu", + "candidate_vs_baseline": compare(candidate_value, baseline_value), + "baseline_vs_traversal": compare(baseline_value, block_outputs[index]), + "candidate_vs_traversal": compare(candidate_value, block_outputs[index]), + "timing": timing, + "dispatch_delta": dispatch_delta, + } + if not all( + row[name]["bf16_exact"] + for name in ("candidate_vs_baseline", "baseline_vs_traversal", "candidate_vs_traversal") + ): + raise RuntimeError(f"production FC2 schedule is not byte-exact in block {index}") + if dispatch_delta["successes"] == 0 or dispatch_delta["fallbacks"] != 0: + raise RuntimeError(f"production FC2 schedule did not dispatch cleanly in block {index}") + rows.append(row) + finally: + if previous is None: + os.environ.pop(flag, None) + else: + os.environ[flag] = previous + return rows + + +def shape_gate(extension, gate_up, fc2, config, args): + base_rows = gate_up.shape[0] - args.text_tokens + rows = [] + for text_tokens in args.shape_text_tokens: + logical_rows = base_rows + text_tokens + if logical_rows <= gate_up.shape[0]: + actual_gate_up = gate_up[:logical_rows].contiguous() + else: + extra = logical_rows - gate_up.shape[0] + actual_gate_up = torch.cat((gate_up, gate_up[:extra]), dim=0).contiguous() + packed = packed_boundary(actual_gate_up, fc2) + run_candidate, _output, workspace = candidate_runner( + extension, packed, fc2, config, args.workspace_bytes, + ) + timing = benchmark_pair( + comfy_packed_runner(packed, fc2, logical_rows), run_candidate, + args.rounds, args.warmup, logical_rows, fc2.out_features, fc2.in_features, + ) + checked = dict( + extension.check( + packed[1], packed[2], fc2.weight, fc2.weight_scale, config, + ) + ) + row = { + "text_tokens": text_tokens, + "logical_rows": logical_rows, + "packed_rows": packed[1].shape[0], + "parity": timing["parity"], + "timing": timing, + "checked": checked, + "supplied_workspace_bytes": workspace.numel(), + } + row["passed"] = row["parity"]["bf16_exact"] and bool(checked.get("valid")) + rows.append(row) + return rows + + +def trajectory_gate(extension, checkpoint, model, config, args, *, production_dispatch=False): packer = H3PromptPacker(checkpoint) torch.manual_seed(args.seed) video, audio, aligned_frames = random_av_latents( @@ -583,11 +697,31 @@ def trajectory_gate(extension, checkpoint, model, config, args): return output[:actual_gate_up.shape[0], :fc2.out_features] return types.MethodType(replacement, fc2) - candidates = [candidate_method(block.mlp.fc2) for block in model.backbone.blocks] + candidates = None if production_dispatch else [ + candidate_method(block.mlp.fc2) for block in model.backbone.blocks + ] + flag = "H3_NVFP4_FC2_LT_SPLITK1" + previous_flag = os.environ.get(flag) + if production_dispatch: + os.environ[flag] = "1" + from h3_blackwell_runtime.fc2_lt import prepare_fc2_lt + + if not prepare_fc2_lt(): + raise RuntimeError("production FC2 extension preparation failed") + if previous_flag is None: + os.environ.pop(flag, None) + else: + os.environ[flag] = previous_flag def run(candidate: bool): - for index, block in enumerate(model.backbone.blocks): - block.mlp.fc2.forward_swiglu = candidates[index] if candidate else originals[index] + if production_dispatch: + if candidate: + os.environ[flag] = "1" + else: + os.environ.pop(flag, None) + else: + for index, block in enumerate(model.backbone.blocks): + block.mlp.fc2.forward_swiglu = candidates[index] if candidate else originals[index] torch.cuda.synchronize() started = time.perf_counter() result = sample_video_res_multistep( @@ -607,10 +741,19 @@ def trajectory_gate(extension, checkpoint, model, config, args): try: with torch.inference_mode(): (reference_video, reference_audio), baseline_seconds = run(False) + if production_dispatch: + from h3_blackwell_runtime.fc2_lt import fc2_lt_status + dispatch_before = fc2_lt_status() (candidate_video, candidate_audio), candidate_seconds = run(True) + if production_dispatch: + dispatch_after = fc2_lt_status() finally: for block, original in zip(model.backbone.blocks, originals, strict=True): block.mlp.fc2.forward_swiglu = original + if previous_flag is None: + os.environ.pop(flag, None) + else: + os.environ[flag] = previous_flag video_parity = compare(candidate_video, reference_video) audio_parity = compare(candidate_audio, reference_audio) @@ -631,6 +774,17 @@ def trajectory_gate(extension, checkpoint, model, config, args): "accepted_swiglu_producer_preserved": True, "accepted_gate_and_residual_path_preserved": True, } + if production_dispatch: + result["production_dispatch"] = True + result["dispatch_delta"] = { + name: dispatch_after[name] - dispatch_before[name] + for name in ("attempts", "successes", "fallbacks") + } + if result["dispatch_delta"]["successes"] == 0 or result["dispatch_delta"]["fallbacks"] != 0: + raise RuntimeError( + f"production FC2 schedule did not dispatch cleanly in trajectory: " + f"{result['dispatch_delta']}" + ) if not result["bf16_exact"]: raise RuntimeError("FC2 library candidate trajectory is not byte-exact") return result @@ -684,7 +838,11 @@ def main() -> None: "explicit_split_k_checks": explicit, } - if args.mode in {"characterize", "sweep", "selected", "profile", "block-gate", "trajectory"}: + if args.mode in { + "characterize", "sweep", "selected", "profile", "block-gate", + "trajectory", "production-block-gate", "production-trajectory", + "shape-gate", + }: config = resolve_candidate(args, enumerated, explicit) result["selected"] = config if args.mode == "characterize": @@ -712,8 +870,20 @@ def main() -> None: result["profile"] = profile_one(extension, packed, fc2, config, gate_ups[probe].shape[0], args) elif args.mode == "block-gate": result["block_gate"] = block_gate(extension, model, block_inputs, block_outputs, gate_ups, timesteps, rotation, segments, config, args) + elif args.mode == "production-block-gate": + result["production_block_gate"] = production_block_gate( + model, block_inputs, block_outputs, timesteps, rotation, segments, args, + ) + elif args.mode == "shape-gate": + result["shape_gate"] = shape_gate( + extension, gate_ups[probe], fc2, config, args, + ) elif args.mode == "trajectory": result["trajectory"] = trajectory_gate(extension, checkpoint, model, config, args) + elif args.mode == "production-trajectory": + result["trajectory"] = trajectory_gate( + extension, checkpoint, model, config, args, production_dispatch=True, + ) result["errors_and_unsupported"] = errors + [{ "feature": "Stream-K",