NVIDIA / NVIDIA/cutlass

[BUG] Cutlass python epilogue doesn't work with BF16

Open
#1,824 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

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

Description

Steps/Code to reproduce bug

import torch
import cutlass.epilogue


def epilogue(accum, bias):
    D = accum + bias
    return D

examples_tensors = dict(
    accum=torch.randn(1024, 1024),
    bias=torch.randn(1024, 1).bfloat16(),
    D=torch.randn(1024, 1024).bfloat16(),
)
cutlass.epilogue.trace(epilogue, examples_tensors)

Error

File ~/miniconda3/envs/dev_nightly/lib/python3.10/site-packages/cutlass/backend/evt/ir/load_nodes.py:187, in ColumnBroadcastImpl.argument_type.<locals>._Argument()
    184 class _Argument(ctypes.Structure):
    185     _fields_ = [
    186         ("ptr_col", ctypes.c_void_p),
--> 187         ("null_default", dtype2ctype[element_type]),
    188         ("dCol", tuple_type)
    189     ]
    190     def __init__(self, kwargs) -> None:
    191         ptr = kwargs[name]

KeyError: <DataType.bf16: 16>

Upon inspecting the code, I found that simply adding DataType.bf16: ctypes.c_uint16 to dtype2ctype will make things work. I verify that the outputs are correct.

https://github.com/NVIDIA/cutlass/blob/1ebda1ccef14df97da9ed098bd20e0c8520d6972/python/cutlass/backend/epilogue.py#L45-L51

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 python/cutlass/backend/epilogue.py at the linked dtype2ctype mapping, then inspect the ColumnBroadcastImpl.argument_type path in python/cutlass/backend/evt/ir/load_nodes.py. Re-run the Python reproduction with BF16 tensors and verify that tracing completes without the KeyError and produces correct outputs.

Written by the indexing model from the issue text.

Assessment

Tech stack
python
Domain
backend
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Stale
Clarity
Clearly specified
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.