NVIDIA CUTLASS #3448 — Why One Tensor Descriptor Could Not Mix 32- and 64-Bit Dynamic Dimensions

CUTLASS CuTeDSL C Export, Part 2 of 2 — Homogeneous shape arrays, shared tensor-list typedefs, and why the exporter rejected an ABI it could not represent

Part 1 examined a width mismatch in CuTeDSL’s generated C interface.

A tensor declared with:

cute.sym_int64()

produced a compiled calling convention that expected:

i64

for each dynamic shape dimension.

But export_to_c() generated:

int32_t dynamic_shapes[...];

for every tensor.

The generated C caller and the compiled function therefore disagreed about the byte layout of the same descriptor.

PR #3448 repaired the uniform-width cases.

All dynamic shape dimensions are 32-bit
→ emit `int32_t`

All dynamic shape dimensions are 64-bit
→ emit `int64_t`

That still left a second question.

What should the exporter do with this tensor?

shape = (
    cute.sym_int32(),
    cute.sym_int64(),
)

The lower symbolic model can preserve the two dimensions independently.

dimension 0
→ 32-bit

dimension 1
→ 64-bit

But the exported C descriptor has only one shape-array declaration:

T dynamic_shapes[2];

The C language requires both array elements to have the same type T.

There is no choice of T that preserves:

slot 0
→ i32

slot 1
→ i64

Choosing int32_t misrepresents the second dimension.

Choosing int64_t misrepresents the first.

PR #3448 therefore did not silently normalize the widths.

It rejected the tensor during C-header generation.

The same rule was extended to tensor lists and tuples, because every element in one exported list shares a single tensor-descriptor typedef.

When an export representation cannot preserve the compiled ABI exactly, rejecting the export is safer than generating a different ABI that merely looks compatible.

The compiled metadata can describe each dimension separately

CuTeDSL’s tensor metadata does not represent the complete shape as one untyped integer array.

It builds one dimension record for each logical shape position.

For a dynamic dimension:

`SymInt`
→ preserve the declared width

anonymous dynamic dimension
→ use 32 bits

A tensor such as:

(
    cute.sym_int32(),
    cute.sym_int64(),
)

can therefore be represented conceptually as:

Dim 0
→ runtime symbol
→ 32 bits

Dim 1
→ runtime symbol
→ 64 bits

The lower metadata path retains those widths independently.

The limitation was not in the symbolic tensor model.

It was in the particular C structure selected by export_to_c().

The exported descriptor uses one homogeneous array

The common tensor-descriptor generator emits dynamic shape slots through one declaration.

Conceptually:

typedef struct {
    void* data;

    shape_type
        dynamic_shapes[n_dynamic_shapes];

    stride_type
        dynamic_strides[n_dynamic_strides];
} TensorDescriptor;

shape_type is one C type string.

`int32_t`

or

`int64_t`

It is not a list of one type per dimension.

The generated representation therefore supports:

[i32, i32, i32]

or:

[i64, i64, i64]

It does not support:

[i32, i64, i32]

without changing the descriptor format itself.

A C array cannot change element type by index

Consider:

int32_t dynamic_shapes[2];

Both fields are 32-bit.

dynamic_shapes[0]
→ int32_t

dynamic_shapes[1]
→ int32_t

Now consider:

int64_t dynamic_shapes[2];

Both fields are 64-bit.

dynamic_shapes[0]
→ int64_t

dynamic_shapes[1]
→ int64_t

C does not allow:

mixed_integer_t dynamic_shapes[2];

where:

element 0
→ int32_t

element 1
→ int64_t

A heterogeneous representation would need a different structure, for example:

typedef struct {
    int32_t dim0;
    int64_t dim1;
} DynamicShape;

But that field layout depends on the exact width sequence of one tensor.

The current generated descriptor is intentionally generic:

T dynamic_shapes[N];

PR #3448 did not redesign it into a per-shape custom structure.

Choosing the widest type would still be an ABI mismatch

One possible reaction is:

If any dimension is 64-bit,
make the whole array `int64_t`.

