feat(nvidia): add cutlass_scaled_mm operator

[voltjia]

Summary

• Resolve CUTLASS 4.4.1 through an existing NvidiaCutlass CMake package or a pinned, SHA256-verified FetchContent archive; no CUTLASS submodule is required.
• Add one representative public operator, cutlass_scaled_mm(a, b, scale_a, scale_b, out_dtype, bias, out), backed by a fused CUTLASS INT8 GEMM epilogue.
• Expose CutlassScaledMm in C++ and cutlass_scaled_mm in Python, matching vLLM rather than introducing a provider-neutral ScaledMm alias.
• Cover scalar/per-token scale_a, scalar/per-channel scale_b, required FP16/BF16 out_dtype, optional bias, padded layouts, non-default streams, and multi-GPU device restoration.

Motivation

This is the first intentionally narrow third-party provider migration slice from InfiniCore into InfiniOps. It validates the dependency, build, and provider structure with one representative CUTLASS operator before additional CUTLASS operators are proposed.

The public name, parameter order through bias, and W8A8 semantics follow vLLM's public cutlass_scaled_mm wrapper. vLLM allocates and returns the output tensor; InfiniOps follows its repository convention by accepting that tensor as a trailing explicit out and requires out.dtype == out_dtype.

The dependency is consumed as a private CMake target rather than a repository submodule. Source archives and sdists can therefore configure without Git submodule metadata, while offline builds can set standard CMake variables such as NvidiaCutlass_DIR or FETCHCONTENT_SOURCE_DIR_CUTLASS.

N/A - no linked issue.

API alignment

InfiniOps API	Alignment target	Evidence
cutlass_scaled_mm(a, b, scale_a, scale_b, out_dtype, bias, out)	vLLM cutlass_scaled_mm(a, b, scale_a, scale_b, out_dtype, bias=None) -> Tensor; the returned tensor is represented by trailing out under the InfiniOps execution convention	vLLM public API, registered low-level op, vLLM tests

This operator is not torch.nn.functional.scaled_mm. PyTorch's public function follows the ATen _scaled_mm_v2 FP8/FP4-oriented contract. This PR implements vLLM's W8A8 INT8 path: INT8 A/B, scalar or per-token A scales, scalar or per-channel B scales, required FP16/BF16 out_dtype, and optional bias. The generated PyTorch internal::_scaled_mm remains separate, and this PR does not modify scripts/torch_ops.yaml.

CUTLASS acquisition

Path	Behavior
Existing CMake package	Reuse exact NvidiaCutlass 4.4.1 target nvidia::cutlass::cutlass via NvidiaCutlass_ROOT, NvidiaCutlass_DIR, or CMAKE_PREFIX_PATH
Default fallback	Download immutable commit 087c84df with a fixed SHA256
Offline/source override	Use CMake's standard FETCHCONTENT_SOURCE_DIR_CUTLASS=/path/to/cutlass

The fallback only populates the source archive and exposes its headers through an interface target. It deliberately does not execute CUTLASS's top-level CMake project because CUTLASS 4.4.1 requires CMake 3.19 while InfiniOps supports CMake 3.18, and loading the subproject would also add unrelated install rules. This is the same dependency-lifecycle direction used by vLLM and TensorRT-LLM.

Type of Change

• feat - new feature / new operator / new platform
• fix - bug fix
• perf - performance improvement (no behavioral change)
• refactor - code restructuring without behavior change
• test - adding or fixing tests only
• docs - documentation only
• build / ci - build system or CI configuration
• chore - tooling, formatting, or other non-code changes
• Breaking change

Platforms Affected

• CPU (WITH_CPU)
• NVIDIA (WITH_NVIDIA)
• Iluvatar (WITH_ILUVATAR)
• MetaX (WITH_METAX)
• Cambricon (WITH_CAMBRICON)
• Moore (WITH_MOORE)
• Ascend (WITH_ASCEND)
• PyTorch C++ bindings (WITH_TORCH)
• Build system / CMake / CI
• Python bindings / user-facing API

Smoke Test Result

Remote environment: ssh nvidia, accelerator-dev/nvidia:latest (sha256:dd94fce2f83a180e2271f02f69713de5360aaf916b01e1f7f54173519fd54efd), NVIDIA A100-SXM4-80GB, PyTorch 2.10.0a0+b4e4ee81d3.nv25.12, CUDA 13.1.

# Source archive, standard local-source override, and smoke operator allowlist.
cmake -S /workspace/src -B /workspace/build -G Ninja \
  -DWITH_CPU=ON -DWITH_NVIDIA=ON -DWITH_TORCH=ON \
  -DGENERATE_PYTHON_BINDINGS=ON \
  -DINFINI_OPS_SMOKE_BUILD=ON \
  -DFETCHCONTENT_SOURCE_DIR_CUTLASS=/opt/cutlass \
  -DFETCHCONTENT_FULLY_DISCONNECTED=ON
cmake --build /workspace/build --target ops -j 16

python -m pytest tests/test_cutlass_scaled_mm.py \
  tests/test_generate_wrappers.py --devices nvidia -q
79 passed in 4.94s

python -m pytest tests --devices nvidia -m smoke -q
48 passed, 12 skipped, 6924 deselected in 12.62s

Test Results on Supported Platforms

Platform	Affected	Build / Smoke Result	Full Result / Notes
NVIDIA	Yes	Source-archive smoke build, focused tests, and main/shadow CI passed	Latest head f19b654
Iluvatar	No	N/A - not affected	N/A
MetaX	No	N/A - not affected	N/A
Cambricon	No	N/A - not affected	N/A
Moore	No	N/A - not affected	N/A
Ascend	No	N/A - not affected	N/A

Additional dependency validation

# Baseline before the provider change:
git archive source configuration failed because third_party/cutlass was absent.

# Current head, forced FetchContent fallback:
source archive configured and built the infiniops target successfully.

# Current head, invalid HTTP/HTTPS proxy plus local source override:
source archive configured and built the infiniops target successfully without downloading.

# Installed artifact inspection:
no include/cutlass or include/cute directories
no CUTLASS DT_NEEDED entry in libinfiniops.so

# GitHub Actions on f19b654:
format, documentation, main CI, and CI v2 shadow passed

Benchmark / Performance Impact

No performance claim in this first integration PR. Scale, bias, and output conversion are fused into the CUTLASS epilogue, so the implementation does not materialize an M x N INT32 accumulator or launch a second post-processing kernel. Architecture-specific tile tuning should be handled in a later performance-only PR.

Notes for Reviewers

• Scope is deliberately limited to CUTLASS plus one representative INT8 CutlassScaledMm; no FlashAttention, MATE, or MoE operator is included.
• The argument order through bias, required out_dtype, and W8A8 behavior follow vLLM's public wrapper; only the returned tensor is represented as trailing out under the InfiniOps execution convention.
• PyTorch torch.nn.functional.scaled_mm and generated internal::_scaled_mm remain separate because their _scaled_mm_v2 FP8/FP4 contract is not interchangeable with this W8A8 INT8 operator.
• CUTLASS is private, compile-time-only, and no longer requires a submodule checkout. The pinned fallback and standard local override are both covered by source-archive builds.
• CUTLASS templates live only in cutlass.cu; only that source receives SM75 SASS/PTX flags, so existing CUDA architecture configuration for other sources is unchanged.