[festeraeb]

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 use CUBLAS_GEMM_DEFAULT instead of CUBLAS_GEMM_DEFAULT_TENSOR_OP (the P100 has no tensor cores; the Strata engine already uses CUBLAS_GEMM_DEFAULT here, and on HIP both names map to the same algorithm).
• src/poly/kernels.cu: the --kv int8 readers added in 91fe653 call __dp4a directly, which compute capability 6.0 lacks, so a build with 60 in CMAKE_CUDA_ARCHITECTURES failed with identifier "__dp4a" is undefined. They now use STRATA_DP4A from strata/kernels/dp4a.hpp, like the i-quant kernels; sm_61 and newer compile to __dp4a exactly as before.
• docs/OLDER_GPUS.md: a short section with the cards, CUDA version, models, numbers and the STRATA_ISA_FLOOR=avx note. Existing text unchanged.

Tested (CUDA 12.8, gcc 14, -DSTRATA_EXPERIMENTAL_SM60=ON, CMAKE_CUDA_ARCHITECTURES="60;61;75", today's block sizes and launch bounds):

• Tesla P100, Xeon E5-2450 v2 host (AVX only, -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.
• Tesla P100, 2x Xeon Gold 6148 host: GLM-5.3 Flash REAP50 Q3_K_M (exercises the top-k selection), 4,768-token prompt then 40 tokens at 1.4 tok/s; 128 tokens at 3.7 tok/s (most experts on the CPU).
• No CUDA errors in either run; both answers read correctly.

Not tested on sm_60: DeepSeek, Qwen3 Coder Next, Gemma 4 and other models.

[Copilot]

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

[VecSzn]

Thanks. The code change is right and I'd like to take it: the Strata project's own engine already uses CUBLAS_GEMM_DEFAULT for these products, and on HIP both names map to the same algorithm. On an RTX 5090 it gives the same greedy tokens as main after a 3,000-token prompt for Qwen3.6, LFM2.5 and GLM 5.3 Flash, and the prompt reads are the same or slightly faster (Qwen3.6 8,014-8,348 against 7,210-8,461 tok/s, three rounds each).

The documentation needs changes before I can merge it:

• README.md: please drop the new section. PolyStrata isn't a fork of yours, the README has six translations that would need the same section, and hardware notes like this live in docs/OLDER_GPUS.md.
• docs/OLDER_GPUS.md: please keep the existing text, which records what was and wasn't run, and add a short paragraph with what you ran: the cards, the CUDA version, the models and files, your numbers, and the -DSTRATA_ISA_FLOOR=avx note (docs/INSTALL.md describes that option too, so a link to it is enough). "Full hardware support has been verified" says more than one machine's run can.
• The "Research & Roadmap" section describes work that isn't in the repository. Please leave it out; it can go in a Discussion, or come with the PR that implements it.
• The build block's --target strata-poly, 60;61;70 and CUDA 12.8 are fine if that's how you built it.
• The commit's author is "Documentation Bot doc-bot@cesarops.com". Could you author it under your own name?

If you'd rather keep it short, I can also merge the kernels.cu change on its own with you as its author and write the OLDER_GPUS.md lines from your numbers. Just say which.

[VecSzn]

Thanks for the two new commits. The trouble is that they lower limits on every card, not only on Pascal:

• IDX_SELECT 1024 -> 256 is the block size of GLM's and DeepSeek's top-k selection. The pools it picks stay the same, but a quarter of the threads scan every pool of a long prompt.
• native_add_rms_norm_q8_1 and gemv_kernel capped at 256 threads, __launch_bounds__(1024) -> 256, and (ROWS <= 2 ? 4 : 1) -> 1 for the MMVQ kernels change the launch or the register budget on every architecture.

Please keep today's values on sm_70 and up and use the smaller ones only where they are needed: an #if __CUDA_ARCH__ < 700 around the attributes, and for the host-side launch sizes a choice from the device's compute capability. Or, if you think the smaller values cost nothing, show decode and long-prompt numbers for GLM or DeepSeek and a 35B-A3B model on a recent card before and after.

src/kernels/cuda/native_mmvq.cu is shared with the Strata project's engine, so a change there conflicts at every merge from it. If Pascal needs it, it would be worth sending there too.

The earlier points still stand: the README section, keeping docs/OLDER_GPUS.md's existing text, the roadmap section, and the commit author.

[VecSzn]

Closing for now but these new limits still apply for every card, not just Pascal. If you come back with them only below sm_70, just reopen it.

[festeraeb]

Ok sorry was deep into another project and just saw this i have dyslexia and dysgraphia so i have an agent write my .md work not aire i have an sm70 to test on

[VecSzn]

@festeraeb no need to apologize, and it's fine to have an agent writing the .md files. You don't need to have an sm_70 card for this: just keep today's values for sm_70 and up (put an #if __CUDA_ARCH__ < 700 around the attributes, and use the launch sizes from the device's compute capability), test it on your Pascal card, and I'll check here that the newer cards give the same tokens and speed. The CUBLAS_GEMM_DEFAULT change is good on its own. If you like, you can open a PR with just that and the OLDER_GPUS.md paragraph, and it can go into the next release with you as the author.

[festeraeb]

I am rewriting with a probe so its dynamic for any card. Ill have it for you later. Yahoo Mail: Search, Organize, Conquer On Fri, Oct 23, 2026 at 10:27 PM, ***@***.***> wrote: VecSzn left a comment (simplexai-labs/PolyStrata#43) @festeraeb no need to apologize, and it's fine to have an agent writing the .md files. You don't need to have an sm_70 card for this: just keep today's values for sm_70 and up (put an #if __CUDA_ARCH__ < 700 around the attributes, and use the launch sizes from the device's compute capability), test it on your Pascal card, and I'll check here that the newer cards give the same tokens and speed. The CUBLAS_GEMM_DEFAULT change is good on its own. If you like, you can open a PR with just that and the OLDER_GPUS.md paragraph, and it can go into the next release with you as the author. — Reply to this email directly, view it on GitHub, or unsubscribe. You are receiving this because you were mentioned.Message ID: ***@***.***>

[festeraeb]

I couldn't reopen this one from my side (GitHub only lets the person who closed it), so the reworked version is #87: rebased on current main, today's limits on every card, the CUBLAS_GEMM_DEFAULT change, a STRATA_DP4A fix that makes the sm_60 build compile again, and a short OLDER_GPUS.md paragraph with what I ran on two P100s.