That avoids truncating the 64-bit dimension.

It does not preserve the compiled ABI of the 32-bit dimension.

Suppose the compiled function expects:

[i32, i64]

and the C header emits:

[i64, i64]

The first slot is already wrong.

Compiled slot 0:
4-byte integer

Generated slot 0:
8-byte integer

The next field does not begin at the same representation boundary.

A raw descriptor passed through void** is not converted field by field.

The compiled wrapper interprets the caller’s bytes using the exact LLVM structure it was compiled for.

wider source field
≠
automatic ABI conversion

An ABI is not preserved merely because every numerical value can fit inside the wider type.

Choosing the narrowest type would also be wrong

The opposite policy is:

Use `int32_t`
because it preserves the narrow dimension
and most tensor shapes are small.

That produces:

[i32, i32]

for a compiled contract expecting:

[i32, i64]

The second dimension is now represented with the wrong width.

Even when its current numerical value fits in 32 bits, the field boundaries still differ.

value fits
≠
descriptor matches

And when the value exceeds the 32-bit range, the mismatch also becomes an ordinary numerical truncation problem.

The exporter could not solve the mismatch with padding alone

Another possible idea is:

Keep the first slot 32-bit,
insert padding,
then place a 64-bit second slot.

That could represent one particular width sequence:

[i32, padding, i64]

But the generated declaration would no longer be:

T dynamic_shapes[N];

It would need a custom struct whose fields depend on:

tensor rank,
which dimensions are dynamic,
and the width of each dynamic dimension.

For example:

typedef struct {
    int32_t dynamic_shape_0;
    int64_t dynamic_shape_1;
    int32_t dynamic_shape_2;
} ShapeFields;

That is a possible future ABI design.

It is not the representation implemented by CUTLASS’s current C exporter.

The patch correctly addressed the representation that actually exists rather than pretending it could encode per-slot heterogeneity.

The exporter now collects the set of dynamic widths

The added helper walks only dimensions marked dynamic.

Conceptually:

widths = set()

for each shape dimension:

    if the dimension is static:
        continue

    if the dimension is a SymInt:
        widths.add(dimension.width)
    else:
        widths.add(32)

The possible outcomes are:

widths = {}
→ no dynamic shape array will be emitted

widths = {32}
→ use `int32_t`

widths = {64}
→ use `int64_t`

widths = {32, 64}
→ reject

The merged implementation raises DSLRuntimeError when more than one dynamic shape width appears.

Static dimensions do not participate in the conflict

Consider:

shape = (
    128,
    cute.sym_int64(),
)

The first dimension is static.

It does not need a runtime descriptor slot.

The dynamic shape array contains only:

the 64-bit symbolic dimension

Therefore:

int64_t dynamic_shapes[1];

is sufficient.

The width-consistency check applies to runtime slots, not every conceptual dimension in the tensor.

Static dimension
→ embedded in compilation

Dynamic dimension
→ crosses the C ABI

Only the second category needs one representable exported type.

Anonymous dynamic dimensions still use the 32-bit rule

A dynamic dimension does not always originate from:

cute.sym_int32()

or:

cute.sym_int64()

A DLPack-backed tensor can be marked dynamically laid out without explicit SymInt dimensions.

For those anonymous dynamic shapes, the lower metadata path uses 32 bits.

The C exporter mirrors that rule.

Anonymous dynamic shape
→ 32-bit metadata symbol
→ `int32_t` C slot

The patch therefore does not interpret:

dynamic

as synonymous with:

64-bit

The width is derived from the calling-convention category that produced the dimension.

Why rejecting the tensor is fail-closed behavior

The exporter had three broad choices.

Choice 1 — Force every slot to 32 bits

One compiled 64-bit slot is misrepresented.

Choice 2 — Force every slot to 64 bits

One compiled 32-bit slot is misrepresented.

Choice 3 — Reject the export

No C header is emitted for a descriptor
the current representation cannot express.

The third choice is the only one that does not silently publish an incompatible ABI.

The error is raised during header generation, before an external C application compiles against the wrong structure.

