NVIDIA / NVIDIA/cutlass

About the make_tensor function

Open
#2,084 5 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

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

Description

`template <class ProblemShape, class CtaTiler,
class TA, class AStride,
class TB, class BStride,
class TC, class CStride>
global static
void
mak_tensor(ProblemShape shape_MNK, CtaTiler cta_tiler,
TA const* A, AStride dA,
TB const* B, BStride dB,
TC const* C, CStride dC
)
{
using namespace cute;

// Preconditions

//
// Full and Tiled Tensors
//

// Represent the full tensors
Tensor mA = make_tensor(make_gmem_ptr(A), select<0,2>(shape_MNK), dA); // (M,K)
Tensor mB = make_tensor(make_gmem_ptr(B), select<1,2>(shape_MNK), dB); // (N,K)
Tensor mC = make_tensor(make_gmem_ptr(C), select<0,1>(shape_MNK), dC); // (M,N)

// Get the appropriate blocks for this thread block
auto cta_coord = make_coord(blockIdx.x, blockIdx.y, _); // (m,n,k)
Tensor gA = local_tile(mA, cta_tiler, cta_coord, Step<_1, X,_1>{}); // (BLK_M,BLK_K,k)
Tensor gB = local_tile(mB, cta_tiler, cta_coord, Step< X,_1,_1>{}); // (BLK_N,BLK_K,k)
Tensor gC = local_tile(mC, cta_tiler, cta_coord, Step<_1,_1, X>{}); // (BLK_M,BLK_N)
#if 1
if(thread0()) {
print(" mB : "); print( mB); print("\n");
// print(" gB : "); print( gB); print("\n");
// print(" sB : "); print( (sB)); print("\n");
// print("tBgB : "); print(tBgB); print("\n");
// print("tBsB : "); print(tBsB); print("\n");
// print("tArA : "); print(tArA); print("\n");
}
#endif
// (BLK_N,BLK_K)
}`

cudaErrorLaunchFailure: unspecified launch failure

I don't know why make_tensor call inside kernel function must provide ld step parameter, otherwise it will report error, but call outside kernel function can not provide ld parameter, I don't know why, hope you can give me an answer, thank you

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_tensor calls in the provided kernel and compare them with the outside-kernel call, then trace the resulting CUDA launch failure. Check the tensor shape and stride arguments and use the reported cudaErrorLaunchFailure to narrow down the differing behavior. Done means explaining why the kernel call requires a leading-dimension stride and why the other call does not.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc
Issue type
Bug
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.