NVIDIA / NVIDIA/cutlass

pycute left_inverse returns a different canonical layout than CuTe C++

Open
#3,485 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

CUTLASS C++ duplicate
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Which component has the problem?

pycute (python/pycute, the pure-Python reference implementation), not the CuTe DSL

Describe the bug

pycute.left_inverse returns a different canonical layout than cute::left_inverse in include/cute/layout.hpp for non-contiguous and hierarchical inputs. Both results are valid left inverses on the image of the layout (applying them to layout(i) recovers i), so this is a fidelity gap between the Python reference and the C++ implementation rather than a broken inverse. Anyone using pycute to cross-check CuTe semantics gets different shapes/strides exactly in the interesting cases.

python/pycute/layout.py implements

def left_inverse(layout):
  ...
  return right_inverse(make_layout(layout, complement(layout)))

while the C++ side uses a dedicated sort-by-stride algorithm (left_inverse at include/cute/layout.hpp, with a divisibility static assert and an explicit stride-0 skip).

Steps to reproduce

Python (main):

from pycute import *
print(left_inverse(Layout(4,2)))
print(left_inverse(Layout((4,2),(1,16))))
print(left_inverse(Layout(((2,2),(2,2)),((1,4),(8,32)))))

C++ (host-only, g++ against include/):

#include <cute/tensor.hpp>
#include <iostream>
using namespace cute;
int main() {
  print(left_inverse(Layout<Shape<_4>,        Stride<_2>>{}));        print("\n");
  print(left_inverse(Layout<Shape<_4,_2>,     Stride<_1,_16>>{}));    print("\n");
  print(left_inverse(Layout<Shape<Shape<_2,_2>,Shape<_2,_2>>,
                        Stride<Stride<_1,_4>,Stride<_8,_32>>>{}));    print("\n");
}

Comparison gathered on main:

input C++ left_inverse pycute left_inverse
4:1 _4:_1 4:1 (agree)
4:2 (_2,_4):(_0,_1) (2,4):(4,1)
1:0 _1:_0 1:0 (agree)
(8,4):(1,8) _32:_1 32:1 (agree)
(8,4):(4,1) (_4,_8):(_8,_1) (4,8):(8,1) (agree)
(2,4,6):(1,2,8) _48:_1 48:1 (agree)
(2,4,6):(4,1,8) (_4,_2,_6):(_2,_1,_8) (4,2,6):(2,1,8) (agree)
(4,2):(1,16) (_16,_2):(_1,_4) (4,4,2):(1,8,4)
(2,5):(5,1) (_5,_2):(_2,_1) (5,2):(2,1) (agree)
((2,2),(2,2)):((1,4),(8,32)) (_4,_8,_2):(_1,_2,_8) (2,2,4,2,2):(1,16,2,32,8)

The C++ output for 4:2 contains an explicit leading stride-0 mode that the Python result lacks entirely, and the hierarchical case differs in both rank and strides.

Expected behavior

pycute.left_inverse should return the same canonical layout as cute::left_inverse, matching shape and stride structure, so the Python reference can be used to verify CuTe behavior. Alternatively, if the two algorithms are intentionally allowed to differ, the divergence should be documented in the pycute sources.

Environment details
  • Bare metal Windows host; C++ side compiled host-only with g++ against include/cute/layout.hpp from current main
  • Pure Python comparison needs no GPU

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start with left_inverse in python/pycute/layout.py and the corresponding implementation in include/cute/layout.hpp, then run the Python and host-only C++ reproductions listed in the issue. Compare the canonical shape and stride structures for the non-contiguous and hierarchical cases, including the explicit stride-0 mode. Done means pycute matches the C++ outputs, or the intentional divergence is documented in the pycute sources.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, python
Domain
backend
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Active
Clarity
Clearly specified
Newbie friendliness
64/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.