Unsupported representation

→ explicit export error

instead of

→ apparently valid header with wrong byte layout

Rejection preserves the original compiled contract

The patch does not rewrite:

cute.sym_int32()

into:

cute.sym_int64()

Nor does it narrow the 64-bit dimension.

The compiled function remains exactly what the source requested.

32-bit symbol
→ remains 32-bit

64-bit symbol
→ remains 64-bit

Only the C export is refused because its current homogeneous-array representation cannot expose that exact sequence.

This preserves a clear ownership boundary.

Compiler
→ honors the declared symbolic widths

C exporter
→ exports only layouts it can represent exactly

A future representation could support mixed widths

The rejection is not a claim that mixed-width shapes are inherently invalid.

The lower metadata already demonstrates that they can exist.

A future exporter might represent them through:

one named field per dynamic dimension,
a packed byte buffer plus explicit width metadata,
separate 32-bit and 64-bit arrays,
a generated per-signature native structure,
or a canonical descriptor using one width after changing the compiled ABI too.

Each option would require both sides to agree on a new representation.

PR #3448 does not introduce such a representation.

Its contract is narrower:

The current single-array C descriptor supports one dynamic shape width per tensor.

Tensor lists introduce another shared-type boundary

CuTeDSL also exports function arguments shaped like:

list[cute.Tensor]

or:

tuple[cute.Tensor, ...]

The C header creates one element descriptor typedef:

typedef struct {
    void* data;
    T dynamic_shapes[N];
    S dynamic_strides[M];
} k_Tensor_mAs_t;

It then creates a list wrapper:

typedef struct {
    k_Tensor_mAs_t elements[COUNT];
    int32_t count;
} k_TensorList_mAs_t;

Every element in that array has the same C type:

element 0
→ k_Tensor_mAs_t

element 1
→ k_Tensor_mAs_t

element 2
→ k_Tensor_mAs_t

Therefore, all tensors in the list must share the descriptor layout represented by that typedef.

One list cannot contain two descriptor layouts under one typedef

Suppose the Python argument contains:

Tensor 0

dynamic shape width
→ 64 bits

Tensor 1

dynamic shape width
→ 32 bits

The generated list has only one element type.

If that type declares:

int64_t dynamic_shapes[2];

Tensor 1 no longer matches its compiled 32-bit descriptor.

If it declares:

int32_t dynamic_shapes[2];

Tensor 0 no longer matches its compiled 64-bit descriptor.

The problem is the same as mixed widths inside one tensor, moved one level higher.

Within one tensor:

one shape array
→ one element type

Within one tensor list:

one descriptor typedef
→ one shape-array type for every element

The list wrapper passes each element separately

The generated wrapper packs:

&list->elements[0]

&list->elements[1]

...

into the generic argument array.

That does not make the element layouts independent.

Each address still points to memory constructed using the same C typedef.

Different pointer value

≠

different descriptor type

The external function can pass several descriptors.

The header still defines all of them through one shared element layout.

The old consistency check compared masks and stride policy

Before #3448, the list exporter already required every tensor element to share:

dynamic_shapes_mask

dynamic_strides_mask

_use_32bit_stride

Those checks protected several structural facts.

Dynamic shape mask

Which logical dimensions appear in the runtime shape array?

Dynamic stride mask

Which logical strides appear in the runtime stride array?

Stride-width policy

Are dynamic stride slots 32-bit or 64-bit?

These checks were necessary.

They were not sufficient.

Two tensors could have the same dynamic shape mask while assigning different widths to those shape slots.

The same mask does not imply the same field type

Consider:

Tensor A

dynamic_shapes_mask
→ [true, true]

shape widths
→ [64, 64]

Tensor B

dynamic_shapes_mask
→ [true, true]

shape widths
→ [32, 32]

The masks are identical.

Both tensors
→ have two runtime shape fields

But the field types are different.

Tensor A
→ two 64-bit fields

Tensor B
→ two 32-bit fields

A descriptor-consistency check that compares only masks can establish:

same field count and positions

It cannot establish:

same byte layout

