[BUG] `Aborted (core dumped)` on incorrect copy layouts
Nobody has claimed this yet.
- 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
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- 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