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

CUTLASS CuTeDSL C Export, Part 1 of 2 — sym_int64(), _mlir_ciface, and a tensor descriptor whose generated C layout no longer matched the compiled ABI

The generated C code compiled successfully.

The compiled CuTeDSL function still expected a different tensor descriptor.

A tensor with two 64-bit symbolic dimensions produced an internal calling convention containing:

i64

i64

But export_to_c() generated a C header containing:

int32_t dynamic_shapes[2];

The Python compiler and the C caller were describing the same argument with different field widths.

CuTeDSL compilation

dynamic shape 0
→ 64 bits

dynamic shape 1
→ 64 bits

Generated C header

dynamic shape 0
→ 32 bits

dynamic shape 1
→ 32 bits

This was not merely a risk of truncating a very large dimension.

The two sides disagreed about the byte layout of the descriptor itself.

For two dynamic dimensions:

Generated C shape array

2 × 4 bytes
=
8 bytes

while the compiled function expected:

Compiled ABI shape payload

2 × 8 bytes
=
16 bytes

The C caller allocated only half the required byte width for the dynamic-shape portion of the descriptor.

Any fields following that array also began at different offsets from the ones expected by the compiled function.

NVIDIA CUTLASS PR #3448 fixed the mismatch by deriving the C shape type from the same symbolic-width information already used by the lower calling-convention path.

cute.sym_int32()
→ `int32_t` header slot
→ `i32` ABI slot

cute.sym_int64()
→ `int64_t` header slot
→ `i64` ABI slot

The patch was merged into CUTLASS main and added eight focused tests. Four of those tests failed against the old implementation and passed with the repair.

An exported C type is not documentation beside the ABI. It is part of the ABI.

What `export_to_c()` is responsible for

CuTeDSL can compile a Python-defined entry point and export a C-facing wrapper.

A simplified flow is:

Python CuTeDSL function

        ↓

compile()

        ↓

MLIR / LLVM calling convention

        ↓

export_to_c()

        ↓

generated C header and wrapper

        ↓

external C or C++ caller

The external caller does not construct CuTeDSL Python objects.

It follows the generated header.

For a tensor argument, that header defines a descriptor containing runtime information such as:

data pointer

dynamic shape dimensions

dynamic strides

A generated descriptor can look conceptually like:

typedef struct {
    void* data;

    int64_t dynamic_shapes[2];

    int64_t dynamic_strides[1];
} k_Tensor_mA_t;

The C caller fills this structure and passes its address to the generated wrapper.

The wrapper packs the argument addresses into:

void* args[];

and forwards them to the compiled C interface:

void _mlir_<prefix>__mlir_ciface_<function>(
    void** args,
    int32_t num_args
);

Because the final bridge uses void**, the C compiler does not compare the generated tensor typedef directly against the LLVM descriptor type.

The generated header is responsible for making those byte layouts agree.

The wrapper trusts the header’s descriptor

The generated wrapper receives a typed pointer such as:

k_Tensor_mA_t* mA

It then places that pointer into the generic argument array.

void* args[] = {
    mA,
    &stream,
    &ret
};

The compiled interface later interprets the memory at mA according to the descriptor type encoded in its calling convention.

The wrapper does not dynamically inspect:

How large is `dynamic_shapes`?

Are its entries 32-bit or 64-bit?

Do the C offsets match the LLVM struct offsets?

It assumes that the header generator and compiler agreed beforehand.

The boundary is therefore:

C header generator
→ declares bytes

C caller
→ writes bytes

compiled function
→ interprets bytes

A width mismatch can survive C compilation because all three layers are individually well-formed.

They simply do not describe the same object.

Static dimensions and dynamic dimensions travel differently

Consider a tensor shape:

[128, 256]

Both dimensions are static.

The compiler already knows them.

The generated runtime descriptor does not need shape slots for those values.

static dimension
→ embedded in compiled type or metadata
→ no runtime field required

Now consider:

[M, N]

where M and N are not known until the wrapper is called.

dynamic dimension
→ supplied by the caller
→ stored in the runtime descriptor

A descriptor therefore includes only the dimensions whose values must cross the ABI boundary.

The test added by #3448 verifies both cases:

fully static shape
→ no `dynamic_shapes` array

two dynamic dimensions
→ `dynamic_shapes[2]`

The number of slots was already correct.

The bug concerned the width of each slot.

`SymInt` carries an explicit width

CuTeDSL provides symbolic integer types including:

cute.sym_int32()

and:

cute.sym_int64()

They do not merely say:

This dimension is dynamic.

They also say:

This runtime dimension crosses the boundary
as a 32-bit or 64-bit integer.