PR #3448 adds dynamic shape width to the list-consistency contract.

The first tensor defines the shared list type

The merged code calculates:

shape_type =
    _get_dynamic_shape_c_type(first_tensor)

Then it checks every later element.

Does this tensor have the same dynamic shape mask?

Does it have the same dynamic stride mask?

Does it use the same stride-width policy?

Does it produce the same dynamic shape C type?

If any condition differs:

reject the list export

Only after all elements agree does the exporter generate one shared descriptor typedef.

Validate every element

        ↓

one common descriptor type

        ↓

list array of that descriptor

This is another validation-before-generation boundary

The pattern is similar to the Intel Triton RemoveMasks fix from the previous series.

Determine every object that will share one representation.

Validate the complete set.

Only then emit the shared structure.

For a tensor list:

All elements will use one C typedef.

Therefore:

Every element must satisfy that typedef’s layout.

Generating the typedef from the first tensor and assuming the rest are compatible would make the first element’s layout an accidental policy for the complete list.

The merged code checks the complete list before publishing the ABI.

One mismatch rejects the complete list export

Suppose:

Tensor 0
→ 64-bit shape slots

Tensor 1
→ 64-bit shape slots

Tensor 2
→ 32-bit shape slots

The exporter does not generate:

one 64-bit list type

and silently reinterpret Tensor 2

Nor does it split the Python list into two C list types.

The function signature requested one list argument.

The current C ABI represents that argument through one homogeneous descriptor array.

Therefore:

one incompatible element

→ complete list export rejected

That is the correct fail-closed result for the current representation.

Error messages describe the unsupported contract

For mixed widths inside one tensor, the helper raises a message equivalent to:

All dynamic shape dimensions of tensor argument <name>
must have the same width,
but 32 and 64 were found.

For a tensor list, the consistency error now states that every element must share:

dynamic shape mask

dynamic stride mask

stride-width policy

dynamic shape width

The diagnostic does not merely say:

C export failed.

It identifies the representational invariant that was violated.

That matters because the user can choose a supported alternative:

use consistently 32-bit symbols,
use consistently 64-bit symbols,
or redesign the argument boundary.

The tests cover both single-tensor and list rejection

The focused test file includes:

Mixed widths in one tensor

shape = (
    cute.sym_int32(),
    cute.sym_int64(),
)

Expected:

DSLRuntimeError

message mentions same-width requirement

Mismatched widths across list elements

element 0
→ sym_int64 dimensions

element 1
→ sym_int32 dimensions

Expected:

DSLRuntimeError

message mentions dynamic shape width

These are negative tests.

They prove that the exporter does not silently produce a header for an unrepresentable ABI.

The list success test protects the positive path

The test suite also creates a list whose elements all use:

cute.sym_int64()

Expected header:

int64_t dynamic_shapes[2];

That test failed before the patch because list descriptors still used the old hard-coded int32_t type.

After the repair:

every element agrees on 64-bit width

→ shared typedef uses `int64_t`

→ export succeeds

The pair of tests establishes both sides.

uniform list
→ accepted

heterogeneous-width list
→ rejected

A rejection test is part of ABI correctness

Tests are often written only for supported outputs.

Input A
→ expected header A

For an ABI generator, unsupported inputs also need an explicit contract.

Input cannot be represented exactly
→ generation must fail

Without the rejection test, a future change could reintroduce:

choose the first width

or

choose the widest width

and still generate syntactically valid C.

The negative test preserves:

unrepresentable
→ no ABI emitted

The eight-test matrix separated unchanged behavior from new behavior

The complete focused matrix was:

Case	Old implementation	Fixed implementation
Uniform `sym_int32` tensor	Pass	Pass
Uniform `sym_int64` tensor header	Fail	Pass
Compiled `sym_int64` ABI inspection	Pass	Pass
Fully static tensor	Pass	Pass
Anonymous DLPack dynamics	Pass	Pass
Mixed-width tensor rejection	Fail	Pass
Uniform 64-bit tensor list	Fail	Pass
Mismatched-width tensor list rejection	Fail	Pass

