Skip to content

[Dev] Fix FP8 Codegen Issue for tilelang backend #291

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

Merged
LeiWang1999 merged 20 commits into from
Feb 4, 2025
Merged

[Dev] Fix FP8 Codegen Issue for tilelang backend #291

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
a1e7c81
Refactor: Rename tensor core related classes and imports to use tile …
LeiWang1999 Feb 2, 2025
566fc67
lint fix
LeiWang1999 Feb 2, 2025
563a704
Fix: Update Matmul initialization to specify backend and clean up imp…
LeiWang1999 Feb 2, 2025
d95fb42
lint fix
LeiWang1999 Feb 2, 2025
97ecf57
Update subproject commit for TVM dependency
LeiWang1999 Feb 2, 2025
fe44dd7
Refactor: Update imports to use tilelang instead of tvm.tl.language
LeiWang1999 Feb 3, 2025
7dc068c
Refactor: Clean up import statements and formatting in bitblas module
LeiWang1999 Feb 3, 2025
691a0dc
Fix: Add newline injection to .bashrc if the last line is not empty i…
LeiWang1999 Feb 3, 2025
fd4c1a6
Update submodule URLs and branches for TVM and TileLang
LeiWang1999 Feb 3, 2025
e7ad6a9
Update tilelang submodule URL and add new subproject commit
LeiWang1999 Feb 3, 2025
6da2f5f
Update cutlass submodule URL to point to tile-ai repository
LeiWang1999 Feb 3, 2025
30f1a95
Merge branch 'main' of https://github.com/microsoft/BitBLAS into sepa…
LeiWang1999 Feb 3, 2025
8f68896
Refactor: Split class definition for MatmulINT4DequantizeMMAWeightPro…
LeiWang1999 Feb 3, 2025
af0a134
Enhance environment variable handling for TVM and TileLang paths in i…
LeiWang1999 Feb 3, 2025
b9bb657
Remove unnecessary blank line in initialization of TILELANG_IMPORT_PA…
LeiWang1999 Feb 3, 2025
9f7d4c6
Add build_tilelang function to setup.py for TILELANG integration
LeiWang1999 Feb 4, 2025
1cc6886
Add TILELANG build step in setup.py
LeiWang1999 Feb 4, 2025
8a53cd4
Update TileLang subproject and improve type mappings in TLCUDASourceW…
LeiWang1999 Feb 4, 2025
3cac227
Refactor line continuation for better readability in gemv_simt.py
LeiWang1999 Feb 4, 2025
7ce4860
Merge branch 'main' of https://github.com/microsoft/BitBLAS into sepa…
LeiWang1999 Feb 4, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion 3rdparty/tilelang
Open in desktop
Submodule tilelang updated 3 files
+98 −30 src/target/codegen_cuda.cc
+25 −2 src/target/codegen_cuda.h
+23 −0 src/tl_templates/cuda/cuda_fp8.h
8 changes: 4 additions & 4 deletions bitblas/builder/wrapper/tl.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,10 +18,10 @@
class TLCUDASourceWrapper(object):
_TYPE_MAP = {
"float32": "float",
"float16": "half_t",
"bfloat16": "bfloat16_t",
"e4m3_float8": "float_e4m3_t",
"e5m2_float8": "float_e5m2_t",
"float16": "half",
"bfloat16": "__nv_bfloat16",
"e4m3_float8": "__nv_fp8_e4m3",
"e5m2_float8": "__nv_fp8_e5m2",
"float64": "double",
"int64": "int64_t",
"int32": "int",
Expand Down
3 changes: 2 additions & 1 deletion bitblas/ops/general_matmul/tilelang/dense/gemv_simt.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -158,7 +158,8 @@ def main(
)
else:
for ki in T.serial(micro_size_k):
accum_res[0] += A_local[ki] * B_local[ki]
accum_res[0] += A_local[ki].astype(accum_dtype) * B_local[ki].astype(
accum_dtype)

with T.attr(
T.comm_reducer(lambda x, y: x + y, [T.Cast(accum_dtype, 0)]),
Expand Down
4 changes: 2 additions & 2 deletions testing/python/operators/test_general_matmul_fp8.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@ def matmul_torch_forward(M, N, K, A_dtype, W_dtype, accum_dtype, out_dtype, layo
with_zeros=with_zeros,
zeros_mode=zeros_mode,
)
matmul = Matmul(config=matmul_config, enable_tuning=True, backend="tir")
matmul = Matmul(config=matmul_config, enable_tuning=True, backend="tl")

input_shape = (M, K)
weight_shape = (N, K) if layout == "nt" else (K, N)
Expand Down Expand Up @@ -93,7 +93,7 @@ def matmul_torch_forward_weight_dequantize(M, N, K, A_dtype, W_dtype, accum_dtyp
propagate_a=False,
propagate_b=False,
)
matmul = Matmul(config=matmul_config, enable_tuning=False, backend="tir")
matmul = Matmul(config=matmul_config, enable_tuning=False, backend="tl")
input_shape = (M, K)
weight_shape = (N, K) if layout == "nt" else (K, N)

Expand Down
Loading