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

[Bug Fix] Support threads_per_head < 64 for wavefront size of 64 #6622

Merged
merged 11 commits into from
Nov 4, 2024
Merged

[Bug Fix] Support threads_per_head < 64 for wavefront size of 64 #6622

Changes from 8 commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
6d9300c
[Bug Fix] Support threads_per_head < 64 for wavefront size of 64
jagadish-amd Oct 11, 2024
90a036a
Merge branch 'master' into rope-kernel-bug
jagadish-amd Oct 14, 2024
5de1f78
Merge branch 'master' into rope-kernel-bug
jagadish-amd Oct 17, 2024
819fc34
Merge branch 'master' into rope-kernel-bug
jagadish-amd Oct 21, 2024
abdf7e0
Add support only for warpsize 64
jagadish-amd Oct 21, 2024
1a066b8
Merge branch 'master' into rope-kernel-bug
jagadish-amd Oct 22, 2024
0eff67f
Merge branch 'master' into rope-kernel-bug
loadams Oct 23, 2024
1c19fa8
Merge branch 'master' into rope-kernel-bug
jagadish-amd Nov 4, 2024
29800b0
Add test case for Rope
jagadish-amd Nov 4, 2024
33db959
Fix formatting issues
loadams Nov 4, 2024
d61ba53
Add license
loadams Nov 4, 2024
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
10 changes: 9 additions & 1 deletion csrc/transformer/inference/csrc/apply_rotary_pos_emb.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,7 +101,15 @@ __global__ void apply_rotary_pos_half(T* mixed_query,

#if defined(__HIP_PLATFORM_AMD__) and ROCM_WAVEFRONT_SIZE == 64
#define LAUNCH_FOR_ALIGNMENT(ALIGNMENT) \
if (threads_per_head == 64) { \
if (threads_per_head == 4) { \
LAUNCH_ROT_POS_EMB_HALF(4, ALIGNMENT); \
} else if (threads_per_head == 8) { \
LAUNCH_ROT_POS_EMB_HALF(8, ALIGNMENT); \
} else if (threads_per_head == 16) { \
LAUNCH_ROT_POS_EMB_HALF(16, ALIGNMENT); \
} else if (threads_per_head == 32) { \
LAUNCH_ROT_POS_EMB_HALF(32, ALIGNMENT); \
} else if (threads_per_head == 64) { \
LAUNCH_ROT_POS_EMB_HALF(64, ALIGNMENT); \
} else { \
assert(false); \
Expand Down