The old implementation did not fail because it correctly rejected mixed widths.

It failed the rejection tests because it silently emitted a 32-bit descriptor instead.

PR #3448’s summary reports four old-main failures and eight passing tests after the fix.

The unchanged tests matter

The successful pre-existing behavior was preserved.

sym_int32
→ still `int32_t`

static shape
→ still no dynamic shape array

anonymous DLPack dynamic shape
→ still `int32_t`

The patch did not globally widen descriptors or reject every dynamic tensor.

It changed only the cases where the existing exporter disagreed with the calling convention or could not represent it.

The structural test does not need a large kernel

The tests compile an empty kernel body.

That is sufficient because the contract under test exists before kernel arithmetic.

tensor signature

→ compiled descriptor

→ generated C typedef

A GEMM would not add evidence about whether:

T dynamic_shapes[N]

matches:

[i32, i64]

The minimal test isolates the host/compiler boundary.

The patch did not implement a runtime adapter

Another possible solution would be to keep one exported array width and have the wrapper convert values into the internal descriptor.

For example:

C header always uses 64-bit shape values

wrapper reads each value

wrapper narrows selected dimensions to 32 bits

wrapper builds the exact LLVM descriptor

The current wrapper does not do that.

It forwards descriptor pointers through void**.

PR #3448 repairs direct layout compatibility rather than adding a field-wise marshalling layer.

That distinction matters.

Direct-layout ABI
→ caller and callee structures must match

Marshalled ABI
→ wrapper could translate between representations

CUTLASS currently uses the first model for these tensor descriptors.

Direct layout is simple but strict

A direct-layout ABI has advantages.

No per-call descriptor conversion

Simple generated wrapper

Low host overhead

Easy C-side structure construction

Its constraint is equally direct.

Every exported byte

must already be in the form

the compiled function expects

The homogeneous-array limitation is a consequence of that representation.

Failing closed protects its simplicity from becoming silent misinterpretation.

Lists amplify a descriptor mistake

A wrong single tensor descriptor affects one argument.

A list uses the same typedef repeatedly.

wrong element layout

×
number of list elements

Every element can shift or misdescribe its dynamic fields in the same way.

This is why descriptor consistency cannot stop at:

same number of elements

same masks

The shared element type must be byte-identical for every list member.

Width is one part of descriptor identity

Two tensor descriptors are compatible only when the relevant layout facts agree.

For the current list export, those include:

number and positions of dynamic shapes

number and positions of dynamic strides

dynamic shape slot width

dynamic stride slot width

data-pointer field contract

The patch adds the shape-width dimension that was previously missing from the list check.

The broader rule is:

Two descriptors with the same logical field names are not the same ABI type when those fields have different widths.

Value conversion and ABI conversion are different

Python or C++ code can easily convert:

32-bit integer value
→ 64-bit integer value

That is a numerical operation.

A raw descriptor boundary does not perform that operation automatically.

memory containing i32

passed as a pointer

read as i64

is not value promotion.

It is reinterpretation of a different number of bytes.

The exporter must not confuse:

these values are convertible

with:

these structures are layout-compatible

The header generator is part of lowering

The dynamic width originates in a high-level symbolic declaration.

cute.sym_int32()

cute.sym_int64()

It then crosses several compiler layers.

Python symbolic type

→ metadata dimension

→ compiled descriptor field

→ C typedef field

→ external caller bytes

The C header generator is not outside the compiler pipeline.

It is the final lowering stage that translates the internal calling convention into a host-language representation.

A lossy decision there can invalidate all earlier type preservation.

Fail-closed export keeps unsupported ABI state visible

Silently choosing one width would hide the representational gap.

Header generated

C project compiles

runtime behavior depends on wrong layout

Rejecting the export leaves the gap visible where it originates.

Mixed-width symbolic tensor

→ current C descriptor cannot represent it

→ explicit error

That is useful both for correctness and future design.

A later implementation can add mixed-width support knowingly, with:

a new descriptor format,
explicit marshalling,
and matching tests.

