NVIDIA / NVIDIA/cutlass

[QST] Tiled copy misaligned, how to solve it?

Open
#1,561 2 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

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

Description

My code:

using GmemTiledCopyL = decltype(make_tiled_copy(
      Copy_Atom<DefaultCopy, half>{}, Layout<Shape<_1>>{}, Layout<Shape<_1>>{}));
using SmemLayoutL = decltype(Layout<Shape<Int<64>>>{});
__shared__ cute::array_aligned<half, cute::cosize_v<SmemLayoutL>> l;
GmemTiledCopyL gmem_tiled_copy_L;

auto gmem_thr_copy_LD = gmem_tiled_copy_LD.get_thread_slice(tid);
Tensor _L = make_tensor(make_gmem_ptr(reinterpret_cast<const float*>(L),
                          make_shape(C), make_stride(Int<1>{}));
Tensor gL = local_tile(_L, make_tile(Int<64>{}),
                         make_coord(_));  // [64, C / 64]
Tensor sL = make_tensor(make_smem_ptr(shared_storage.l.data()),
                          SmemLayoutL{});  // [64]
Tensor gL_to_sL_src = gmem_thr_copy_LD.partition_S(gL);
Tensor gL_to_sL_dst = gmem_thr_copy_LD.partition_D(sL);
cute::copy(gmem_tiled_copy_L, gL_to_sL_src(_, _, 0), gL_to_sL_dst);

if my C is not divided by 8, then will throw error:
CUDA get error, code:716, msg:misaligned address
How to solve it? 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 by examining the tiled-copy, global-memory tile, and shared-memory layout assumptions shown in the issue, especially how C and the 64-element tile affect alignment. Reproduce the CUDA misaligned-address error with C not divisible by 8, then confirm the documented or accepted behavior for those dimensions and ensure the copy no longer faults.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc, performance
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
25/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.