For example:

a = make_fake_tensor(
    Float32,
    (
        cute.sym_int64(),
        cute.sym_int64(),
    ),
    stride=(
        cute.sym_int64(),
        1,
    ),
)

describes:

shape dimension 0
→ dynamic, 64-bit

shape dimension 1
→ dynamic, 64-bit

stride dimension 0
→ dynamic, 64-bit

stride dimension 1
→ static

The symbolic width is part of the function’s runtime calling convention.

It is not only an analysis hint.

The lower metadata path already preserved that width

The mismatch did not originate in the compiled tensor binding.

CUTLASS’s lower metadata conversion already distinguished the cases.

For every dynamic shape dimension:

dimension is `SymInt`
→ preserve `dim.width`

dimension is another anonymous dynamic value
→ use 32 bits

The relevant logic is effectively:

for each dynamic shape dimension:

    if isinstance(dim, SymInt):
        use dim.width
    else:
        use 32

The stride path similarly derives a runtime width from either the symbolic stride or the tensor’s 32-bit-stride policy.

This is why the compiled interface for two sym_int64() dimensions already contained:

struct<(i64, i64)>

The compiler side was not silently narrowing the dimensions.

The export layer was describing that existing ABI incorrectly.

The original header generator treated shape and stride differently

Before PR #3448, the C header generator already derived a stride type.

Conceptually:

stride_type = (
    "int32_t"
    if tensor_uses_32bit_stride
    else "int64_t"
)

But the shape type was passed as a constant:

shape_type = "int32_t"

The generated descriptor could therefore contain:

typedef struct {
    void* data;

    int32_t dynamic_shapes[2];

    int64_t dynamic_strides[1];
} Tensor;

even when the compiled function expected:

shape:
i64, i64

stride:
i64

The exporter knew that integer width mattered for one half of the descriptor.

It did not derive the same fact for the other half.

The pre-fix call was effectively:

_generate_tensor_element_typedef(
    ...,
    stride_type,
    shape_type="int32_t",
)

PR #3448 replaced the constant with a width derived from the tensor’s actual dynamic shape dimensions.

The tensor typedef is the on-wire layout

The shared C header generator emits the descriptor in a form equivalent to:

typedef struct {
    void* data;

    <shape_type>
        dynamic_shapes[n_dyn_shape];

    <stride_type>
        dynamic_strides[n_dyn_stride];
} TensorDescriptor;

The generator’s own comments describe this as the descriptor layout crossing the C boundary.

It deliberately uses fixed-width C types such as:

int32_t

int64_t

rather than target-dependent names such as:

long

int

size_t

because the exported structure must retain one stable ABI across build and target environments.

But using a fixed-width type is safe only when the selected width is the correct one.

Stable wrong width
≠
correct ABI

The minimal mismatch

The reproducer used two tensors that differed only in symbolic shape width.

32-bit symbolic shape

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

Compiled ABI:

struct<(i32, i32)>

Generated header:

int32_t dynamic_shapes[2];

The two sides agreed.

64-bit symbolic shape

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

Compiled ABI:

struct<(i64, i64)>

Old generated header:

int32_t dynamic_shapes[2];

The two sides disagreed.

The original issue reproduced this directly by printing both the _mlir_ciface type and the header declaration.

The mismatch exists even when the values are small

It may be tempting to think that the bug matters only when a dimension exceeds the signed 32-bit range.

For example:

N > 2,147,483,647

Such a value would indeed be truncated or unrepresentable in an int32_t slot.

But the descriptor mismatch occurs before that numerical limit matters.

Suppose the caller wants:

M = 128

N = 256

Both values fit easily in 32 bits.

The old C header writes:

shape slot 0
→ 4 bytes containing 128

shape slot 1
→ 4 bytes containing 256

The compiled interface expects:

shape slot 0
→ 8 bytes

shape slot 1
→ 8 bytes

The first 64-bit read spans bytes that the C header divided into two separate 32-bit fields.

The second 64-bit read begins where the C layout may already have placed the next descriptor field.

The problem is therefore:

field boundary mismatch

not only:

large integer truncation

A byte-level view

For two dynamic shapes and one dynamic stride, the old header’s payload can be drawn as:

offset 0
┌─────────────────────────────┐
│ data pointer       8 bytes  │
├─────────────────────────────┤
│ shape[0]           4 bytes  │
├─────────────────────────────┤
│ shape[1]           4 bytes  │
├─────────────────────────────┤
│ stride[0]          8 bytes  │
└─────────────────────────────┘

The compiled descriptor expected:

