NVIDIA / NVIDIA/cutlass

[BUG] `Aborted (core dumped)` on incorrect copy layouts

Open
#2,319 3 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage bug inactive-30d inactive-90d
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Hi!

I came across a minor bug while using the new CuTe DSL. When incorrectly specifying a tv layout, CuTe DSL segfaults instead of giving an error.

Reproducer:

import torch
import cutlass
import cutlass.cute as cute
from cutlass.cute.runtime import from_dlpack

M = 64
device = torch.device("cuda")
A = torch.randn(M, M, dtype=torch.float32, device=device)
A_tiler = (M, M)

@cute.kernel
def test_kernel(
    mAT: cute.Tensor,
    sA_layout: cute.Layout,
    copy_A: cute.TiledCopy
):
    tidx, _, _ = cute.arch.thread_idx()

    gAT = cute.local_tile(mAT, tiler=A_tiler, coord=(1, 1))

    smem = cutlass.utils.SmemAllocator()
    sA = smem.allocate_tensor(mAT.element_type, sA_layout, byte_alignment=16)

    thr_copy_A = copy_A.get_slice(tidx)

    tAgAT = thr_copy_A.partition_S(gAT)
    tAsA = thr_copy_A.partition_D(sA)

    cute.copy(copy_A, tAgAT, tAsA)
    return


@cute.jit
def test_fn(mA: cute.Tensor):

    mAT_layout = cute.make_layout((M, M), stride=(1, M))
    mAT = cute.make_tensor(mA.iterator, layout=mAT_layout)

    sA_layout = cute.make_layout((M, M), stride=(M, 1))
    etype_mA = mA.element_type
    smem_size = cute.size_in_bytes(etype_mA, sA_layout)

    tA = cute.make_layout((16, 16), stride=(16, 1))
    vA = cute.make_layout((1, 1))

    atom_async_copy_A = cute.make_copy_atom(
        cute.nvgpu.cpasync.CopyG2SOp(),
        etype_mA, num_bits_per_copy=etype_mA.width,
    )

    # should be cute.make_tiled_copy_tv, "Aborted (core dumped)" as is
    tiled_copy_A = cute.make_tiled_copy(atom_async_copy_A, tA, vA)

    test_kernel(mAT, sA_layout, tiled_copy_A).launch(
        grid=(1, 1, 1),
        block=(256, 1, 1),
        smem=smem_size
    )

A_tensor = from_dlpack(A)

gemm = cute.compile(test_fn, A_tensor)
gemm(A)
torch.cuda.synchronize()

Here, if I use make_tiled_copy_tv (as I actually intended) it works. However, if I use make_tiled_copy it core dumps. I think that it would be nice if using make_tiled_copy, while incorrect, would produce an error message instead of an Aborted (core dumped). It's a bit tricky to debug core dumps 😅

Thanks!

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 the make_tiled_copy and make_tiled_copy_tv entry points shown in the reproducer, then run the provided CuTe DSL script to reproduce the abort. The work is done when the incorrect make_tiled_copy call reports an error instead of terminating with a core dump, while the intended make_tiled_copy_tv path continues to work.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, python
Domain
tooling
Issue type
Bug
Difficulty
3/5
Estimated time
1-2 days
Activity status
Stale
Clarity
Clearly specified
Newbie friendliness
48/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.