Repository navigation
fix(pascal): restore the sm_60 build (STRATA_DP4A in the poly int8 KV readers) and use CUBLAS_GEMM_DEFAULT - #87
Merged
Conversation
… poly cuBLAS calls The tensor-op hint asks cuBLAS for tensor-core algorithms, which the Tesla P100 (sm_60) does not have; the earlier version of this PR saw CUBLAS_STATUS_ARCH_MISMATCH there. The Strata engine already uses CUBLAS_GEMM_DEFAULT for these products, and on HIP both names map to the same algorithm.
…mpiles The --kv int8 readers added in 91fe653 call __dp4a directly, which exists only from sm_61, so -DSTRATA_EXPERIMENTAL_SM60=ON with sm_60 in CMAKE_CUDA_ARCHITECTURES fails with identifier __dp4a is undefined. strata/kernels/dp4a.hpp already has the bit-exact sm_60 fallback the i-quant kernels use; on sm_61 and newer STRATA_DP4A is __dp4a, so other cards compile to the same instruction.
Collaborator
|
Thanks, one thing to fix is the include of |
… builds see STRATA_DP4A The int8 KV readers use STRATA_DP4A outside the GDN MMA block, but dp4a.hpp was only included inside it, and STRATA_POLY_GDN_MMA is defined only for the CUDA strata_poly_mmq target. A HIP build of kernels.cu (and the CUDA poly_attend_test target, which also lacks the define) failed with: identifier "STRATA_DP4A" is undefined. Reported by VecSzn in simplexai-labs#87. Checked on CUDA 12.8 (sm_60/61/75) by compiling poly_attend_test's kernels.cu.o: 2 errors before, clean after. Not built with hipcc (none available here). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Contributor
Author
|
Thanks, good catch. Changes applied and tested with poly_attend_test: compiles clean after on sm_60/61/75. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Follows #43, which was closed (I can't reopen it from my side), reworked as you suggested.
Restores the Pascal sm_60 build of strata-poly (Tesla P100) with two small code changes and a docs paragraph. No block sizes or launch limits change on any card.
Changes
src/poly/kernels.cu: the three poly cuBLAS calls useCUBLAS_GEMM_DEFAULTinstead ofCUBLAS_GEMM_DEFAULT_TENSOR_OP(the P100 has no tensor cores; the Strata engine already usesCUBLAS_GEMM_DEFAULThere, and on HIP both names map to the same algorithm).src/poly/kernels.cu: the--kv int8readers added in 91fe653 call__dp4adirectly, which compute capability 6.0 lacks, so a build with60inCMAKE_CUDA_ARCHITECTURESfailed withidentifier "__dp4a" is undefined. They now useSTRATA_DP4Afromstrata/kernels/dp4a.hpp, like the i-quant kernels; sm_61 and newer compile to__dp4aexactly as before.docs/OLDER_GPUS.md: a short section with the cards, CUDA version, models, numbers and theSTRATA_ISA_FLOOR=avxnote. Existing text unchanged.Why the limit changes from #43 are gone
#43 lowered block sizes and launch bounds on every card to get Pascal running. The real cause turned out not to be Pascal-specific: @llollaalloll tracked it down in #73 (a V100 failing with "too many resources requested for launch":
kda_step_kernelcompiled to 72 registers a thread, and 1,024 threads need more than an SM's 65,536 registers), and 79d51a1 fixed it at the source with__launch_bounds__(1024)for every card. Thanks for that, it is what made this PR small. With it in main, Pascal runs today's values; both runs below used them. If another model hits a launch error on sm_60, the fix should be to bound that one kernel the same way, not to shrink sizes everywhere.Tested (CUDA 12.8, gcc 14,
-DSTRATA_EXPERIMENTAL_SM60=ON,CMAKE_CUDA_ARCHITECTURES="60;61;75", today's block sizes and launch bounds):-DSTRATA_ISA_FLOOR=avx): Tiel Coder 35B-A3B UD-Q5_K_XL, 6,085-token prompt then 40 tokens at 27.7 tok/s; 128 tokens at 39.4 tok/s on a short prompt.Not tested on sm_60: DeepSeek, Qwen3 Coder Next, Gemma 4 and other models.