offset 0
┌─────────────────────────────┐
│ data pointer       8 bytes  │
├─────────────────────────────┤
│ shape[0]           8 bytes  │
├─────────────────────────────┤
│ shape[1]           8 bytes  │
├─────────────────────────────┤
│ stride[0]          8 bytes  │
└─────────────────────────────┘

The difference is not confined to one value.

The field following the shape array also moves.

Old C stride offset
→ after 8 bytes of shape data

Compiled stride offset
→ after 16 bytes of shape data

Once one field has the wrong width, every later interpretation can become displaced.

Exact padding can depend on the emitted structure and target ABI, but the core mismatch is unconditional:

8 bytes supplied for shape slots

16 bytes expected for shape slots

Why the C compiler could not diagnose it

A direct function declaration might allow a C compiler to compare two conflicting structure types.

That was not the interface here.

The generated wrapper declares a typed outer function for the caller, but the compiled bridge itself accepts:

void** args

The caller’s tensor pointer is reduced to one untyped entry in that array.

C compiler sees:

pointer to generated descriptor
→ converted into `void*`

compiled C interface sees:

`void*`
→ interpreted according to internal ABI metadata

There is no common C type shared between those two compilation units that can trigger a static layout error.

The generated header must be correct by construction.

The header and the compiled IR were two views of one contract

The new test pairs two assertions.

Header assertion

self.assertEqual(
    shape_slot_declaration,
    ("int64_t", "2"),
)

Compiled-interface assertion

self.assertIn(
    "struct<(i64, i64)>",
    ciface,
)

The second assertion confirms that the compiled ABI is genuinely 64-bit.

The first confirms that the exported C type now matches it.

This pairing is important.

A header test by itself could pass if both the compiler and exporter had incorrectly narrowed to 32 bits.

An ABI test by itself could pass while the exporter remained wrong.

Together:

compiled consumer width

=

exported producer width

The test file makes this contract explicit.

The repair derives shape width instead of assuming it

PR #3448 adds:

_get_dynamic_shape_c_type(
    tensor,
    argument_name,
)

The helper examines only the dimensions marked dynamic.

For every such dimension:

dimension is `SymInt`
→ use its declared width

dimension is another dynamic form
→ use 32 bits

The set of observed widths then determines the C type.

all dynamic widths are 64
→ `int64_t`

all dynamic widths are 32
→ `int32_t`

no dynamic dimensions
→ no shape array is emitted

The helper is now used when generating descriptors for both:

one tensor argument,
and a list or tuple of tensors.

The same logic remains present in current CUTLASS main.

The fix follows the lower calling convention

The patch did not create a new shape-width policy.

It copied the already-existing calling-convention rule into the C export layer.

Lower metadata path

SymInt
→ preserve width

non-SymInt dynamic dimension
→ 32 bits

C header path after fix

SymInt
→ preserve width

non-SymInt dynamic dimension
→ 32 bits

That is a stronger design than maintaining two independent hard-coded tables.

The producer and consumer derive their representation from the same semantic fact:

What width was assigned to this runtime dimension?

DLPack-dynamic tensors remain 32-bit

Not every dynamic dimension comes from:

cute.sym_int64()

A tensor imported through DLPack can be marked dynamically laid out without containing explicit SymInt objects for each shape dimension.

For those anonymous dynamic dimensions, the lower metadata path uses:

32 bits

The generated header continues to use:

int32_t dynamic_shapes[...];

The patch does not convert all dynamic shapes to 64 bits.

It preserves the distinction already represented by the calling convention.

The DLPack test passed before and after the patch.

The fix was not “always use `int64_t`”

A simple reaction could have been:

Dynamic dimensions can be large.

Therefore declare every slot as `int64_t`.

That would repair the sym_int64() case.

It would create a new mismatch for:

cute.sym_int32()

whose compiled ABI still expects:

i32

The correct rule is not:

largest available type

It is:

the exact type used by the calling convention

ABI compatibility is not monotonic.

A wider caller field is not automatically compatible with a narrower callee field.

Shape width and stride width remain separate decisions

The patch does not force:

shape width
=
stride width

A tensor can have:

64-bit symbolic shapes

and

a stride policy determined separately

The descriptor generator therefore receives:

shape_type

and:

stride_type

as independent arguments.

That separation reflects the compiled descriptor.

Runtime shape cells

and

runtime stride cells

can have different width contracts.

The bug was caused by treating only one of those contracts as data-driven.

Mixed-width dimensions expose a second limitation

The new helper can encounter:

one dynamic shape dimension
→ 32 bits

and:

another dynamic shape dimension
→ 64 bits

inside the same tensor.