It does not inherit a silent accidental convention.

What the patch directly changed for mixed-width handling

The merged patch directly:

gathers the widths of all dynamic shape dimensions,
rejects a tensor when that set contains both 32 and 64,
emits int32_t when all runtime shape slots are 32-bit,
emits int64_t when all runtime shape slots are 64-bit,
derives the shared tensor-list shape type from the first element,
checks every later element for the same shape width,
extends the list diagnostic to include dynamic shape width,
and adds positive and negative tests for both individual tensors and lists.

The current CUTLASS source retains this logic.

What the patch does not establish

The public evidence does not establish that:

mixed-width dynamic dimensions are invalid in CuTeDSL generally,
the lower compiled ABI cannot represent them,
a future C ABI could never support them,
every tensor-list representation must always be homogeneous,
a field-wise marshalling wrapper is impossible,
every host language has the same representation restriction,
the patch adds runtime conversion between integer widths,
or one published application previously failed numerically because of a mixed-width tensor.

The supported conclusion is narrower:

The current exported tensor descriptor stores all dynamic shape values in one homogeneous C array. It cannot represent mixed 32- and 64-bit ABI slots exactly, so the merged patch rejects those tensors and requires every element of one exported tensor list to agree on shape width.

The complete two-part series

Part 1 — Uniform 64-bit dimensions were exported as 32-bit

`sym_int64()`

→ compiled ABI uses `i64`

→ old C header used `int32_t`

→ descriptor layout mismatch

The patch derived the uniform C shape type from the symbolic width already present in the calling convention.

Part 2 — A homogeneous array could not preserve heterogeneous slots

dynamic dimension 0
→ i32

dynamic dimension 1
→ i64

generated representation
→ `T dynamic_shapes[2]`

No single T preserves both fields.

The exporter therefore rejects the tensor, and tensor-list elements must agree before sharing one typedef.

The final lesson is about representation honesty

A compiler may understand a richer type than one export format can express.

That does not justify silently reducing the richer type.

The safe choices are:

extend the export representation

or

reject the unsupported value

The unsafe choice is:

emit a different ABI
and hope the values happen to fit

When one C type is used to represent several compiled fields, those fields must actually share that type.

NVIDIA CUTLASS #3448 repaired the uniform-width mismatch and made the remaining representational limit explicit.

The header now says exactly what the compiled function reads—or refuses to say anything at all.

NVIDIA CUTLASS #3448 — Why the C Header Declared 32-Bit Shape Slots for a Kernel Reading 64-Bit Dimensions

Link the title above to Part 1 after publication.

Related material

Patch status: Merged into NVIDIA CUTLASS main
Affected component: CuTeDSL export_to_c()
Mixed-width input: Dynamic dimensions containing both 32- and 64-bit slots
Current C representation: One homogeneous dynamic_shapes array per tensor descriptor
Why widening was rejected: int64_t[] would still disagree with compiled i32 slots
Why narrowing was rejected: int32_t[] would disagree with compiled i64 slots
Single-tensor policy: Mixed dynamic shape widths raise DSLRuntimeError
Tensor-list policy: All elements must share masks, stride-width policy, and dynamic shape width
Positive list case: Uniform 64-bit elements export an int64_t descriptor
Negative list case: 32-/64-bit element mismatch is rejected
Validation: Eight focused tests; all passed after the patch
Evidence boundary: Export-layout and rejection behavior confirmed; no separate external-C runtime failure was published

This is Part 2 and the final article in the NVIDIA CUTLASS CuTeDSL C-export ABI series.

Part 1 examined why the generated C header and compiled _mlir_ciface assigned different byte widths to the same uniformly 64-bit dynamic dimensions.

This final article examined why the current homogeneous shape array cannot represent per-dimension width heterogeneity and why the exporter had to validate complete tensor and tensor-list layouts before emitting a shared C ABI.

#NVIDIA #CUTLASS #CuTeDSL #CUDA #CABI #SymInt #TensorDescriptor #CodeGeneration #AOTCompilation #FailClosed #CompilerCorrectness #CodeAnalysis