[BUG] Cutlass python epilogue doesn't work with BF16
Open
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.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- 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