NVIDIA / NVIDIA/cutlass

[QST] Gemm results are different with tile_description?

Open
#1,769 3 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?
It seems that add tile_description would make the gemm result different? assert (tensor_D_numpy - tensor_D).max() == 0.0 would pass if I add tile_decription.

import numpy as np
import random
import torch
import cutlass

# This controls whether the C++ GEMM declaration will be printed at each step. 
# Set to `False` to omit this information.
print_module = True

m = 8192
n = 8192
k = 8192
dtype=torch.float16
tensor_A = torch.rand(m, k, device=0, dtype=torch.float16)   
tensor_B = torch.rand(k, n, device=0, dtype=torch.float16)   
tensor_C = torch.zeros(m, n, device=0, dtype=torch.float16)   
tensor_D = torch.zeros(m, n, device=0, dtype=torch.float16)   

alpha = 1
beta = 0.0

# We specify `element_accumulator` here so as to match the kernel run by NumPy below. However,
# specifying `element_accumulator` is not required if it is the same as `element`
plan = cutlass.Gemm(element=dtype, layout=cutlass.LayoutType.RowMajor, element_accumulator=torch.float32)
"""
tile_description = {
    "threadblock_shape":  [128, 256, 32],   # Threadblock shape
    "stages": 3,                # Number of stages
    "wrap_count" : [2, 4, 1],        # Number of warps within each dimension of the threadblock shape
    "instruction_shape":  [16, 8 , 16] ,
    "cluster_shape":  [1, 1 , 1]
}
plan.tile_description = tile_description
"""

plan.run(tensor_A, tensor_B, tensor_C, tensor_D, print_module=print_module)

tensor_D_numpy = (alpha * (tensor_A @ tensor_B)) + (beta * tensor_C)


assert (tensor_D_numpy - tensor_D).max() == 0.0
print(tensor_D)

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 Python reproduction around cutlass.Gemm, plan.tile_description, and plan.run, comparing results with and without the tile description. Determine why the GEMM outputs differ from the NumPy reference and document the expected behavior or correction; the payload names no repository file or test to run.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, numpy, python, pytorch
Domain
hpc, machine-learning, performance
Issue type
Bug
Difficulty
3/5
Estimated time
1-2 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
30/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.