NVIDIA / NVIDIA/cutlass

[QST] The definition of zero stride in CUTE layout algebra

Open
#1,772 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

What is your question?

auto sA_layout = make_layout(make_shape(bM, bK));                 // (m,k) -> smem_idx; m-major  
ThrCopy thr_copy_a = copy_a.get_slice(threadIdx.x); 
Tensor sA = make_tensor(make_smem_ptr(smemA), sA_layout);            // (BLK_M,BLK_K) 
Tensor tAsA = thr_copy_a.partition_D(sA);    

From the above snippet from sgemm_2.cu, I got the following from partitioning sA via ThrCopy:

sA:
smem_ptr[32b](0x6ffff4000000) o (_128,_8):(_1,_128)

tAsA:
smem_ptr[32b](0x6ffff4000ff0) o ((_4,_1),_1,_1):((_1,_0),_0,_0)

How to interpret those _0 in tAsA? In the tutorial, there was a mentioning that _0 in the dimensions of tAsA indicates a dimension with size 1 where the exact stride is not crucial, allowing it to be any value. But if we follow the coordinate-to-index mapping equation (index = shapes .* strides), _0 could also be interpreted as that the tensor is being broadcasted along that dimension?

Will the meaning of _0 change depending on the print scenario or it always means ignore the stride in this dimension?

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 cited sgemm_2.cu snippet and the tutorial’s discussion of _0; inspect the printed sA and tAsA layouts and the partition_D call. Done means the repository documentation or issue response clearly states how zero strides are interpreted and whether that interpretation depends on the printing context.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
documentation
Issue type
Documentation
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.