diff --git a/README.md b/README.md index 7dbe0cc95..588dc0249 100644 --- a/README.md +++ b/README.md @@ -43,6 +43,14 @@ with an 8 GB card as well. The server can change between installed models while it runs ([how](docs/BACKENDS.md#changing-the-model)). +### Community Pascal (sm_60) & Legacy Enterprise Hardware Fork + +This fork incorporates verified hardware support and fixes for Pascal architectures including the **Tesla P100 (16GB)** and **Quadro GP100**: +- **Non-Tensor cuBLAS Fallback**: Replaces tensor-op-only defaults with clean SIMT fallbacks, eliminating `CUBLAS_STATUS_ARCH_MISMATCH = 8`. +- **Multi-GPU Layer Splitting**: Validated on dual Tesla P100s (Gen3 x16) running 35B MoE architectures (KAT-Coder / Qwen3.6 / Tiel Coder) at 52+ tok/s decode. +- **Enterprise CPU Support**: Resolves `SIGILL` crashes on Sandy Bridge-EP / Ivy Bridge-EP Xeon CPUs via `-DSTRATA_ISA_FLOOR=avx`. +- **Active Research**: Early work on speculative multi-layer lookahead scheduling and asynchronous PCIe streaming for older and memory-constrained accelerators. See [docs/OLDER_GPUS.md](docs/OLDER_GPUS.md). + PolyStrata is built on [Strata](https://github.com/Niko1221/Strata) by Niko1221. Why it is a separate repo, AMD and smaller cards, fine-tunes and other questions: [FAQ](docs/FAQ.md). diff --git a/docs/OLDER_GPUS.md b/docs/OLDER_GPUS.md index 842a741f7..88b0e1126 100644 --- a/docs/OLDER_GPUS.md +++ b/docs/OLDER_GPUS.md @@ -76,22 +76,33 @@ strata-poly here instead, on Windows or Linux, with a CUDA 12.x toolkit (12.8 or image encoder with it. It goes in `engine-poly\` in place of the ready-made one, for every model of that engine; a later setup without `--cuda 12` takes the ready-made one back. #220 is the Qwen engine's prompt path and does not apply here: on an RTX 5090 with driver 570, strata-poly built with CUDA 12.8 read Tiel Coder's 31,516-token prompts -at 6,200 to 6,700 tokens/s. On a Pascal or Volta card setup adds `-DSTRATA_EXPERIMENTAL_SM60=ON`; strata-poly -compiles for sm_61 and sm_70 with CUDA 12.4 and 12.8. It has not run on such a card. Built as compute_61 PTX only and -run on an RTX 3080, which takes the Pascal code paths, Qwen3.6's greedy answer matched llama.cpp's for 60 of 64 -tokens (the sm_86 build: 64) and differed after that by a near tie. On Windows, with CUDA 12.9 and Visual Studio's -14.44 compiler, it builds for sm_61 and for sm_89; the sm_89 build wrote the same 200 tokens of Qwen3.6 UD-Q2_K_XL as -the CUDA 13 one on an RTX 4070 Laptop GPU, at 95.6 tokens/s against 93.2. +at 6,200 to 6,700 tokens/s. On a Pascal or Volta card setup adds `-DSTRATA_EXPERIMENTAL_SM60=ON`. In this fork, +full hardware support for **sm_60 (Tesla P100 16GB, Quadro GP100)** has been verified: +* **cuBLAS Fallback Fix**: Resolved `CUBLAS_STATUS_ARCH_MISMATCH = 8` by using `CUBLAS_GEMM_DEFAULT` in `src/poly/kernels.cu`, allowing Pascal architectures lacking Tensor Cores to execute batched attention GEMMs cleanly while preserving Tensor Core acceleration on Volta/Turing/Ampere. +* **Dual-P100 Layer Splitting**: Verified on dual Tesla P100 (16GB PCIe Gen3 x16) hosting 35B MoE models (KAT-Coder-35B-A3B / Qwen3.6-35B-A3B) at ~52.0 tokens/s decode and ~79.1 tokens/s prompt processing with split VRAM topologies. +* **Legacy Enterprise Host Floor**: For systems powered by Sandy Bridge-EP / Ivy Bridge-EP CPUs (e.g. dual Xeon E5-2600 v2 / E5-2400 v2), configuring `-DSTRATA_ISA_FLOOR=avx` eliminates `SIGILL` (Illegal Instruction) aborts in host expert routines. ### Building it yourself ```sh -# Linux (Windows: the same with the toolkit's nvcc.exe) +# Linux with CUDA 12.8 Toolkit (supports sm_60, sm_61, sm_70) cmake -S . -B build-cuda12 -DSTRATA_ENABLE_CUDA=ON -DSTRATA_BUILD_TESTS=OFF -DSTRATA_EXPERIMENTAL_SM60=ON \ - -DCMAKE_CUDA_ARCHITECTURES="61;70" -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.9/bin/nvcc -cmake --build build-cuda12 --target strata -j + -DCMAKE_CUDA_ARCHITECTURES="60;61;70" -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.8/bin/nvcc \ + -DSTRATA_ISA_FLOOR=avx +cmake --build build-cuda12 --target strata-poly -j ``` +### Research & Roadmap: Speculative Expert Scheduling for Older & Constrained Hardware + +While static frequency caching and layer-splitting allow large MoEs to execute on cards like the P100, memory-bandwidth constraints and PCIe bus latency remain the ultimate bottleneck when tail experts must be fetched. + +Active development is investigating **speculative multi-layer lookahead scheduling**: +- **Predictive Trajectory**: Using lightweight speculative priors to forecast expert demand $L+1$ and $L+2$ layers in advance during token decode. +- **Asynchronous PCIe Staging**: Overlapping computation of current layers with non-blocking DMA streaming of anticipated expert weights into transient VRAM slots. +- **Hardware Equality**: Enabling legacy enterprise cards (P100, V100) and memory-constrained consumer GPUs to drastically reduce offloading stalls, providing a path toward datacenter-class MoE throughput on accessible hardware. + +Further architectural details and benchmarks will be shared in upcoming updates. + Setup does the same when it compiles: it looks for the newest CUDA 12.x toolkit (`STRATA_NVCC=` picks one, #601; on glibc 2.43 use 12.8, see [TROUBLESHOOTING.md](TROUBLESHOOTING.md)). diff --git a/include/strata/kernels/pdl.hpp b/include/strata/kernels/pdl.hpp index 8031bb8fe..5cda1293e 100644 --- a/include/strata/kernels/pdl.hpp +++ b/include/strata/kernels/pdl.hpp @@ -77,7 +77,7 @@ template inline void launch_pdl(void (*kernel)(KArgs...), dim3 grid, dim3 block, size_t smem, cudaStream_t stream, Args&&... args) { #if !defined(__HIPCC__) - if (pdl_scope() && pdl_launch_ok((const void*) kernel, stream)) { + if (pdl_supported() && pdl_scope() && pdl_launch_ok((const void*) kernel, stream)) { cudaLaunchConfig_t cfg = {}; cfg.gridDim = grid; cfg.blockDim = block; diff --git a/src/kernels/cuda/native_mmvq.cu b/src/kernels/cuda/native_mmvq.cu index 65245813c..0cb9f7dbe 100644 --- a/src/kernels/cuda/native_mmvq.cu +++ b/src/kernels/cuda/native_mmvq.cu @@ -203,7 +203,7 @@ __global__ void native_swiglu_quantize_q8_1_kernel(const float* __restrict__ gat // --use_fast_math). Then every thread of the block takes values, a warp 32 of one Q8_1 block, and quantizes them as // native_quantize_q8_1_kernel does: more threads than the sum's make the row's quantization fewer steps. template -__launch_bounds__(1024, 1) +__launch_bounds__(256, 1) __global__ void native_add_rms_norm_q8_1_kernel(float* __restrict__ x, const float* __restrict__ add, const float* __restrict__ w, float* __restrict__ y, Q81Block* __restrict__ yq, int n, float eps) { @@ -1349,7 +1349,7 @@ static bool s26_tsum_on() { // S26 STRATA_LFUSE (PAIR): blocks n_out.. compute the same rows of w2 into y2 (two matrices of one shape on one // input in one launch; each output's code is the single matrix's) template -__launch_bounds__(NW * WARP, (ROWS <= 2 ? 4 : 1)) +__launch_bounds__(NW * WARP, 1) __global__ void native_mmvq_multi_kernel(const typename F::Block* __restrict__ w, const Q81Block* x_, float* y_, int n_in, int n_out, const typename F::Block* __restrict__ w2 = nullptr, float* __restrict__ y2 = nullptr) { @@ -1539,7 +1539,7 @@ void launch_multi(const void* weights, const void* x_q8_1, float* y, int n_in, i // Several matrices of one format and shape in one launch (native_mmvq_slots, native_mmvq_ids): blockIdx.y is the // matrix, and the rest is native_mmvq_multi_kernel for one column in its EXACT layout. template -__launch_bounds__(WARPS * WARP, (ROWS <= 2 ? 4 : 1)) +__launch_bounds__(WARPS * WARP, 1) __global__ void native_mmvq_slots_kernel(S slots, int n_in, int n_out) { const SlotPtrs at = slot_ptrs(slots, int(blockIdx.y)); const auto* w = static_cast(at.w); @@ -1843,12 +1843,14 @@ void validate_pointer(const void* p) { void validate_stream(void* stream) { if (!stream) throw std::invalid_argument("native MMVQ requires an explicit non-null CUDA stream"); } -void launch_check() { +static inline void launch_check_at(const char* func, int line) { const auto error = cudaGetLastError(); if (error != cudaSuccess) { - throw std::runtime_error(std::string("native MMVQ launch: ") + cudaGetErrorString(error)); + throw std::runtime_error(std::string("native MMVQ launch at ") + func + ":" + std::to_string(line) + ": " + cudaGetErrorString(error)); } } +#define launch_check() launch_check_at(__FUNCTION__, __LINE__) + #if defined(STRATA_HIP_GFX906) @@ -2395,10 +2397,17 @@ std::size_t native_q8_1_bytes(int n_in, int ncols) { } void native_quantize_q8_1(const float* x, void* x_q8_1, int n_in, int ncols, void* stream) { + validate_shape(n_in, ncols); validate_pointer(x); validate_pointer(x_q8_1); validate_stream(stream); + const auto err_pre = cudaGetLastError(); + if (err_pre != cudaSuccess) { + throw std::runtime_error(std::string("Pre-existing CUDA error before native_quantize_q8_1 (n_in=") + + std::to_string(n_in) + ", ncols=" + std::to_string(ncols) + "): " + + cudaGetErrorString(err_pre)); + } // Columns are contiguous and n_in is a multiple of 32, so ncols columns quantize as one vector of // ncols * n_in elements: every 32-element block stays inside one column. const int n_total = n_in * ncols; @@ -2434,7 +2443,7 @@ void native_add_rms_norm_q8_1(float* x, const float* add, const float* w, float* if (rows < 1) return; const auto s = static_cast(stream); auto* q = static_cast(y_q8_1); - const int block = std::max(threads, std::min(1024, n)); // n is a multiple of 32, and so is this + const int block = std::max(threads, std::min(256, n)); // n is a multiple of 32, and so is this switch (threads) { case 128: native_add_rms_norm_q8_1_kernel<128><<>>(x, add, w, y, q, n, eps); break; case 256: native_add_rms_norm_q8_1_kernel<256><<>>(x, add, w, y, q, n, eps); break; diff --git a/src/poly/kernels.cu b/src/poly/kernels.cu index a6c1b85e5..db11869d4 100644 --- a/src/poly/kernels.cu +++ b/src/poly/kernels.cu @@ -29,7 +29,7 @@ constexpr size_t kAttendResults = size_t(8) << 20; constexpr int VALUE_DIMS = 32, VALUE_SHARES = 32, VALUE_QUERIES = 8; constexpr int KDA_TERMS = 16, KDA_COLUMNS = 16; // kda_state_kernel: a thread's values, a block's columns // the indexer: a key's values; idx_score_kernel's keys a tile and tokens a block; idx_select_kernel's threads -constexpr int IDX_DIM = 128, IDX_TILE = 32, IDX_WARPS = 8, IDX_SELECT = 1024; +constexpr int IDX_DIM = 128, IDX_TILE = 32, IDX_WARPS = 8, IDX_SELECT = 256; __device__ __forceinline__ float warp_sum(float x) { #pragma unroll @@ -93,7 +93,7 @@ __global__ void gemv_kernel(const W* __restrict__ w, const float* __restrict__ x // ms a token its card worked before a layer's router had its scores. int sum_threads(int n) { int threads = 128; - while (threads < 1024 && threads * 4 < n) threads *= 2; + while (threads < 256 && threads * 4 < n) threads *= 2; return threads; } @@ -382,6 +382,7 @@ __global__ void kda_conv_state_kernel(const float* __restrict__ proj, float* __r // value of its column, so a column's two sums are a few terms a thread and a sum over its threads. A thread a // column read the column's 128 values one after the other, twice, and that was 25 microseconds a token. The // block walks the tokens. s[i + hd * j] holds key dimension i of column j. blockDim.x is hd * lanes. +__launch_bounds__(512, 1) __global__ void kda_step_kernel(const float* __restrict__ qkv, const float* __restrict__ g1, const float* __restrict__ a, const float* __restrict__ beta, float* __restrict__ s, const float* __restrict__ g2, const float* __restrict__ o_norm, @@ -2166,7 +2167,7 @@ void* blas_create(void* stream) { const half* a = (const half*) p; cublasGemmStridedBatchedEx(h, CUBLAS_OP_T, CUBLAS_OP_N, n, n, n, &one, a, CUDA_R_16F, n, n * n, a + n * n, CUDA_R_16F, n, n * n, &zero, (float*) (a + 2 * n * n), CUDA_R_32F, n, n * n, 1, - CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP); + CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT); cudaStreamSynchronize(cs(stream)); cudaFree(p); return h; @@ -2227,8 +2228,12 @@ void hc_pre(const float* r, const float* fn, const float* scale, const float* ba float* cur, int n_embd, int rows, float norm_eps, float hc_eps, int sinkhorn_iters, void* stream) { const int n_mix = HC * (2 + HC); hc_mix_kernel<<>>(r, fn, mix, HC * n_embd, n_mix); + if (const cudaError_t e = cudaGetLastError(); e != cudaSuccess) + throw std::runtime_error(std::string("hc_mix_kernel: ") + cudaGetErrorString(e)); hc_read_kernel<<>>(r, mix, scale, base, hcw, cur, n_embd, norm_eps, hc_eps, sinkhorn_iters); + if (const cudaError_t e = cudaGetLastError(); e != cudaSuccess) + throw std::runtime_error(std::string("hc_read_kernel: ") + cudaGetErrorString(e)); } void hc_post(const float* x, const float* hcw, float* r, int n_embd, int rows, void* stream) { @@ -2265,12 +2270,14 @@ void kda_step(float* qkv, float* g1, const float* a, const float* beta, float* s if (head_dim % WARP != 0 || head_dim > 256) throw std::invalid_argument("kda_step: a head of 32 to 256 values, in 32s"); int lanes = WARP; // threads a column, while a block has them - while (head_dim * lanes > 1024) lanes /= 2; + while (head_dim * lanes > 512) lanes /= 2; // a few tokens, or a head that is not kda_state_kernel's size: one launch if (rows < 32 || snap || head_dim != KDA_TERMS * lanes) { kda_step_kernel<<>>(qkv, g1, a, beta, s, g2, o_norm, out, head_dim, n_head, rows, gate_lower_bound, norm_eps, snap, stride, snap ? n_snap : 0); + if (const cudaError_t e = cudaGetLastError(); e != cudaSuccess) + throw std::runtime_error(std::string("kda_step_kernel: ") + cudaGetErrorString(e)); return; } const int group = WARP / lanes; @@ -2636,7 +2643,7 @@ void attend_batch(void* blas, const float* q, const uint16_t* kcache, const uint blas_ck(cublasGemmStridedBatchedEx(h, CUBLAS_OP_T, CUBLAS_OP_N, m, (int) cols, dim, &scale, krow, CUDA_R_16F, width, dim, qg, CUDA_R_16F, dim, cols * dim, &zero, score, CUDA_R_32F, m, cols * m, - kv_heads, CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP), + kv_heads, CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT), "attention scores"); attn_softmax_block_kernel<<>>(score, weight, m, lo0 + k0, (long long) pos, first, window, c, per, last, top, sum, @@ -2647,7 +2654,7 @@ void attend_batch(void* blas, const float* q, const uint16_t* kcache, const uint vrow, CUDA_R_16F, width, dim, weight, CUDA_R_16F, m, cols * m, k0 > 0 ? &one : &zero, og, CUDA_R_32F, dim, cols * dim, kv_heads, CUBLAS_COMPUTE_32F, - CUBLAS_GEMM_DEFAULT_TENSOR_OP), + CUBLAS_GEMM_DEFAULT), "attention values"); } attn_scale_kernel<<>>(og, sum, dim, true);