Skip to content

[CuTeDSL] Compare missing runtime arguments by identity - #3634

Open
voipmonitor wants to merge 1 commit into
NVIDIA:mainfrom
voipmonitor:fix/runtime-argument-sentinel-identity
Open

voipmonitor wants to merge 1 commit into
NVIDIA:mainfrom
voipmonitor:fix/runtime-argument-sentinel-identity

Conversation

@voipmonitor

Copy link
Copy Markdown

Summary

Use object identity to detect unbound runtime-argument slots in
ExecutionArgs.get_rectified_args. The sentinel is a private object();
comparing supplied operand values with it is not part of argument binding.

self._missing in rectified invokes operand __eq__. This can raise for array
values, perform unnecessary tensor/DSL work, or report Missing required arguments: [] when a supplied object compares equal to everything. Identity
testing avoids these effects and agrees with the existing identity test used
to enumerate missing argument names.

The runtime change is one line. Positional/keyword binding, default values,
genuine missing-argument errors, argument packing and kernel execution are
otherwise unchanged. This does not add a prepared-launch API or cache argument
values.

Reproducer

import inspect
from cutlass.base_dsl.jit_executor import ExecutionArgs

class Operand:
    def __eq__(self, other):
        raise AssertionError("binding must not evaluate operand equality")

def signature(value):
    pass

operand = Operand()
binder = ExecutionArgs(inspect.signature(signature), "signature")
result = binder.get_rectified_args((operand,), {})
assert result[0] is operand

The equality-based implementation raises; the identity-based implementation
preserves the supplied object.

Validation

Status: implemented and host-qualified. Seven added unittest methods cover
positional/keyword/default binding, equality-raising and always-equal objects,
NumPy arrays, CuTe Int32 values, genuinely missing arguments, duplicate and
unexpected arguments, excess positional arguments and empty signatures.

The tests pass on Python 3.12 with CuTe DSL 4.6.2 and NumPy 2.1.0, using the
same one-line patch already deployed in a serving image. The executable AST of
the patched upstream binding method matches that deployed method exactly.
Extracting the method from base 147295a3d4b75f3aeff247c25b8927cea9a7006a
reproduces five failing and six erroring subcases; the patched source method
passes all tests with the same installed dependencies.

python test/python/CuTeDSL/test_execution_args.py -v
git diff --check

A full CuTe build at the upstream base was not performed: its complete
jit_executor.py requires adapter APIs absent from the installed package.
The source-method test isolates this pure-Python binding change; it does not
qualify unrelated compiler/runtime changes. No GPU tests or end-to-end speedup
claims are attached to this upstream submission.

Human review and AI assistance

Martin Vit (@voipmonitor) confirmed that he manually reviewed the one-line
runtime implementation and requested its upstream submission. PR preparation
and regression tests were assisted by OpenAI Codex.

Runtime argument binding detects its private missing-slot sentinel by identity
instead of invoking operand equality. This preserves supplied tensor-like
objects and avoids false missing-argument diagnostics. Binding, defaults and
error codes for genuinely invalid calls remain unchanged.

Seven host-only unittest methods cover equality-sensitive values and ordinary
binding/error behavior. The upstream method was tested in isolation with CuTe
DSL 4.6.2 native dependencies; no complete upstream compiler build was tested.

Martin Vit manually reviewed the one-line runtime change and authorized its
upstream submission. Preparation and tests were assisted by OpenAI Codex.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant