Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

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

Open
gau-nernst opened this issue Sep 18, 2024 · 2 comments · May be fixed by #1832
Open

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

gau-nernst opened this issue Sep 18, 2024 · 2 comments · May be fixed by #1832
Labels
? - Needs Triage bug Something isn't working

Comments

@gau-nernst
Copy link

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.

cutlass/python/cutlass/backend/epilogue.py

Lines 45 to 51 in 1ebda1c

dtype2ctype = {
DataType.f16: ctypes.c_uint16,
DataType.f32: ctypes.c_float,
DataType.f64: ctypes.c_double,
DataType.s8: ctypes.c_int8,
DataType.s32: ctypes.c_int32
}

@gau-nernst gau-nernst added ? - Needs Triage bug Something isn't working labels Sep 18, 2024
@jackkosaian
Copy link
Contributor

Thanks. Would you be willing to contribute a PR to fix this?

@gau-nernst
Copy link
Author

Sure! I will do it tmr.

@gau-nernst gau-nernst linked a pull request Sep 20, 2024 that will close this issue
Open
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
? - Needs Triage bug Something isn't working
Projects
None yet
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

Add BF16 dtype mapping for python EVT gau-nernst/cutlass
2 participants
@jackkosaian @gau-nernst