The current exported C descriptor represents dynamic shapes through one homogeneous array:

<one_element_type>
    dynamic_shapes[count];

One C array cannot declare:

slot 0 as `int32_t`

and

slot 1 as `int64_t`

without changing the descriptor representation itself.

PR #3448 therefore rejects that configuration rather than silently choosing one width.

mixed 32-/64-bit dynamic shape dimensions

→ `DSLRuntimeError`

That fail-closed boundary—and why tensor lists need the same width contract—is the subject of Part 2.

The regression matrix contained eight cases

The added test file contains:

Case	Old implementation	Fixed implementation
`sym_int32` shape slots	Pass	Pass
`sym_int64` shape slots	Fail	Pass
`sym_int64` compiled ABI is `i64`	Pass	Pass
Fully static shape	Pass	Pass
DLPack dynamic shape	Pass	Pass
Mixed 32-/64-bit shape rejection	Fail	Pass
Tensor list with 64-bit shapes	Fail	Pass
Tensor list with mismatched widths	Fail	Pass

The PR summary states that four of the eight tests failed on the old main revision:

64-bit header slot width

mixed-width rejection

64-bit tensor-list descriptor

mismatched-width tensor-list rejection

The remaining cases passed before and after.

The passing ABI test isolated the exporter

One particularly useful case was:

test_sym_int64_shape_slots_match_abi

It checks that the compiled _mlir_ciface already contains:

struct<(i64, i64)>

That test passed on the old implementation.

The failure was therefore not:

CuTeDSL lost the symbolic width everywhere.

It was:

The compiled path preserved the width.

The C exporter did not.

This narrows the defect to the interface between two correct or partially correct representations.

Python symbolic type
→ correct

compiled ABI
→ correct

generated C typedef
→ incorrect

The patch validates structure, not an external C program

The public regression tests:

compile the CuTeDSL entry,
inspect the generated C header,
inspect the compiled _mlir_ciface type,
and check expected rejection cases.

They do not show a separate C application filling the old descriptor and producing one captured numerical failure or crash.

The evidence directly proves:

header field width
≠
compiled field width

The consequence for a conforming C caller follows from that structural mismatch.

But the public patch does not quantify:

which exact wrong shape values a particular caller observed,
whether one environment crashed,
or whether one kernel produced a numerical error before the repair.

The article should keep that boundary explicit.

The source-level bug was sufficient to invalidate the ABI

A separate runtime crash is not required to establish that the descriptor was wrong.

The ABI contract itself is byte-exact.

Caller writes 8 bytes

Callee expects 16 bytes

Those layouts are not compatible.

A particular run may appear to work because:

neighboring bytes happen to contain convenient values,
the kernel does not use every dynamic dimension,
the wrong value does not affect its empty body,
or the test entry launches a no-op kernel.

That does not make the interface valid.

No visible failure in one call

≠

matching ABI

The no-op test kernel was enough to inspect the contract

The regression kernel itself contains no arithmetic.

@cute.kernel
def _k(mA):
    pass

That is deliberate.

The test was not trying to show a tensor-computation error.

It was isolating:

function compilation

descriptor generation

header emission

calling-convention type

A larger GEMM or convolution would add unrelated behavior without strengthening the structural proof.

The smallest kernel made the mismatch easier to locate.

A generated header is compiler output

Developers sometimes treat a generated header as secondary material:

The real compiler output is the binary.

The header is only a convenience.

For an external C caller, the header is executable interface information.

It determines:

structure size,
field alignment,
field offsets,
integer width,
wrapper signature,
and argument packing.

A mistake in that file can invalidate otherwise correct generated code.

correct kernel binary

+

incorrect host descriptor

→ incorrect system

The host interface and device code form one compilation product.

Fixed-width types make the contract visible

The use of:

int32_t

and:

int64_t

is valuable because their size is explicit.

There is no ambiguity about whether:

long

means 32 or 64 bits on a given platform.

The bug was visible as:

int32_t
vs.
i64

The failure was not caused by platform-dependent integer naming.

It was caused by choosing the wrong explicit width.

This is a useful ABI-review principle:

Fixed-width types remove ambiguity. They do not remove the obligation to choose the correct width.

The descriptor is a producer-consumer contract

The complete data path is:

CuTe symbolic dimension

        ↓

lower metadata binding

        ↓

compiled LLVM descriptor

        ↓

C header generation

        ↓

C caller structure

        ↓

`void**` argument packing

        ↓

compiled interface interpretation

The same width must survive every transition.

Before #3448:

64
→ 64
→ 64
→ 32
→ 32
→ interpreted as 64

After #3448:

