Skip to content
Merged
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
7f6b327
feat(hipblaslt): add rocBLAS wvSpltK M=1 decode GEMV via GFA
jinchen62 Sep 17, 2026
e48a882
replace AssertFree0SizeEqual with AssertSizeEqual
jinchen62 Sep 18, 2026
505d398
fix(hipblaslt): restore TensileLite coverage for AssertSizeEqual
jinchen62 Sep 18, 2026
03d9a19
Merge remote-tracking branch 'upstream/develop' into gemv_kernel
jinchen62 Sep 21, 2026
8e22a46
correct out-of-place beta and unguarded shapes in wvSpltK
jinchen62 Sep 21, 2026
abff96d
add wvSpltK M=2/M=4 and fix out-of-place beta
jinchen62 Sep 21, 2026
f9e3952
Merge remote-tracking branch 'upstream/develop' into gemv_kernel
jinchen62 Sep 21, 2026
eff4671
Merge branch 'develop' into gemv_kernel
jinchen62 Sep 22, 2026
98dbf95
predicate the packed layout of the wvSpltK kernels
jinchen62 Sep 24, 2026
9183020
Merge remote-tracking branch 'upstream/develop' into gemv_kernel
jinchen62 Sep 24, 2026
77af803
Merge branch 'develop' into gemv_kernel
jinchen62 Sep 28, 2026
42d4acf
add kernel source
jinchen62 Sep 30, 2026
bcddc08
fix test
jinchen62 Sep 30, 2026
e0e5a86
serve M <= 4 from wvSpltK_hf_m4
jinchen62 Sep 30, 2026
919e02c
add wvSpltK BF16 TN kernels for n = 1, 2, 4 on gfx950
jinchen62 Sep 30, 2026
e553758
reject empty problems in the m4 wvSpltK kernels
jinchen62 Sep 30, 2026
d79264b
emit the wvSpltK kernels at code object v4
jinchen62 Sep 30, 2026
092d3ca
route skinny GEMMs to wvSpltK through Range logic
jinchen62 Sep 30, 2026
2be650d
drop the wvSpltK kernel sources
jinchen62 Sep 30, 2026
752f43c
Merge remote-tracking branch 'upstream/develop' into gemv_kernel
jinchen62 Sep 30, 2026
69b3db6
skip the wvSpltK range-logic test without the logic tree
jinchen62 Oct 1, 2026
724681b
Merge branch 'develop' into gemv_kernel
jinchen62 Oct 1, 2026
35801a5
Merge remote-tracking branch 'upstream/develop' into gemv_kernel
jinchen62 Oct 5, 2026
f892742
pass ldb to the FP16 NN wvSpltK kernels and rename them
jinchen62 Oct 5, 2026
379a7ba
skip custom kernels for pointer-array GEMMs in the heuristic
jinchen62 Oct 5, 2026
fd2f7a1
test the skinny GEMMs library logic routes to wvSpltK
jinchen62 Oct 5, 2026
61e25ac
cap the wvSpltK ranges at K = 4096 except FP16 NN m = 1
jinchen62 Oct 5, 2026
84bdfee
drop the FP16 NN wvSpltK kernels
jinchen62 Oct 5, 2026
5e7ce69
Merge branch 'develop' into gemv_kernel
jinchen62 Oct 5, 2026
466dd36
cap the ComputeUnits grid at the caller's CU budget
jinchen62 Oct 5, 2026
5b4f2c8
move the code-object guidance out of the Triton section
jinchen62 Oct 5, 2026
9e24e6d
Merge branch 'develop' into gemv_kernel
jinchen62 Oct 7, 2026
a738f4e
Merge remote-tracking branch 'upstream/develop' into gemv_kernel
jinchen62 Oct 7, 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
19 changes: 19 additions & 0 deletions projects/hipblaslt/clients/tests/data/matmul_gtest.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -104,6 +104,25 @@ Tests:
matrix_size:
- { M: 128, N: 128, K: 0 }

# Skinny GEMMs that library logic routes to the rocBLAS wvSpltK custom kernels on
# gfx950: padded leading dimensions they must honor, and general-batched
# (pointer-array) problems they cannot run, which must fall back to other kernels.
- name: matmul_skinny_bf16_tn
category: pre_checkin
function:
matmul: *hpa_bf16_precision
transA: T
transB: N
matrix_size:
- { M: 1024, N: 1, K: 1024, lda: 1024, ldb: 1024, ldc: 1024, ldd: 1024 }
- { M: 1024, N: 2, K: 1024, lda: 1032, ldb: 1040, ldc: 1048, ldd: 1056 }
- { M: 1024, N: 3, K: 1024, lda: 1024, ldb: 1032, ldc: 1024, ldd: 1024 }
- { M: 1024, N: 4, K: 1024, lda: 1024, ldb: 1024, ldc: 1024, ldd: 1024 }
alpha: 1
beta: [0, 2]
batch_mode: [0, 1]
batch_count: 1

# Tests confirm no NaN propagation when alpha = 0, 2 and beta = 0. Value .NaN is converted into zero
- {name: alpha_beta_zero_NaN, category: pre_checkin, precision: *real_precisions,
function: matmul, transA: N, transB: N, M: 256, N: 128, K: 64, alpha: [ .NaN, 2 ], beta: [ .NaN, 2 ] }
Expand Down
Loading
Loading