diff --git a/docs/beellama-features.md b/docs/beellama-features.md index 2f2fe7d076b0..c217a9eb6e57 100644 --- a/docs/beellama-features.md +++ b/docs/beellama-features.md @@ -198,6 +198,12 @@ CDNA fast routing is compiled and selected by capability but remains experimental until hardware parity and performance results are published. MUSA explicitly remains on the portable route. +On HIP, Bee reports `integrated = false`, backing out the upstream APU +zero-copy host-buffer path after async-execution corruption was observed +(PPL 5.9243 -> 8.51+ without `HIP_LAUNCH_BLOCKING`). This changes APU +tensor placement off host-mapped memory and therefore VRAM headroom; the +trade-off is unmeasured (no APU hardware available). + Set `GGML_KVARN_DEBUG_ROUTES=1` to log the selected CUDA/HIP route, compute capability, rotated/original domain, K/V bit widths, query and KV counts, attached exact-tail rows and type for integrated entries, entry path, and diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 880db1a77ae7..e6497c768744 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -307,7 +307,14 @@ static ggml_cuda_device_info ggml_cuda_init() { info.default_tensor_split[id] = total_vram; total_vram += device_vram; #if defined(GGML_USE_HIP) - info.devices[id].integrated = prop.integrated; + // Fork divergence from PR #24233: integrated=true enables the CUDA + // host-buffer path (zero-copy UMA weights) on APUs, which corrupts + // full-model results under async execution on this box (PPL 5.9243 + // -> 8.51+ without HIP_LAUNCH_BLOCKING). The back-out restores + // async-safe operation. Placement changes with it: supports_buft gates + // the CUDA-host buffer type on `integrated`, so APU tensors move off + // host-mapped memory and VRAM headroom changes (unmeasured, no APU HW here). + info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034) #else info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034) #endif diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 976251a1ba60..3e8400fcbd06 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -27,9 +27,24 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) case GGML_TYPE_NVFP4: return vec_dot_nvfp4_q8_1; case GGML_TYPE_Q2_K: return vec_dot_q2_K_q8_1; case GGML_TYPE_Q3_K: return vec_dot_q3_K_q8_1; - case GGML_TYPE_Q4_K: return vec_dot_q4_K_q8_1; - case GGML_TYPE_Q5_K: return vec_dot_q5_K_q8_1; - case GGML_TYPE_Q6_K: return vec_dot_q6_K_q8_1; + case GGML_TYPE_Q4_K: +#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4) + return vec_dot_q4_K_q8_1_vdr4; +#else + return vec_dot_q4_K_q8_1; +#endif + case GGML_TYPE_Q5_K: +#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4) + return vec_dot_q5_K_q8_1_vdr4; +#else + return vec_dot_q5_K_q8_1; +#endif + case GGML_TYPE_Q6_K: +#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4) + return vec_dot_q6_K_q8_1_vdr2; +#else + return vec_dot_q6_K_q8_1; +#endif case GGML_TYPE_IQ2_XXS: return vec_dot_iq2_xxs_q8_1; case GGML_TYPE_IQ2_XS: return vec_dot_iq2_xs_q8_1; case GGML_TYPE_IQ2_S: return vec_dot_iq2_s_q8_1; @@ -82,6 +97,7 @@ enum mmvq_parameter_table_id { MMVQ_PARAMETERS_GCN, MMVQ_PARAMETERS_RDNA2, MMVQ_PARAMETERS_RDNA3_0, + MMVQ_PARAMETERS_RDNA3_5, MMVQ_PARAMETERS_RDNA4, MMVQ_PARAMETERS_GB10 }; @@ -91,7 +107,9 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() { return MMVQ_PARAMETERS_RDNA4; #elif defined(RDNA3_0) return MMVQ_PARAMETERS_RDNA3_0; -#elif defined(RDNA2) || defined(RDNA3_5) +#elif defined(RDNA3_5) + return MMVQ_PARAMETERS_RDNA3_5; +#elif defined(RDNA2) return MMVQ_PARAMETERS_RDNA2; #elif defined(GCN) || defined(CDNA) return MMVQ_PARAMETERS_GCN; @@ -111,7 +129,10 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) { if (GGML_CUDA_CC_IS_RDNA3_0(cc)) { return MMVQ_PARAMETERS_RDNA3_0; } - if (GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc)) { + if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { + return MMVQ_PARAMETERS_RDNA3_5; + } + if (GGML_CUDA_CC_IS_RDNA2(cc)) { return MMVQ_PARAMETERS_RDNA2; } if (GGML_CUDA_CC_IS_GCN(cc) || GGML_CUDA_CC_IS_CDNA(cc)) { @@ -260,7 +281,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type case GGML_TYPE_Q3_K: return 4; case GGML_TYPE_Q4_0: return 7; case GGML_TYPE_Q4_1: return 7; - case GGML_TYPE_Q4_K: return 4; + case GGML_TYPE_Q4_K: return 7; case GGML_TYPE_Q5_0: return 7; case GGML_TYPE_Q5_1: return 7; case GGML_TYPE_Q6_0: return 6; @@ -508,6 +529,24 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d } return 1; } + if (table_id == MMVQ_PARAMETERS_RDNA3_5) { + // gfx1151 (Strix Halo iGPU): nwarps=1 (the RDNA2 table) underutilizes the + // wave32 datapath on the large-K decode matmuls; nwarps=8 (the RDNA3_0 + // table) over-parallelizes the small ones. Swept 2025-08: nwarps=2 wins + // (~+0.6% decode on Qwen3.6-35B-A3B Q8_0), nwarps=4 regresses. + // Apply to the whole mmvq range (ncols_dst 1..8), not just decode: the + // speculative verify batch (n_draft+1 tokens) must use the same nwarps + // as decode so its per-row dot-product accumulation is bit-identical. + if (ncols_dst <= MMVQ_MAX_BATCH_SIZE) { + switch (type) { + case GGML_TYPE_Q8_0: + return 2; + default: + return 1; + } + } + return 1; + } if (table_id == MMVQ_PARAMETERS_TURING) { if (ncols_dst == 1) { switch (type) { diff --git a/ggml/src/ggml-cuda/top-k.cu b/ggml/src/ggml-cuda/top-k.cu index f9d262aee30b..32516f08639f 100644 --- a/ggml/src/ggml-cuda/top-k.cu +++ b/ggml/src/ggml-cuda/top-k.cu @@ -1,5 +1,8 @@ #include "argsort.cuh" #include "top-k.cuh" +#if defined(GGML_USE_HIP) +#include +#endif #ifdef GGML_CUDA_USE_CUB # include @@ -147,11 +150,732 @@ static bool ggml_cuda_top_k_tiled(ggml_cuda_pool & pool, const float * src, int static __device__ __forceinline__ uint32_t top_k_float_to_ordered(float value) { const uint32_t bits = __float_as_uint(value); - const uint32_t mask = (uint32_t) (-(int32_t) (bits >> 31)) | 0x80000000U; - return bits ^ mask; + return (bits & 0x80000000U) != 0 ? ~bits : bits | 0x80000000U; +} + +template +static __global__ void top_k_nary_search_cuda( + const float * __restrict__ src, + const int2 * __restrict__ src_pairs, + int * __restrict__ dst, + int2 * __restrict__ dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int k, + int nrows, + bool first_pass, + bool last_pass) { + const int tid = threadIdx.x; + const int lane = tid % warpSize; + const int warp = tid / warpSize; + const int warp_count = BLOCK_SIZE / warpSize; + + __shared__ int2 candidates[BLOCK_SIZE]; + __shared__ uint32_t counts[64]; + __shared__ uint32_t selected_bucket; + __shared__ uint32_t selected_total; + __shared__ uint32_t warp_offsets[32]; + __shared__ uint32_t warp_equal_offsets[32]; + + for (int row = blockIdx.y; row < nrows; row += gridDim.y) { + const int col = blockIdx.x * BLOCK_SIZE + tid; + const bool valid = col < ncols_input; + int2 value; + if (valid) { + value = first_pass + ? make_int2(col, __float_as_int(src[(size_t) row * ncols_input + col])) + : src_pairs[(size_t) row * ncols_input + col]; + } else { + value = make_int2(original_ncols, (int) 0xff800000U); + } + candidates[tid] = value; + __syncthreads(); + + const int limit = min(k, ncols_input - blockIdx.x * BLOCK_SIZE); + if (k == 1) { +#pragma unroll + for (int stride = BLOCK_SIZE / 2; stride >= 1; stride /= 2) { + if (tid < stride) { + const int2 a = candidates[tid]; + const int2 b = candidates[tid + stride]; + if (a.x >= original_ncols || + (b.x < original_ncols && __int_as_float(b.y) > __int_as_float(a.y))) { + candidates[tid] = b; + } + } + __syncthreads(); + } + } else { + constexpr int radix_bits = 6; + constexpr int radix_size = 1 << radix_bits; + int shift = 32 - radix_bits; + uint32_t mask = ((1U << radix_bits) - 1) << shift; + uint32_t range_min = 0; + uint32_t range_max = 0xff800000U; + uint32_t total = 0; + + while (mask != 0) { + __syncthreads(); + if (tid < radix_size) { + counts[tid] = 0; + } + __syncthreads(); + + // Seed per iteration: both ballots below stay zero when the current range holds + // fewer than `limit` keys (e.g. all-+inf/NaN row, ordered keys >= 0xff800000 + // excluded above), so no lane writes the shared selection. Zeroes keep the + // broadcast read deterministic instead of uninitialized (first pass) / stale. + if (tid == 0) { + selected_bucket = 0; + selected_total = 0; + } + + const uint32_t key = top_k_float_to_ordered(__int_as_float(value.y)); + if (valid && key >= range_min && key < range_max) { + atomicAdd(&counts[(key & mask) >> shift], 1U); + } + __syncthreads(); + + if (tid < 32) { + uint32_t high_partial = counts[radix_size - 1 - tid]; + uint32_t low_partial = counts[31 - tid]; + for (int offset = 1; offset < 32; offset *= 2) { + const uint32_t high_previous = __shfl_up(high_partial, offset, 32); + const uint32_t low_previous = __shfl_up(low_partial, offset, 32); + if (tid >= offset) { + high_partial += high_previous; + low_partial += low_previous; + } + } + high_partial += total; + low_partial += __shfl(high_partial, 31, 32); + const unsigned long long high_selected = __ballot(high_partial >= (uint32_t) limit); + if (high_selected != 0) { + const int first = __ffsll(high_selected) - 1; + if (tid == first) { + selected_bucket = radix_size - 1 - first; + selected_total = high_partial; + } + } else { + const unsigned long long low_selected = __ballot(low_partial >= (uint32_t) limit); + const int first = __ffsll(low_selected) - 1; + if (tid == first) { + selected_bucket = 31 - first; + selected_total = low_partial; + } + } + } + __syncthreads(); + + const uint32_t bucket = selected_bucket; + total = selected_total; + range_max = range_min + ((bucket + 1) << shift); + range_min = range_min + (bucket << shift); + if (total == (uint32_t) limit) { + break; + } + total -= counts[bucket]; + mask >>= radix_bits; + shift -= radix_bits; + if (shift < 0) { + shift = 0; + } + } + + const uint32_t key = top_k_float_to_ordered(__int_as_float(value.y)); + const bool above = valid && key > range_min; + const bool equal = valid && key == range_min; + const unsigned long long above_mask = __ballot(above); + const unsigned long long equal_mask = __ballot(equal); + if (lane == 0) { + warp_offsets[warp] = __popcll(above_mask); + warp_equal_offsets[warp] = __popcll(equal_mask); + } + __syncthreads(); + + uint32_t above_base = 0; + uint32_t equal_base = 0; + uint32_t above_total = 0; + for (int i = 0; i < warp_count; ++i) { + if (i < warp) { + above_base += warp_offsets[i]; + equal_base += warp_equal_offsets[i]; + } + above_total += warp_offsets[i]; + } + equal_base += above_total; + + const unsigned long long lane_mask = lane == 0 ? 0 : (1ULL << lane) - 1; + if (above) { + candidates[above_base + __popcll(above_mask & lane_mask)] = value; + } + const uint32_t equal_index = equal_base + __popcll(equal_mask & lane_mask); + if (equal && equal_index < (uint32_t) limit) { + candidates[equal_index] = value; + } + __syncthreads(); + } + + if (tid < k) { + if (last_pass) { + dst[(size_t) row * k + tid] = candidates[tid].x; + } else { + const int output_col = blockIdx.x * k + tid; + if (output_col < ncols_output) { + dst_pairs[(size_t) row * ncols_output + output_col] = candidates[tid]; + } + } + } + __syncthreads(); + } } -struct top_k_radix_state { +template +static __device__ __forceinline__ int2 top_k_one_reduce(int2 * candidates, int2 value, int original_ncols) { + const int tid = threadIdx.x; + if constexpr (USE_SHUFFLE) { + const int lane = tid & 31; + const int wave = tid >> 5; + constexpr int wave_count = BLOCK_SIZE / 32; + +# pragma unroll + for (int offset = 16; offset >= 1; offset /= 2) { + const int2 other = make_int2(__shfl_down(value.x, offset, 32), __shfl_down(value.y, offset, 32)); + if (lane + offset < 32 && + (value.x >= original_ncols || + (other.x < original_ncols && __int_as_float(other.y) > __int_as_float(value.y)))) { + value = other; + } + } + if (lane == 0) { + candidates[wave] = value; + } + __syncthreads(); + + if (wave == 0) { + value = lane < wave_count ? candidates[lane] : make_int2(original_ncols, (int) 0xff800000U); +# pragma unroll + for (int offset = 16; offset >= 1; offset /= 2) { + const int2 other = make_int2(__shfl_down(value.x, offset, 32), __shfl_down(value.y, offset, 32)); + if (lane + offset < 32 && + (value.x >= original_ncols || + (other.x < original_ncols && __int_as_float(other.y) > __int_as_float(value.y)))) { + value = other; + } + } + } + return value; + } else { + candidates[tid] = value; + __syncthreads(); + +# pragma unroll + for (int stride = BLOCK_SIZE / 2; stride >= 1; stride /= 2) { + if (tid < stride) { + const int2 a = candidates[tid]; + const int2 b = candidates[tid + stride]; + if (a.x >= original_ncols || (b.x < original_ncols && __int_as_float(b.y) > __int_as_float(a.y))) { + candidates[tid] = b; + } + } + __syncthreads(); + } + return candidates[0]; + } +} + +template +static __global__ void top_k_one_first_cuda(const float * __restrict__ src, + int2 * __restrict__ dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int nrows) { + const int tid = threadIdx.x; + __shared__ int2 candidates[USE_SHUFFLE ? BLOCK_SIZE / 32 : BLOCK_SIZE]; + + for (int row = blockIdx.y; row < nrows; row += gridDim.y) { + const int col = blockIdx.x * (ITEMS_PER_THREAD * BLOCK_SIZE) + tid; + int2 value = col < ncols_input ? make_int2(col, __float_as_int(src[(size_t) row * ncols_input + col])) : + make_int2(original_ncols, (int) 0xff800000U); +# pragma unroll + for (int item = 1; item < ITEMS_PER_THREAD; ++item) { + const int other_col = col + item * BLOCK_SIZE; + if (other_col < ncols_input) { + const int2 other = make_int2(other_col, __float_as_int(src[(size_t) row * ncols_input + other_col])); + if (value.x >= original_ncols || __int_as_float(other.y) > __int_as_float(value.y)) { + value = other; + } + } + } + const int2 result = top_k_one_reduce(candidates, value, original_ncols); + if (tid == 0) { + dst_pairs[(size_t) row * ncols_output + blockIdx.x] = result; + } + __syncthreads(); + } +} + +template +static __global__ void top_k_one_first_last_cuda(const float * __restrict__ src, + int * __restrict__ dst, + int original_ncols, + int ncols_input, + int nrows) { + const int tid = threadIdx.x; + __shared__ int2 candidates[USE_SHUFFLE ? BLOCK_SIZE / 32 : BLOCK_SIZE]; + + for (int row = blockIdx.y; row < nrows; row += gridDim.y) { + const int col = blockIdx.x * (ITEMS_PER_THREAD * BLOCK_SIZE) + tid; + int2 value = col < ncols_input ? make_int2(col, __float_as_int(src[(size_t) row * ncols_input + col])) : + make_int2(original_ncols, (int) 0xff800000U); +# pragma unroll + for (int item = 1; item < ITEMS_PER_THREAD; ++item) { + const int other_col = col + item * BLOCK_SIZE; + if (other_col < ncols_input) { + const int2 other = make_int2(other_col, __float_as_int(src[(size_t) row * ncols_input + other_col])); + if (value.x >= original_ncols || __int_as_float(other.y) > __int_as_float(value.y)) { + value = other; + } + } + } + const int2 result = top_k_one_reduce(candidates, value, original_ncols); + if (tid == 0) { + dst[row] = result.x; + } + __syncthreads(); + } +} + +template +static __global__ void top_k_one_middle_cuda(const int2 * __restrict__ src_pairs, + int2 * __restrict__ dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int nrows) { + const int tid = threadIdx.x; + __shared__ int2 candidates[USE_SHUFFLE ? BLOCK_SIZE / 32 : BLOCK_SIZE]; + + for (int row = blockIdx.y; row < nrows; row += gridDim.y) { + const int col = blockIdx.x * (ITEMS_PER_THREAD * BLOCK_SIZE) + tid; + int2 value = col < ncols_input ? src_pairs[(size_t) row * ncols_input + col] : + make_int2(original_ncols, (int) 0xff800000U); +# pragma unroll + for (int item = 1; item < ITEMS_PER_THREAD; ++item) { + const int other_col = col + item * BLOCK_SIZE; + if (other_col < ncols_input) { + const int2 other = src_pairs[(size_t) row * ncols_input + other_col]; + if (value.x >= original_ncols || __int_as_float(other.y) > __int_as_float(value.y)) { + value = other; + } + } + } + const int2 result = top_k_one_reduce(candidates, value, original_ncols); + if (tid == 0) { + dst_pairs[(size_t) row * ncols_output + blockIdx.x] = result; + } + __syncthreads(); + } +} + +template +static __global__ void top_k_one_last_cuda(const int2 * __restrict__ src_pairs, + int * __restrict__ dst, + int original_ncols, + int ncols_input, + int nrows) { + const int tid = threadIdx.x; + __shared__ int2 candidates[USE_SHUFFLE ? BLOCK_SIZE / 32 : BLOCK_SIZE]; + + for (int row = blockIdx.y; row < nrows; row += gridDim.y) { + const int col = blockIdx.x * (ITEMS_PER_THREAD * BLOCK_SIZE) + tid; + int2 value = col < ncols_input ? src_pairs[(size_t) row * ncols_input + col] : + make_int2(original_ncols, (int) 0xff800000U); +# pragma unroll + for (int item = 1; item < ITEMS_PER_THREAD; ++item) { + const int other_col = col + item * BLOCK_SIZE; + if (other_col < ncols_input) { + const int2 other = src_pairs[(size_t) row * ncols_input + other_col]; + if (value.x >= original_ncols || __int_as_float(other.y) > __int_as_float(value.y)) { + value = other; + } + } + } + const int2 result = top_k_one_reduce(candidates, value, original_ncols); + if (tid == 0) { + dst[row] = result.x; + } + __syncthreads(); + } +} + +template +static __global__ void top_k_radix_select_cuda( + const float * __restrict__ src, + int * __restrict__ dst, + int ncols, + int nrows, + int k) { + constexpr int RADIX_BITS = 8; + constexpr int RADIX_SIZE = 1 << RADIX_BITS; + + const int tid = threadIdx.x; + __shared__ uint32_t histogram[RADIX_SIZE]; + __shared__ uint32_t selected_bucket; + __shared__ uint32_t count_above; + __shared__ uint32_t output_count; + + for (int row = blockIdx.x; row < nrows; row += gridDim.x) { + const float * row_src = src + (size_t) row * ncols; + int * row_dst = dst + (size_t) row * k; + uint32_t prefix = 0; + uint32_t desired = k; + +#pragma unroll + for (int shift = 32 - RADIX_BITS; shift >= 0; shift -= RADIX_BITS) { + for (int bin = tid; bin < RADIX_SIZE; bin += BLOCK_SIZE) { + histogram[bin] = 0; + } + __syncthreads(); + + const uint32_t high_mask = shift == 32 - RADIX_BITS ? 0 : 0xffffffffU << (shift + RADIX_BITS); + const uint32_t prefix_high = prefix & high_mask; + for (int col = tid; col < ncols; col += BLOCK_SIZE) { + const uint32_t key = top_k_float_to_ordered(row_src[col]); + if ((key & high_mask) == prefix_high) { + atomicAdd(&histogram[(key >> shift) & (RADIX_SIZE - 1)], 1U); + } + } + __syncthreads(); + + if (tid == 0) { + uint32_t above = 0; + uint32_t bucket = 0; + for (int bin = RADIX_SIZE - 1; bin >= 0; --bin) { + const uint32_t count = histogram[bin]; + if (above + count >= desired) { + bucket = bin; + break; + } + above += count; + } + selected_bucket = bucket; + count_above = above; + } + __syncthreads(); + + prefix |= selected_bucket << shift; + desired -= count_above; + __syncthreads(); + } + + if (tid == 0) { + output_count = 0; + } + __syncthreads(); + + for (int col = tid; col < ncols; col += BLOCK_SIZE) { + if (top_k_float_to_ordered(row_src[col]) > prefix) { + row_dst[atomicAdd(&output_count, 1U)] = col; + } + } + __syncthreads(); + + for (int col = tid; col < ncols; col += BLOCK_SIZE) { + if (top_k_float_to_ordered(row_src[col]) == prefix) { + const uint32_t output = atomicAdd(&output_count, 1U); + if (output < (uint32_t) k) { + row_dst[output] = col; + } + } + } + __syncthreads(); + } +} + +static void top_k_radix_select_cuda( + const float * src, int * dst, int ncols, int nrows, int k, cudaStream_t stream) { + constexpr int BLOCK_SIZE = 1024; + const int grid_size = std::min(nrows, 65535); + top_k_radix_select_cuda<<>>(src, dst, ncols, nrows, k); +} + +static int top_k_floor_log2(int value) { + int result = 0; + while (value > 1) { + value >>= 1; + ++result; + } + return result; +} + +static int top_k_ceil_log2(int value) { + const int floor = top_k_floor_log2(value); + return value == (1 << floor) ? floor : floor + 1; +} + +static int top_k_nary_block_log2(int ncols, int k) { + const int min_block = std::max(top_k_floor_log2(k) + 1, 6); + if (min_block > 10) { + return -1; + } + const int max_block = std::min(std::max(top_k_floor_log2(k) + 2, 8), 10); + int block = std::min(std::max(top_k_ceil_log2(ncols), min_block), max_block); + if (ncols > (1 << block)) { + for (int candidate = block; candidate <= 10; ++candidate) { + if (ncols <= (1 << candidate)) { + block = candidate; + break; + } + } + } + return block; +} + +template +static void top_k_one_cuda_launch(const float * src, + const int2 * src_pairs, + int * dst, + int2 * dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int nrows, + bool first_pass, + bool last_pass, + cudaStream_t stream) { + constexpr int TILE_SIZE = BLOCK_SIZE * ITEMS_PER_THREAD; + const dim3 grid((ncols_input + TILE_SIZE - 1) / TILE_SIZE, std::min(nrows, 65535), 1); + if (first_pass && last_pass) { + top_k_one_first_last_cuda + <<>>(src, dst, original_ncols, ncols_input, nrows); + } else if (first_pass) { + top_k_one_first_cuda + <<>>(src, dst_pairs, original_ncols, ncols_input, ncols_output, nrows); + } else if (last_pass) { + top_k_one_last_cuda + <<>>(src_pairs, dst, original_ncols, ncols_input, nrows); + } else { + top_k_one_middle_cuda + <<>>(src_pairs, dst_pairs, original_ncols, ncols_input, ncols_output, nrows); + } +} + +template +static void top_k_one_cuda_launch(int block_log2, + const float * src, + const int2 * src_pairs, + int * dst, + int2 * dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int nrows, + bool first_pass, + bool last_pass, + cudaStream_t stream) { + switch (block_log2) { + case 6: + top_k_one_cuda_launch<64, 1, USE_SHUFFLE>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + case 7: + top_k_one_cuda_launch<128, 1, USE_SHUFFLE>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + case 8: + top_k_one_cuda_launch<256, 1, USE_SHUFFLE>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + case 9: + top_k_one_cuda_launch<512, 1, USE_SHUFFLE>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + case 10: + top_k_one_cuda_launch<1024, 1, USE_SHUFFLE>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + default: + GGML_ABORT("invalid HIP TOP_K block size"); + } +} + +enum top_k_one_mode { + TOP_K_ONE_SHARED, + TOP_K_ONE_SHUFFLE, + TOP_K_ONE_SHUFFLE_2, + TOP_K_ONE_SHUFFLE_4, +}; + +static top_k_one_mode top_k_one_select_mode(int ncols, int nrows) { + if (ncols <= 32 || (nrows == 1 && ncols > 131072)) { + return TOP_K_ONE_SHARED; + } + if (ncols <= 4096) { + return TOP_K_ONE_SHUFFLE; + } + return nrows > 1 ? TOP_K_ONE_SHUFFLE_4 : TOP_K_ONE_SHUFFLE_2; +} + +static int top_k_one_items_per_thread(top_k_one_mode mode) { + switch (mode) { + case TOP_K_ONE_SHUFFLE_2: + return 2; + case TOP_K_ONE_SHUFFLE_4: + return 4; + default: + return 1; + } +} + +static int top_k_one_block_log2(top_k_one_mode mode, int ncols) { + switch (mode) { + case TOP_K_ONE_SHUFFLE_2: + return 7; + case TOP_K_ONE_SHUFFLE_4: + return 6; + default: + return top_k_nary_block_log2(ncols, 1); + } +} + +static void top_k_one_cuda_launch(top_k_one_mode mode, + int block_log2, + const float * src, + const int2 * src_pairs, + int * dst, + int2 * dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int nrows, + bool first_pass, + bool last_pass, + cudaStream_t stream) { + switch (mode) { + case TOP_K_ONE_SHARED: + top_k_one_cuda_launch(block_log2, src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + case TOP_K_ONE_SHUFFLE: + top_k_one_cuda_launch(block_log2, src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + case TOP_K_ONE_SHUFFLE_2: + top_k_one_cuda_launch<128, 2, true>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + case TOP_K_ONE_SHUFFLE_4: + top_k_one_cuda_launch<64, 4, true>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, + ncols_output, nrows, first_pass, last_pass, stream); + break; + } +} + +template +static void top_k_nary_search_cuda_launch( + const float * src, + const int2 * src_pairs, + int * dst, + int2 * dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int k, + int nrows, + bool first_pass, + bool last_pass, + cudaStream_t stream) { + const dim3 grid((ncols_input + BLOCK_SIZE - 1) / BLOCK_SIZE, std::min(nrows, 65535), 1); + top_k_nary_search_cuda<<>>( + src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, ncols_output, + k, nrows, first_pass, last_pass); +} + +static void top_k_nary_search_cuda_launch( + int block_log2, + const float * src, + const int2 * src_pairs, + int * dst, + int2 * dst_pairs, + int original_ncols, + int ncols_input, + int ncols_output, + int k, + int nrows, + bool first_pass, + bool last_pass, + cudaStream_t stream) { + switch (block_log2) { + case 6: + top_k_nary_search_cuda_launch<64>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, ncols_output, k, nrows, first_pass, last_pass, stream); + break; + case 7: + top_k_nary_search_cuda_launch<128>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, ncols_output, k, nrows, first_pass, last_pass, stream); + break; + case 8: + top_k_nary_search_cuda_launch<256>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, ncols_output, k, nrows, first_pass, last_pass, stream); + break; + case 9: + top_k_nary_search_cuda_launch<512>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, ncols_output, k, nrows, first_pass, last_pass, stream); + break; + case 10: + top_k_nary_search_cuda_launch<1024>(src, src_pairs, dst, dst_pairs, original_ncols, ncols_input, ncols_output, k, nrows, first_pass, last_pass, stream); + break; + default: + GGML_ABORT("invalid HIP TOP_K block size"); + } +} + +static void top_k_small_cuda( + ggml_cuda_pool & pool, + const float * src, + int * dst, + int ncols, + int nrows, + int k, + cudaStream_t stream) { + const top_k_one_mode one_mode = top_k_one_select_mode(ncols, nrows); + const int one_items_per_thread = top_k_one_items_per_thread(one_mode); + int block_log2 = k == 1 ? top_k_one_block_log2(one_mode, ncols) : top_k_nary_block_log2(ncols, k); + if (block_log2 < 0) { + top_k_radix_select_cuda(src, dst, ncols, nrows, k, stream); + return; + } + + const int block_size = (k == 1 ? one_items_per_thread : 1) << block_log2; + const int first_output = (ncols / block_size) * k + std::min(k, ncols % block_size); + const size_t scratch_elements = (size_t) first_output * nrows; + ggml_cuda_pool_alloc scratch_alloc(pool, 2 * scratch_elements); + int2 * scratch[2] = {scratch_alloc.get(), scratch_alloc.get() + scratch_elements}; + + int ncols_input = ncols; + int buffer = 0; + bool first_pass = true; + while (ncols_input > k || first_pass) { + block_log2 = k == 1 ? top_k_one_block_log2(one_mode, ncols_input) : top_k_nary_block_log2(ncols_input, k); + const int current_block_size = (k == 1 ? one_items_per_thread : 1) << block_log2; + const int ncols_output = (ncols_input / current_block_size) * k + std::min(k, ncols_input % current_block_size); + const bool last_pass = ncols_output == k; + if (k == 1) { + top_k_one_cuda_launch(one_mode, block_log2, src, first_pass ? nullptr : scratch[buffer], dst, + last_pass ? nullptr : scratch[buffer ^ 1], ncols, ncols_input, ncols_output, nrows, + first_pass, last_pass, stream); + } else { + top_k_nary_search_cuda_launch( + block_log2, src, first_pass ? nullptr : scratch[buffer], dst, + last_pass ? nullptr : scratch[buffer ^ 1], ncols, ncols_input, + ncols_output, k, nrows, first_pass, last_pass, stream); + } + ncols_input = ncols_output; + first_pass = false; + buffer ^= 1; + } +} + +struct top_k_parallel_radix_state { uint32_t prefix; uint32_t prefix_mask; int rank; @@ -159,7 +883,7 @@ struct top_k_radix_state { int equal_count; }; -static __global__ void top_k_radix_init(top_k_radix_state * states, int nrows, int k) { +static __global__ void top_k_parallel_radix_init(top_k_parallel_radix_state * states, int nrows, int k) { const int row = blockIdx.x * blockDim.x + threadIdx.x; if (row < nrows) { states[row] = {0, 0, k, 0, 0}; @@ -167,9 +891,9 @@ static __global__ void top_k_radix_init(top_k_radix_state * states, int nrows, i } template -static __global__ void top_k_radix_histogram( +static __global__ void top_k_parallel_radix_histogram( const float * __restrict__ src, - const top_k_radix_state * __restrict__ states, + const top_k_parallel_radix_state * __restrict__ states, int * __restrict__ block_histograms, int ncols, int blocks_per_row, @@ -185,7 +909,7 @@ static __global__ void top_k_radix_histogram( histogram[tid] = 0; __syncthreads(); - const top_k_radix_state state = states[row]; + const top_k_parallel_radix_state state = states[row]; for (int col = row_block * BLOCK_SIZE + tid; col < ncols; col += blocks_per_row * BLOCK_SIZE) { @@ -202,9 +926,9 @@ static __global__ void top_k_radix_histogram( } template -static __global__ void top_k_radix_select( +static __global__ void top_k_parallel_radix_select( const int * __restrict__ block_histograms, - top_k_radix_state * __restrict__ states, + top_k_parallel_radix_state * __restrict__ states, int blocks_per_row, int shift) { constexpr int NBINS = 1 << RADIX_BITS; @@ -222,7 +946,7 @@ static __global__ void top_k_radix_select( __syncthreads(); if (tid == 0) { - top_k_radix_state state = states[row]; + top_k_parallel_radix_state state = states[row]; int bin = NBINS - 1; while (bin > 0 && histogram[bin] < state.rank) { state.rank -= histogram[bin--]; @@ -233,7 +957,7 @@ static __global__ void top_k_radix_select( } } -static __global__ void top_k_radix_reset_counters(top_k_radix_state * states, int nrows) { +static __global__ void top_k_parallel_radix_reset_counters(top_k_parallel_radix_state * states, int nrows) { const int row = blockIdx.x * blockDim.x + threadIdx.x; if (row < nrows) { states[row].greater_count = 0; @@ -242,10 +966,10 @@ static __global__ void top_k_radix_reset_counters(top_k_radix_state * states, in } template -static __global__ void top_k_radix_gather( +static __global__ void top_k_parallel_radix_gather( const float * __restrict__ src, int * __restrict__ dst, - top_k_radix_state * __restrict__ states, + top_k_parallel_radix_state * __restrict__ states, int ncols, int k, int blocks_per_row) { @@ -254,7 +978,7 @@ static __global__ void top_k_radix_gather( const int tid = threadIdx.x; const float * row_src = src + (size_t) row * ncols; int * row_dst = dst + (size_t) row * k; - top_k_radix_state * state = &states[row]; + top_k_parallel_radix_state * state = &states[row]; for (int col = row_block * BLOCK_SIZE + tid; col < ncols; @@ -272,7 +996,7 @@ static __global__ void top_k_radix_gather( } } -static void top_k_radix_cuda( +static void top_k_parallel_radix_cuda( ggml_cuda_pool & pool, const float * src, int * dst, int ncols, int nrows, int k, cudaStream_t stream) { constexpr int BLOCK_SIZE = 256; @@ -280,29 +1004,48 @@ static void top_k_radix_cuda( constexpr int NBINS = 1 << RADIX_BITS; const int blocks_per_row = std::min((ncols + 1023) / 1024, 64); - ggml_cuda_pool_alloc states_alloc(pool, nrows); + ggml_cuda_pool_alloc states_alloc(pool, nrows); ggml_cuda_pool_alloc histograms_alloc(pool, (size_t) nrows * blocks_per_row * NBINS); - top_k_radix_state * states = states_alloc.get(); + top_k_parallel_radix_state * states = states_alloc.get(); int * histograms = histograms_alloc.get(); - top_k_radix_init<<<(nrows + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(states, nrows, k); + top_k_parallel_radix_init<<<(nrows + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(states, nrows, k); const dim3 row_grid(blocks_per_row * nrows); for (int shift = 32 - RADIX_BITS; shift >= 0; shift -= RADIX_BITS) { - top_k_radix_histogram + top_k_parallel_radix_histogram <<>>( src, states, histograms, ncols, blocks_per_row, shift); - top_k_radix_select + top_k_parallel_radix_select <<>>(histograms, states, blocks_per_row, shift); } - top_k_radix_reset_counters + top_k_parallel_radix_reset_counters <<<(nrows + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(states, nrows); - top_k_radix_gather + top_k_parallel_radix_gather <<>>( src, dst, states, ncols, k, blocks_per_row); } +static bool top_k_use_small_kernel(int ncols, int nrows, int k) { + if (k == 1) { + return true; + } +#if HIP_VERSION >= 71500000 + if (ncols <= 1024) { + return true; + } + const uint64_t elements = (uint64_t) ncols * nrows; + if (k <= 32) { + return elements <= (1U << 20); + } + return nrows == 1 && ncols <= (1U << 17); +#else + GGML_UNUSED(ncols); + GGML_UNUSED(nrows); + return false; +#endif +} #endif // !defined(GGML_CUDA_USE_CUB) && defined(GGML_USE_HIP) void ggml_cuda_op_top_k(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { @@ -359,8 +1102,10 @@ void ggml_cuda_op_top_k(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { } #else // GGML_CUDA_USE_CUB #if defined(GGML_USE_HIP) - if (ncols > 1024) { - top_k_radix_cuda(pool, src0_d, dst_d, ncols, nrows, k, stream); + if (top_k_use_small_kernel(ncols, nrows, k)) { + top_k_small_cuda(pool, src0_d, dst_d, ncols, nrows, k, stream); + } else if (ncols > 1024) { + top_k_parallel_radix_cuda(pool, src0_d, dst_d, ncols, nrows, k, stream); } else { #endif // defined(GGML_USE_HIP) ggml_cuda_pool_alloc temp_dst_alloc(pool, ncols * nrows); diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index bdb42bc9a551..ce17b659e4a1 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -417,7 +417,16 @@ template static __device__ __forceinline__ float vec_dot_q5_1_q8_1_imp return sumi*d5d8 + m5s8 / (QI5_1 / vdr); } +#if defined(RDNA4) || defined(RDNA3_0) +// VDR=4 measured on gfx1200/gfx1201 (RX 9000): decode -10..-33% on the +// compute-bound shapes, neutral on the DRAM-bound lm_head. RDNA3_0 (gfx1100- +// gfx1103, RX 7900 XTX) verified 2026-08-28: tg128 123.74 -> 127.7x (+3.x%), +// PPL 24.4430 vs 24.44xx (near-lossless), greedy byte-identical. RDNA3_5 +// (gfx115x) keeps VDR=2 pending verification on those GPUs. +#define VDR_Q8_0_Q8_1_MMVQ 4 +#else #define VDR_Q8_0_Q8_1_MMVQ 2 +#endif #define VDR_Q8_0_Q8_1_MMQ 8 template static __device__ __forceinline__ T vec_dot_q8_0_q8_1_impl( @@ -678,7 +687,11 @@ static __device__ __forceinline__ float vec_dot_q3_K_q8_1_impl_mmq( return d3*d8 * sumi; } +#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4) +#define VDR_Q4_K_Q8_1_MMVQ 4 +#else #define VDR_Q4_K_Q8_1_MMVQ 2 +#endif #define VDR_Q4_K_Q8_1_MMQ 8 // contiguous v/x values @@ -706,6 +719,34 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq( return dm4f.x*sumf_d - dm4f.y*sumf_m; } +// VDR=4 variant: processes two adjacent 16-element chunks (32 elements). +// Both chunks share the q8_1 block pair, the two sub-scales, the two mins and +// the d8 values, so the loads are amortized over twice the dp4a work of the +// VDR=2 kernel. +static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq4( + const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc, + const uint8_t * __restrict__ m, const half2 & dm4, const float * __restrict__ d8) { + + float sumf_d = 0.0f; + float sumf_m = 0.0f; + +#pragma unroll + for (int i = 0; i < 2*QR4_K; ++i) { + const int v0i = (v[2*(i>>1)] >> (4*(i&1))) & 0x0F0F0F0F; + const int v1i = (v[2*(i>>1) + 1] >> (4*(i&1))) & 0x0F0F0F0F; + + const int dot1 = ggml_cuda_dp4a(v1i, u[2*i+1], ggml_cuda_dp4a(v0i, u[2*i+0], 0)); // SIMD dot product + const int dot2 = ggml_cuda_dp4a(0x01010101, u[2*i+1], ggml_cuda_dp4a(0x01010101, u[2*i+0], 0)); // sum of u + + sumf_d += d8[i&1] * (dot1 * sc[i&1]); + sumf_m += d8[i&1] * (dot2 * m[i&1]); // multiply constant part of q4_K with sum of q8_1 values + } + + const float2 dm4f = __half22float2(dm4); + + return dm4f.x*sumf_d - dm4f.y*sumf_m; +} + // contiguous v/x + u/y values static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_mmq( const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc, @@ -734,7 +775,11 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_mmq( return dm4f.x*sumf_d - dm4f.y*sumf_m; } +#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4) +#define VDR_Q5_K_Q8_1_MMVQ 4 +#else #define VDR_Q5_K_Q8_1_MMVQ 2 +#endif #define VDR_Q5_K_Q8_1_MMQ 8 // contiguous v/x values @@ -769,6 +814,41 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_vmmq( return dm5f.x*sumf_d - dm5f.y*sumf_m; } +// VDR=4 variant: processes two adjacent 16-element chunks (32 elements). +// Both chunks share the q8_1 block pair, the two sub-scales, the two mins and +// the d8 values, so the loads are amortized over twice the dp4a work of the +// VDR=2 kernel. +static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_vmmq4( + const int * __restrict__ vl, const int * __restrict__ vh, const int * __restrict__ u, const uint8_t * __restrict__ sc, + const uint8_t * __restrict__ m, const half2 & dm5, const float * __restrict__ d8) { + + float sumf_d = 0.0f; + float sumf_m = 0.0f; + +#pragma unroll + for (int i = 0; i < 2*QR5_K; ++i) { + const int vl0i = (vl[2*(i>>1)] >> (4*(i&1))) & 0x0F0F0F0F; + const int vl1i = (vl[2*(i>>1) + 1] >> (4*(i&1))) & 0x0F0F0F0F; + + const int vh0i = ((vh[2*(i>>1)] >> (i&1)) << 4) & 0x10101010; + const int vh1i = ((vh[2*(i>>1) + 1] >> (i&1)) << 4) & 0x10101010; + + const int v0i = vl0i | vh0i; + const int v1i = vl1i | vh1i; + + const int dot1 = ggml_cuda_dp4a(v0i, u[2*i+0], ggml_cuda_dp4a(v1i, u[2*i+1], 0)); // SIMD dot product + const int dot2 = ggml_cuda_dp4a(0x01010101, u[2*i+0], ggml_cuda_dp4a(0x01010101, u[2*i+1], 0)); // sum of u + + sumf_d += d8[i&1] * (dot1 * sc[i&1]); + sumf_m += d8[i&1] * (dot2 * m[i&1]); + + } + + const float2 dm5f = __half22float2(dm5); + + return dm5f.x*sumf_d - dm5f.y*sumf_m; +} + // contiguous v/x + u/y values static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_mmq( const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc, @@ -797,7 +877,11 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_mmq( return dm4f.x*sumf_d - dm4f.y*sumf_m; } +#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4) +#define VDR_Q6_K_Q8_1_MMVQ 2 +#else #define VDR_Q6_K_Q8_1_MMVQ 1 +#endif #define VDR_Q6_K_Q8_1_MMQ 8 // contiguous v/x values @@ -823,6 +907,35 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmvq( return d*sumf; } +// VDR=2 variant: processes two adjacent 8-element chunks (16 elements). +// Both chunks share the q8_1 block pair, the two sub-scales and the d8 values, +// so the loads are amortized over twice the dp4a work of the VDR=1 kernel. +static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmvq_vdr2( + const int & vl0, const int & vl1, const int & vh0, const int & vh1, + const int * __restrict__ u, const int8_t * __restrict__ scales, + const float & d, const float * __restrict__ d8) { + + float sumf = 0.0f; + +#pragma unroll + for (int i = 0; i < 2*QR6_K; ++i) { + const int sc = scales[4*(i&1)]; + + const int vl = (i < QR6_K) ? vl0 : vl1; + const int vh = (i < QR6_K) ? vh0 : vh1; + + const int vil = (vl >> (4*(i&1))) & 0x0F0F0F0F; + + const int vih = ((vh >> (4*(i&1))) << 4) & 0x30303030; + + const int vi = __vsubss4((vil | vih), 0x20202020); // vi = (vil | vih) - 32 + + sumf += d8[i&1] * (ggml_cuda_dp4a(vi, u[i], 0) * sc); // SIMD dot product + } + + return d*sumf; +} + // contiguous v/x + u/y values static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmq( const int * __restrict__ v, const int * __restrict__ u, const int8_t * __restrict__ sc, @@ -1262,6 +1375,59 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1( return vec_dot_q4_K_q8_1_impl_vmmq(v, u, sc, m, bq4_K->dm, d8); } +// VDR=4 entry point: iqs must be a multiple of 4 (the mmvq kernel strides kqs +// by VDR). Processes 32 elements per call, splitting the ql/u loads over two +// 16-element chunks that share the q8_1 block pair, the sub-scale/min pair and +// the d8 values. +static __device__ __forceinline__ float vec_dot_q4_K_q8_1_vdr4( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_q4_K * bq4_K = (const block_q4_K *) vbq + kbx; + + int v[4]; + int u[4*QR4_K]; + float d8[2*QR4_K]; + + // iqs is in 0,4..28. The two 16-element chunks at iqs and iqs+2 land in the + // same (iqs/4, (iqs%4)/2) cell, so bq8_offset and the scale pair are shared. + const int bq8_offset = QR4_K * ((iqs/2) / (QI8_1/2)); + const int i8 = (iqs/2) % 4; + + const int * q4 = (const int *)(bq4_K->qs + 16 * bq8_offset + 4 * i8); + v[0] = q4[0]; + v[1] = q4[4]; + v[2] = q4[1]; + v[3] = q4[5]; + + const uint16_t * scales = (const uint16_t *)bq4_K->scales; + uint16_t aux[2]; + const int j = bq8_offset/2; + if (j < 2) { + aux[0] = scales[j+0] & 0x3f3f; + aux[1] = scales[j+2] & 0x3f3f; + } else { + aux[0] = ((scales[j+2] >> 0) & 0x0f0f) | ((scales[j-2] & 0xc0c0) >> 2); + aux[1] = ((scales[j+2] >> 4) & 0x0f0f) | ((scales[j-0] & 0xc0c0) >> 2); + } + const uint8_t * sc = (const uint8_t *)aux; + const uint8_t * m = sc + 2; + + // i8 is even, so i8 and i8+1 are adjacent int32 groups in each q8_1 block + d8[0] = __low2float(bq8_1[bq8_offset + 0].ds); + d8[1] = __low2float(bq8_1[bq8_offset + 1].ds); + +#pragma unroll + for (int i = 0; i < 2*QR4_K; ++i) { + const block_q8_1 * bq8i = bq8_1 + bq8_offset + (i&1); + + const int * q8 = (const int *)bq8i->qs + i8 + (i>>1); + u[2*i+0] = q8[0]; + u[2*i+1] = q8[4]; + } + + return vec_dot_q4_K_q8_1_impl_vmmq4(v, u, sc, m, bq4_K->dm, d8); +} + static __device__ __forceinline__ float vec_dot_q5_K_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { @@ -1308,6 +1474,64 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1( return vec_dot_q5_K_q8_1_impl_vmmq(vl, vh, u, sc, m, bq5_K->dm, d8); } +// VDR=4 entry point: iqs must be a multiple of 4 (the mmvq kernel strides kqs +// by VDR). Processes 32 elements per call, splitting the ql/qh/u loads over two +// 16-element chunks that share the q8_1 block pair, the sub-scale/min pair and +// the d8 values. +static __device__ __forceinline__ float vec_dot_q5_K_q8_1_vdr4( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_q5_K * bq5_K = (const block_q5_K *) vbq + kbx; + + int vl[4]; + int vh[4]; + int u[4*QR5_K]; + float d8[2*QR5_K]; + + const int bq8_offset = QR5_K * ((iqs/2) / (QI8_1/2)); + const int i8 = (iqs/2) % 4; + + const int * ql = (const int *)(bq5_K->qs + 16 * bq8_offset + 4 * i8); + const int * qh = (const int *)(bq5_K->qh + 4 * i8); + + vl[0] = ql[0]; + vl[1] = ql[4]; + vl[2] = ql[1]; + vl[3] = ql[5]; + + vh[0] = qh[0] >> bq8_offset; + vh[1] = qh[4] >> bq8_offset; + vh[2] = qh[1] >> bq8_offset; + vh[3] = qh[5] >> bq8_offset; + + const uint16_t * scales = (const uint16_t *)bq5_K->scales; + uint16_t aux[2]; + const int j = bq8_offset/2; + if (j < 2) { + aux[0] = scales[j+0] & 0x3f3f; + aux[1] = scales[j+2] & 0x3f3f; + } else { + aux[0] = ((scales[j+2] >> 0) & 0x0f0f) | ((scales[j-2] & 0xc0c0) >> 2); + aux[1] = ((scales[j+2] >> 4) & 0x0f0f) | ((scales[j-0] & 0xc0c0) >> 2); + } + const uint8_t * sc = (const uint8_t *)aux; + const uint8_t * m = sc + 2; + + d8[0] = __low2float(bq8_1[bq8_offset + 0].ds); + d8[1] = __low2float(bq8_1[bq8_offset + 1].ds); + +#pragma unroll + for (int i = 0; i < 2*QR5_K; ++i) { + const block_q8_1 * bq8i = bq8_1 + bq8_offset + (i&1); + + const int * q8 = (const int *)bq8i->qs + i8 + (i>>1); + u[2*i+0] = q8[0]; + u[2*i+1] = q8[4]; + } + + return vec_dot_q5_K_q8_1_impl_vmmq4(vl, vh, u, sc, m, bq5_K->dm, d8); +} + static __device__ __forceinline__ float vec_dot_q6_K_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { @@ -1334,6 +1558,41 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1( return vec_dot_q6_K_q8_1_impl_mmvq(vl, vh, u, scales, bq6_K->d, d8); } +// VDR=2 entry point: iqs must be even (the mmvq kernel strides kqs by VDR). +// Processes 16 elements per call, splitting the ql/qh/u loads over two chunks. +static __device__ __forceinline__ float vec_dot_q6_K_q8_1_vdr2( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_q6_K * bq6_K = (const block_q6_K *) vbq + kbx; + + // both chunks share these offsets: iqs+1 lands in the same (iqs/16, (iqs%16)/8) cell + const int bq8_offset = 2 * QR6_K * (iqs / (QI6_K/2)) + (iqs % (QI6_K/2)) / (QI6_K/4); + const int scale_offset = (QI6_K/4) * (iqs / (QI6_K/2)) + (iqs % (QI6_K/2)) / (QI6_K/8); + const int vh_shift = 2 * ((iqs % (QI6_K/2)) / (QI6_K/4)); + const int vh_idx = (QI6_K/4) * (iqs / (QI6_K/2)) + iqs % (QI6_K/4); + + const int vl0 = get_int_b4(bq6_K->ql, iqs); + const int vl1 = get_int_b4(bq6_K->ql, iqs + 1); + const int vh0 = get_int_b2(bq6_K->qh, vh_idx) >> vh_shift; + const int vh1 = get_int_b2(bq6_K->qh, vh_idx + 1) >> vh_shift; + + const int8_t * scales = bq6_K->scales + scale_offset; + + int u[2*QR6_K]; + float d8[QR6_K]; + + // iqs is even, so iqs%QI8_1 and (iqs%QI8_1)+1 are adjacent int32 groups in the block + const int i8 = iqs % QI8_1; + u[0] = get_int_b4(bq8_1[bq8_offset + 0].qs, i8); + u[1] = get_int_b4(bq8_1[bq8_offset + 2].qs, i8); + u[2] = get_int_b4(bq8_1[bq8_offset + 0].qs, i8 + 1); + u[3] = get_int_b4(bq8_1[bq8_offset + 2].qs, i8 + 1); + d8[0] = __low2float(bq8_1[bq8_offset + 0].ds); + d8[1] = __low2float(bq8_1[bq8_offset + 2].ds); + + return vec_dot_q6_K_q8_1_impl_mmvq_vdr2(vl0, vl1, vh0, vh1, u, scales, bq6_K->d, d8); +} + #define VDR_IQ2_XXS_Q8_1_MMVQ 2 #define VDR_IQ2_XXS_Q8_1_MMQ 2 diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index e3db6816757b..47d416ed56ca 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -10083,6 +10083,12 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q8_0, GGML_TYPE_F32, 128, 128, false, 8192, 1, 5120)); // Llama-4-Maverick-17B-128E-PAB-Q8_0 test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 8192, 1, 5120, {128, 1}, {1, 1})); test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 8192, 512, 5120, {128, 1}, {1, 1})); + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 17408, 512, 5120, {1, 1}, {1, 1})); // q8_0 at ffn shape + + // Qwen3.6-27B Q6_K prefill shapes (perf tuning targets): + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 17408, 512, 5120, {1, 1}, {1, 1})); // ffn_up/ffn_gate + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 512, 17408, {1, 1}, {1, 1})); // ffn_out + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 512, 5120, {1, 1}, {1, 1})); // attn qkv #endif for (ggml_type type_a : all_types) { @@ -10920,7 +10926,8 @@ static std::vector> make_test_cases_eval() { for (float max_bias : { 0.0f, 8.0f }) { if (!mask && max_bias > 0.0f) continue; for (float logit_softcap : {0.0f, 10.0f}) { - if (hsk != 128 && logit_softcap != 0.0f) continue; + // The mma kernel instantiates logit_softcap for heads 128/256/512 only. + if (hsk != 128 && hsk != 256 && hsk != 512 && logit_softcap != 0.0f) continue; for (int nh : { 1, 4 }) { if (nh == 1 && hsk != 320 && hsk != 576) continue; for (int nr3 : { 1, 3, }) { @@ -11454,6 +11461,30 @@ static std::vector> make_test_cases_perf() { GGML_TYPE_F32, {n_kv, 512, 64, 1}, false, {2, 1, 0, 3})); } + // Qwen3.6-27B Q6_K prefill shapes: + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 17408, 512, 5120, {1, 1}, {1, 1})); // ffn_up/ffn_gate + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 5120, 512, 17408, {1, 1}, {1, 1})); // ffn_out + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 512, 5120, {1, 1}, {1, 1})); // attn qkv + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 17408, 512, 5120, {1, 1}, {1, 1})); // q8_0 at ffn shape + // f16/f32 references at the same shapes: + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 17408, 512, 5120, {1, 1}, {1, 1})); + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 5120, 512, 17408, {1, 1}, {1, 1})); + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 17408, 512, 5120, {1, 1}, {1, 1})); + + // Flash attention perf: head 256 (default-on WMMA) vs 512/320/576 (DKQ != DV). + // A/B with GGML_CUDA_FA_WMMA_256=0 (forces the tile kernel for head > 128 on RDNA4). + for (const auto & fa : {std::tuple{256, 256, 8, 4, 256}, + {512, 512, 4, 8, 128}, + {320, 256, 4, 32, 128}, + {576, 512, 4, 4, 128}, + {192, 128, 4, 8, 128}, + {512, 512, 4, 8, 1}, + {576, 512, 4, 4, 1}}) { + const auto [hsk, hsv, nh, nr2, nb] = fa; + test_cases.emplace_back(new test_flash_attn_ext(hsk, hsv, nh, {nr2, 1}, 16384, nb, true, false, 0, 0, + GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); + } + // Conv2d: K=CRS=NPQ=4096 matmul performance uint32_t iwh_idx = 0; uint32_t kwh_idx = 1; @@ -11624,9 +11655,33 @@ static std::vector> make_test_cases_perf() { } } + // Qwen3.6-27B decode shapes (n=1, mmvq path), Q6_K vs Q8_0: + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 1, 5120, {1, 1}, {1, 1})); // ffn + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 1, 5120, {1, 1}, {1, 1})); // lm_head + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 12288, 1, 5120, {1, 1}, {1, 1})); // qkv + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 1, 5120, {1, 1}, {1, 1})); // ssm z + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 1024, 1, 5120, {1, 1}, {1, 1})); // kv proj + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 10240, 1, 5120, {1, 1}, {1, 1})); // ffn + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 248320, 1, 5120, {1, 1}, {1, 1})); // lm_head + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 12288, 1, 5120, {1, 1}, {1, 1})); // qkv + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 6144, 1, 5120, {1, 1}, {1, 1})); // ssm z + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 1024, 1, 5120, {1, 1}, {1, 1})); // kv proj + + // same shapes, Q4_K/Q5_K decode (n=1, mmvq) and prefill (mmq) rows: + for (ggml_type type_a : {GGML_TYPE_Q4_K, GGML_TYPE_Q5_K}) { + test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 10240, 1, 5120, {1, 1}, {1, 1})); // ffn + test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 248320, 1, 5120, {1, 1}, {1, 1})); // lm_head + test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 12288, 1, 5120, {1, 1}, {1, 1})); // qkv + test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 6144, 1, 5120, {1, 1}, {1, 1})); // ssm z + test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 1024, 1, 5120, {1, 1}, {1, 1})); // kv proj + for (int bs : {16, 128, 512}) { + test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 10240, bs, 5120, {1, 1}, {1, 1})); // prefill + } + } + // qwen3-30b-a3b - for (int bs : {1, 4, 8, 32, 64, 128, 256, 512}) { - for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) { + for (int bs : {1, 4, 5, 6, 7, 8, 32, 64, 128, 256, 512}) { + for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) { for (ggml_type type_b : {GGML_TYPE_F32}) { test_cases.emplace_back(new test_mul_mat_id(type_a, type_b, 128, 8, false, 768, bs, 2048)); test_cases.emplace_back(new test_mul_mat_id_fusion(type_a, type_b, 128, 8, false, 768, bs, 2048, 1)); @@ -11634,8 +11689,8 @@ static std::vector> make_test_cases_perf() { } } - for (int bs : {1, 4, 8, 32, 64, 128, 256, 512}) { - for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) { + for (int bs : {1, 4, 5, 6, 7, 8, 32, 64, 128, 256, 512}) { + for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) { for (ggml_type type_b : {GGML_TYPE_F32}) { test_cases.emplace_back(new test_mul_mat_id(type_a, type_b, 32, 4, false, 1792, bs, 2048)); test_cases.emplace_back(new test_mul_mat_id_fusion(type_a, type_b, 32, 4, false, 1792, bs, 2048, 1));