64
→ 64
→ 64
→ 64
→ 64
→ interpreted as 64

The repair did not change the kernel’s meaning.

It restored the descriptor’s identity across language boundaries.

What the patch directly fixed

PR #3448 directly:

imports the SymInt type into the C header generator,
derives dynamic shape width from each tensor’s declared shape dimensions,
emits int64_t for uniformly 64-bit symbolic shape slots,
retains int32_t for 32-bit or anonymous dynamic dimensions,
applies the same logic to individual tensors and tensor lists,
rejects mixed-width dynamic shape arrays,
extends tensor-list consistency checks to include shape width,
and adds eight focused tests.

The patch does not alter:

kernel arithmetic,
tensor element type,
pointer address space,
dynamic-stride width policy,
device launch dimensions,
or the compiled SymInt ABI itself.

The compiled ABI already retained the 64-bit width.

What the patch does not establish

The public evidence does not establish that:

every external C caller experienced a crash,
one specific kernel produced incorrect numerical output,
every descriptor field in CuTeDSL has now been audited,
arbitrary mixed-width shapes are supported,
changing the header affects device performance,
every host ABI lays out unrelated structures identically,
or tensor descriptors can freely mix per-slot C integer types.

The supported conclusion is narrower:

export_to_c() declared every dynamic shape slot as 32-bit even when the compiled CuTeDSL calling convention expected 64-bit SymInt dimensions. The merged patch derives the C field width from the same symbolic-width contract and rejects representations that the homogeneous C array cannot express.

The complete failure chain

The bug can be written as:

Tensor shape uses `cute.sym_int64()`
        ↓
lower tensor binding preserves 64-bit width
        ↓
compiled `_mlir_ciface` expects `i64` shape cells
        ↓
`export_to_c()` hard-codes `int32_t`
        ↓
C caller allocates 4 bytes per shape cell
        ↓
wrapper passes descriptor through `void**`
        ↓
compiled function interprets 8 bytes per cell
        ↓
caller and callee disagree about field boundaries

The fixed path is:

Inspect dynamic dimensions
        ↓
preserve each `SymInt` width
        ↓
require one representable array width
        ↓
emit matching fixed-width C type
        ↓
header and compiled descriptor agree

The final lesson is that the header cannot guess the ABI

The old exporter did not lack the information it needed.

The width was already present:

in the `SymInt`

in the lower metadata

in the compiled interface

The export path replaced that information with a constant.

Known 64-bit contract

→ overwritten by `int32_t`

The repair did not require a new runtime protocol.

It required the C header generator to stop inventing one.

When a compiler exports a descriptor across a language boundary, every field must be derived from the calling convention that will consume it.

NVIDIA CUTLASS #3448 restored that connection for dynamic tensor shapes.

Part 2: why one shape array could not contain two widths

The second and final article examines the cases the patch refuses.

shape dimension 0
→ `sym_int32()`

shape dimension 1
→ `sym_int64()`

The lower symbolic model can describe both widths.

The generated C descriptor has only:

T dynamic_shapes[N];

with one element type T.

Part 2 examines:

why choosing int32_t would truncate or misdescribe the 64-bit slot,
why choosing int64_t would still disagree with the 32-bit ABI slot,
why a list of tensors shares one descriptor typedef,
why masks and stride policy were not enough for list consistency,
and why the patch rejects mixed-width tensors instead of silently normalizing them.

Related material

Patch status: Merged into NVIDIA CUTLASS main
Affected component: CuTeDSL export_to_c()
Affected field: Tensor descriptor dynamic_shapes array
Old header behavior: Always emitted int32_t shape slots
Compiled ABI behavior: Preserved SymInt width, including i64
Minimal mismatch: sym_int64(), sym_int64() produced i64, i64 internally but int32_t[2] in C
Repair: Derive the C shape type from dynamic-dimension widths
32-bit behavior: Preserved for sym_int32() and anonymous DLPack dynamics
Mixed-width behavior: Rejected rather than silently normalized
Validation: Eight focused tests; four failed on the old implementation and passed after the fix
Runtime evidence boundary: Structural header/IR mismatch confirmed; no separate external-C numerical failure was published

This is Part 1 of a two-part series on NVIDIA CUTLASS CuTeDSL’s C-export ABI boundary.

This article examined how the compiled function and generated header assigned different byte widths to the same dynamic tensor dimensions, and why even small shape values could not make the two descriptor layouts compatible.

Part 2 examines why one homogeneous C array cannot represent mixed 32- and 64-bit ABI slots and why tensor-list elements must agree on dynamic shape width before they can share one exported descriptor typedef.

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