Skip to content

[Enhancement] Fallback transposed_ldmatrix into SM75_U16x4_LDSM_N when warp_n is 8 #498

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
merged 8 commits into from
May 17, 2025
Merged

[Enhancement] Fallback transposed_ldmatrix into SM75_U16x4_LDSM_N when warp_n is 8 #498

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
e82712f
Remove debug print statement from block_sparse_attn_triton.py and imp…
LeiWang1999 May 14, 2025
2417094
Enhance the autotuner module by adding a timeout handler for function…
LeiWang1999 May 14, 2025
e105d12
Merge branch 'main' of https://github.com/tile-ai/tilelang into HEAD
LeiWang1999 May 15, 2025
6eeeeaa
Add merge shared memory allocations pass and related configurations
LeiWang1999 May 16, 2025
aa8ab88
Reduce num_stages parameter in GEMM functions from 3 to 1 for improve…
LeiWang1999 May 16, 2025
73fb9a0
Update Copy type in OperandTraits for GEMM templates to use condition…
LeiWang1999 May 17, 2025
cd3fe03
Merge branch 'main' of https://github.com/tile-ai/tilelang into 0516_…
LeiWang1999 May 17, 2025
08f2418
lint fix
LeiWang1999 May 17, 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
1 change: 1 addition & 0 deletions examples/deepseek_mla/test_example_mla_decode.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@


@tilelang.testing.requires_cuda
@tilelang.testing.requires_cuda_compute_version_ge(9, 0)
def test_example_mla_decode():
with mock.patch.object(sys, 'argv', ["example_mla_decode.py"]):
example_mla_decode.main()
Expand Down
6 changes: 4 additions & 2 deletions src/tl_templates/cuda/gemm_sm80.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -98,7 +98,8 @@ struct OperandTraits<16, N, K, false, num_warp_n,
Swizzle<2, 3, 3>{}, Layout<Shape<_32, _8>, Stride<_1, _32>>{}));
using Layout = decltype(tile_to_shape(LayoutAtom{}, Shape<Int<N>, Int<K>>{},
Step<_2, _1>{}));
using Copy = SM75_U16x8_LDSM_T;
using Copy = typename std::conditional<N == 8 * num_warp_n, SM75_U16x4_LDSM_N,
SM75_U16x8_LDSM_N>::type;
};

template <int N, int K, int num_warp_n>
Expand All @@ -108,7 +109,8 @@ struct OperandTraits<16, N, K, false, num_warp_n,
Swizzle<3, 3, 3>{}, Layout<Shape<_64, _8>, Stride<_1, _64>>{}));
using Layout = decltype(tile_to_shape(LayoutAtom{}, Shape<Int<N>, Int<K>>{},
Step<_2, _1>{}));
using Copy = SM75_U16x8_LDSM_T;
using Copy = typename std::conditional<N == 8 * num_warp_n, SM75_U16x4_LDSM_N,
SM75_U16x8_LDSM_N>::type;
};

template <int N, int K, int num_warp_n>
Expand Down
6 changes: 4 additions & 2 deletions src/tl_templates/cuda/gemm_sm89.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -201,7 +201,8 @@ struct OperandTraits<16, N, K, false, num_warp_n,
Swizzle<2, 3, 3>{}, Layout<Shape<_32, _8>, Stride<_1, _32>>{}));
using Layout = decltype(tile_to_shape(LayoutAtom{}, Shape<Int<N>, Int<K>>{},
Step<_2, _1>{}));
using Copy = SM75_U16x8_LDSM_T;
using Copy = typename std::conditional<N == 8 * num_warp_n, SM75_U16x4_LDSM_N,
SM75_U16x8_LDSM_N>::type;
};

template <int N, int K, int num_warp_n>
Expand All @@ -211,7 +212,8 @@ struct OperandTraits<16, N, K, false, num_warp_n,
Swizzle<3, 3, 3>{}, Layout<Shape<_64, _8>, Stride<_1, _64>>{}));
using Layout = decltype(tile_to_shape(LayoutAtom{}, Shape<Int<N>, Int<K>>{},
Step<_2, _1>{}));
using Copy = SM75_U16x8_LDSM_T;
using Copy = typename std::conditional<N == 8 * num_warp_n, SM75_U16x4_LDSM_T,
SM75_U16x8_LDSM_T>::type;
};

template <int N, int K, int num_warp_n>
Expand Down
6 changes: 4 additions & 2 deletions src/tl_templates/cuda/gemm_sm90.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -255,7 +255,8 @@ struct OperandTraits<16, N, K, false, num_warp_n,
Swizzle<2, 3, 3>{}, Layout<Shape<_32, _8>, Stride<_1, _32>>{}));
using Layout = decltype(tile_to_shape(LayoutAtom{}, Shape<Int<N>, Int<K>>{},
Step<_2, _1>{}));
using Copy = SM75_U16x8_LDSM_T;
using Copy = typename std::conditional<N == 8 * num_warp_n, SM75_U16x4_LDSM_N,
SM75_U16x8_LDSM_N>::type;
};

template <int N, int K, int num_warp_n>
Expand All @@ -265,7 +266,8 @@ struct OperandTraits<16, N, K, false, num_warp_n,
Swizzle<3, 3, 3>{}, Layout<Shape<_64, _8>, Stride<_1, _64>>{}));
using Layout = decltype(tile_to_shape(LayoutAtom{}, Shape<Int<N>, Int<K>>{},
Step<_2, _1>{}));
using Copy = SM75_U16x8_LDSM_T;
using Copy = typename std::conditional<N == 8 * num_warp_n, SM75_U16x4_LDSM_N,
SM75_U16x8_LDSM_N>::type;
};

template <int N, int K, int num_warp_n>
Expand Down
Loading