Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
33 commits
Select commit Hold shift + click to select a range
24f2c83
Add ROCm GLM routed IQ2 down support
Jul 22, 2026
d7a8909
Fix ROCm distributed GLM slice residency
Jul 22, 2026
84d2de4
Allow resident ROCm GLM distributed slices
Jul 22, 2026
fd6d056
Add non-invasive ROCm graph dumps
Jul 22, 2026
34d8046
Add GLM batch FFN stage dumps
Jul 22, 2026
784bc77
Add resident IQ2 MoE sorted-path isolation
Jul 22, 2026
507fd7c
Fix resident ROCm IQ2 routed MoE correctness
Jul 22, 2026
80afbe7
Fix ROCm IQ2 expert tile LUT initialization
Jul 22, 2026
2dc022d
Tune GLM resident ROCm slice memory reserve
Jul 22, 2026
7a424a4
Add grouped GLM value projection path
Jul 25, 2026
fa8b0b9
Optimize ROCm indexer top-k 2048
Jul 25, 2026
f6f5b92
Optimize GLM ROCm prefill projections
Jul 25, 2026
9463fd7
Rename DS4_ROCM_GLM_GROUPED_QK_LOW to DS4_ROCM_GLM_GROUPED_QK_LOW
Jul 25, 2026
26e8e88
Add ROCm GLM causal prefill GEMM path
Jul 25, 2026
c9e1013
Add optimized ROCm GLM slice decode path
Jul 25, 2026
0e609ab
Add F32 ROCm GLM causal attention GEMMs
Jul 25, 2026
b7f5cdd
remove FP32 hipblas path
Jul 25, 2026
312934c
Add opt-in ROCm GLM F16 compact cache
Jul 25, 2026
5e61ad5
Remove ineffective ROCm GLM F16 compact cache
Jul 26, 2026
5920d68
Add opt-in ROCm GLM performance experiments
Jul 26, 2026
41a3c79
Remove ineffective ROCm GLM QKV experiment
Jul 26, 2026
7a84671
Add opt-in GLM wave value projection
Jul 26, 2026
34b3862
Default GLM wave value projection
Jul 26, 2026
ea2e766
Default GLM selected attention GEMM
Jul 26, 2026
764987b
Record GLM selected attention deployment
Jul 26, 2026
dd837f0
Default GLM causal attention GEMM
Jul 26, 2026
cfb1ff9
Record GLM causal attention deployment
Jul 26, 2026
4d7fbbf
Add GLM selected attention head tiling
Jul 26, 2026
367a64c
Extend GLM selected attention head tiling
Jul 26, 2026
4e80dd7
Default GLM selected attention to tile 16
Jul 26, 2026
4ab7572
Add GLM indexed decode tuning path
Jul 26, 2026
4f11251
Remove ineffective GLM group8 decode experiment
Jul 26, 2026
2510636
Remove internal GLM validation artifacts
Jul 27, 2026
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
538 changes: 444 additions & 94 deletions ds4.c

Large diffs are not rendered by default.

4 changes: 2 additions & 2 deletions ds4_rocm.cu
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@
#include <sys/stat.h>
#include <time.h>
#include <unistd.h>
#include <algorithm>
#include <unordered_map>
#include <vector>

Expand All @@ -52,8 +53,7 @@ enum {
* decode calls to the online attention kernel so this fixed buffer never
* becomes an out-of-bounds write at long context. */
DS4_ROCM_ATTENTION_SCORE_CAP = 8192u,
DS4_ROCM_ATTENTION_RAW_SCORE_CAP = 256u,
DS4_ROCM_TOPK_MERGE_GROUP = 8u
DS4_ROCM_ATTENTION_RAW_SCORE_CAP = 256u
};

struct ds4_gpu_tensor {
Expand Down
Loading