diff --git a/pkgs-set/pkgs/llama-cpp-bonsai-mmq-pipeline.patch b/pkgs-set/pkgs/llama-cpp-bonsai-mmq-pipeline.patch new file mode 100644 index 0000000..2a1d1bb --- /dev/null +++ b/pkgs-set/pkgs/llama-cpp-bonsai-mmq-pipeline.patch @@ -0,0 +1,339 @@ +diff --git a/ggml/src/ggml-cuda/cp-async.cuh b/ggml/src/ggml-cuda/cp-async.cuh +index 63d0c482..a9b47bd0 100644 +--- a/ggml/src/ggml-cuda/cp-async.cuh ++++ b/ggml/src/ggml-cuda/cp-async.cuh +@@ -55,3 +55,23 @@ static __device__ __forceinline__ void cp_async_wait_all() { + NO_DEVICE_CODE; + #endif // CP_ASYNC_AVAILABLE + } ++ ++// Closes the group of cp.async copies issued so far by this thread. ++static __device__ __forceinline__ void cp_async_commit_group() { ++#ifdef CP_ASYNC_AVAILABLE ++ asm volatile("cp.async.commit_group;"); ++#else ++ NO_DEVICE_CODE; ++#endif // CP_ASYNC_AVAILABLE ++} ++ ++// Makes each thread wait until at most n of its committed cp.async groups are still pending. ++// Like cp_async_wait_all, this does NOT synchronize threads. ++template ++static __device__ __forceinline__ void cp_async_wait_group() { ++#ifdef CP_ASYNC_AVAILABLE ++ asm volatile("cp.async.wait_group %0;" : : "n"(n)); ++#else ++ NO_DEVICE_CODE; ++#endif // CP_ASYNC_AVAILABLE ++} +diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh +index 91087109..3b70f694 100644 +--- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh ++++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh +@@ -262,96 +262,125 @@ template static __device__ __forceinline_ + } + + #if !defined(GGML_USE_HIP) +-template +-static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_ptq1_0(const char * __restrict__ x, +- int * __restrict__ x_tile, +- const int kbx0, +- const int i_max, +- const int stride) { +- constexpr int warp_size = ggml_cuda_get_physical_warp_size(); +- constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; +- constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); +- constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); ++// The PTQ1_0 x tile is loaded in two steps so the global loads for the next k-iteration can be issued before the ++// current one is computed: fetch() reads the packed words and scales into registers, store() decodes them into the tile. ++template struct ggml_cuda_mmq_ptq1_0_x_stage { ++ static constexpr int warp_size = ggml_cuda_get_physical_warp_size(); ++ static constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; ++ static constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); ++ static constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); ++ ++ static constexpr int blocks_per_iter = MMQ_ITER_K / QK_PTQ1_0; ++ static constexpr int threads_per_block = 8; ++ static constexpr int threads_per_row = blocks_per_iter * threads_per_block; ++ static constexpr int nrows = warp_size / threads_per_row; ++ static constexpr int scale_entries_per_block = QK_PTQ1_0 / QK8_1; ++ static constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; ++ static constexpr int rows_per_warp = warp_size / scale_entries_per_row; ++ static constexpr int nqs = I / (nrows * nwarps); ++ static constexpr int nd = I / (rows_per_warp * nwarps); ++ ++ uint32_t packed[nqs]; ++ ggml_half d[nd]; ++ ++ static __device__ __forceinline__ int row_qs(const int i0, const int i_max) { ++ const int i = i0 + threadIdx.y * nrows + threadIdx.x / threads_per_row; ++ return fallback ? min(i, i_max) : i; ++ } + ++ static __device__ __forceinline__ int row_d(const int i0, const int i_max) { ++ const int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / scale_entries_per_row; ++ return fallback ? min(i, i_max) : i; ++ } ++ ++ __device__ __forceinline__ void fetch(const char * __restrict__ x, const int kbx0, const int i_max, const int stride) { ++ const int txi = threadIdx.x % threads_per_row; ++ const int kbx = txi / threads_per_block; ++ const int lane = txi % threads_per_block; ++# pragma unroll ++ for (int n = 0; n < nqs; ++n) { ++ const block_ptq1_0 * bxi = (const block_ptq1_0 *) x + kbx0 + row_qs(n * nrows * nwarps, i_max) * stride + kbx; ++ packed[n] = get_int_b4(bxi->qs, lane < 7 ? lane : 6); ++ } ++ const int scale_block = (threadIdx.x % scale_entries_per_row) / scale_entries_per_block; ++# pragma unroll ++ for (int n = 0; n < nd; ++n) { ++ const block_ptq1_0 * bxi = ++ (const block_ptq1_0 *) x + kbx0 + row_d(n * nwarps * rows_per_warp, i_max) * stride + scale_block; ++ d[n] = bxi->d; ++ } ++ } ++ ++ __device__ __forceinline__ void store(int * __restrict__ x_tile, const int i_max) const { + # if defined(TURING_MMA_AVAILABLE) +- int * x_qs = (int *) x_tile; +- float * x_df = (float *) (x_qs + 2 * MMQ_TILE_NE_K); ++ int * x_qs = (int *) x_tile; ++ float * x_df = (float *) (x_qs + 2 * MMQ_TILE_NE_K); + # else +- constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); +- int * x_qs = (int *) x_tile; +- float * x_df = (float *) (x_qs + txs.qs); ++ constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); ++ int * x_qs = (int *) x_tile; ++ float * x_df = (float *) (x_qs + txs.qs); + # endif +- +- constexpr int blocks_per_iter = MMQ_ITER_K / QK_PTQ1_0; +- constexpr int threads_per_block = 8; +- constexpr int threads_per_row = blocks_per_iter * threads_per_block; +- constexpr int nrows = warp_size / threads_per_row; +- +- const int txi = threadIdx.x % threads_per_row; +- const int kbx = txi / threads_per_block; +- const int lane = txi % threads_per_block; ++ const int txi = threadIdx.x % threads_per_row; ++ const int kbx = txi / threads_per_block; ++ const int lane = txi % threads_per_block; + + # pragma unroll +- for (int i0 = 0; i0 < I; i0 += nrows * nwarps) { +- int i = i0 + threadIdx.y * nrows + threadIdx.x / threads_per_row; +- if (fallback) { +- i = min(i, i_max); +- } +- +- const block_ptq1_0 * bxi = (const block_ptq1_0 *) x + kbx0 + i * stride + kbx; ++ for (int n = 0; n < nqs; ++n) { ++ const int i = row_qs(n * nrows * nwarps, i_max); + # if defined(TURING_MMA_AVAILABLE) +- int * row = x_qs + i * sram_stride + kbx * (QK_PTQ1_0 / 4); ++ int * row = x_qs + i * sram_stride + kbx * (QK_PTQ1_0 / 4); + # else +- int * row = x_qs + i * (2 * MMQ_TILE_NE_K + 1) + kbx * (QK_PTQ1_0 / 4); ++ int * row = x_qs + i * (2 * MMQ_TILE_NE_K + 1) + kbx * (QK_PTQ1_0 / 4); + # endif + +- // Branch-free unpack: all 8 lanes run the same 5-iteration trit loop on their own 32-bit word. Only smem store offsets differ, so the warp does not diverge. +- const uint32_t packed = get_int_b4(bxi->qs, lane < 7 ? lane : 6); +- uint32_t v_lo = __byte_perm(packed, 0, 0x4140); // bytes 0,1 as 16-bit lanes +- uint32_t v_hi = __byte_perm(packed, 0, 0x4342); // bytes 2,3 +- const bool full_lane = lane < 6; +- v_hi = full_lane ? v_hi : v_lo; // lane 6: both halves walk qh0/qh1 +- const int dst_base = lane < 4 ? lane : 16 + lane; // lanes 4,5 -> 20,21 +- const int dst_stride = lane < 4 ? 4 : 2; +- int q[5]; ++ // Branch-free unpack: all 8 lanes run the same 5-iteration trit loop on their own 32-bit word. Only smem store offsets differ, so the warp does not diverge. ++ uint32_t v_lo = __byte_perm(packed[n], 0, 0x4140); // bytes 0,1 as 16-bit lanes ++ uint32_t v_hi = __byte_perm(packed[n], 0, 0x4342); // bytes 2,3 ++ const bool full_lane = lane < 6; ++ v_hi = full_lane ? v_hi : v_lo; // lane 6: both halves walk qh0/qh1 ++ const int dst_base = lane < 4 ? lane : 16 + lane; // lanes 4,5 -> 20,21 ++ const int dst_stride = lane < 4 ? 4 : 2; ++ int q[5]; + # pragma unroll +- for (int t = 0; t < 5; ++t) { +- const uint32_t w_lo = v_lo * 3; +- const uint32_t w_hi = v_hi * 3; +- v_lo = w_lo & 0x00FF00FF; +- v_hi = w_hi & 0x00FF00FF; +- q[t] = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); +- if (full_lane) { +- row[dst_base + t * dst_stride] = q[t]; ++ for (int t = 0; t < 5; ++t) { ++ const uint32_t w_lo = v_lo * 3; ++ const uint32_t w_hi = v_hi * 3; ++ v_lo = w_lo & 0x00FF00FF; ++ v_hi = w_hi & 0x00FF00FF; ++ q[t] = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); ++ if (full_lane) { ++ row[dst_base + t * dst_stride] = q[t]; ++ } ++ } ++ if (lane == 6) { ++ // q[t] = {qh0.t, qh1.t, qh0.t, qh1.t}; the old layout is {qh0.t, qh1.t, qh0.t+1, qh1.t+1}. ++ row[30] = __byte_perm(q[0], q[1], 0x5410); ++ row[31] = __byte_perm(q[2], q[3], 0x5410); + } + } +- if (lane == 6) { +- // q[t] = {qh0.t, qh1.t, qh0.t, qh1.t}; the old layout is {qh0.t, qh1.t, qh0.t+1, qh1.t+1}. +- row[30] = __byte_perm(q[0], q[1], 0x5410); +- row[31] = __byte_perm(q[2], q[3], 0x5410); +- } +- } +- +- constexpr int scale_entries_per_block = QK_PTQ1_0 / QK8_1; +- constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; +- constexpr int rows_per_warp = warp_size / scale_entries_per_row; +- const int ksx = threadIdx.x % scale_entries_per_row; +- const int scale_block = ksx / scale_entries_per_block; + ++ const int ksx = threadIdx.x % scale_entries_per_row; + # pragma unroll +- for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { +- int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / scale_entries_per_row; +- if (fallback) { +- i = min(i, i_max); +- } +- +- const block_ptq1_0 * bxi = (const block_ptq1_0 *) x + kbx0 + i * stride + scale_block; ++ for (int n = 0; n < nd; ++n) { ++ const int i = row_d(n * nwarps * rows_per_warp, i_max); + # if defined(TURING_MMA_AVAILABLE) +- x_df[i * sram_stride + ksx] = bxi->d; ++ x_df[i * sram_stride + ksx] = d[n]; + # else +- x_df[i * (2 * MMQ_TILE_NE_K / QI8_0) + i / (QI8_0 / 2) + ksx] = bxi->d; ++ x_df[i * (2 * MMQ_TILE_NE_K / QI8_0) + i / (QI8_0 / 2) + ksx] = d[n]; + # endif ++ } + } ++}; ++ ++template ++static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_ptq1_0(const char * __restrict__ x, ++ int * __restrict__ x_tile, ++ const int kbx0, ++ const int i_max, ++ const int stride) { ++ ggml_cuda_mmq_ptq1_0_x_stage stage; ++ stage.fetch(x, kbx0, i_max, stride); ++ stage.store(x_tile, i_max); + } + #endif + +diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp +index 462c2b71..2e610278 100644 +--- a/tests/test-backend-ops.cpp ++++ b/tests/test-backend-ops.cpp +@@ -10604,6 +10604,14 @@ static std::vector> make_test_cases_perf() { + } + } + ++ // Ternary Bonsai 2 27B (qwen35, PTQ1_0) projections at the prefill ubatch (MMQ path) ++ for (ggml_type type_a : {GGML_TYPE_PTQ1_0, GGML_TYPE_Q4_0}) { ++ for (auto [m, k] : std::vector>{{17408, 5120}, {5120, 17408}, {10240, 5120}, {5120, 6144}, ++ {6144, 5120}, {12288, 5120}, {1024, 5120}}) { ++ test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, m, 1024, k, {1, 1}, {1, 1})); ++ } ++ } ++ + // 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}) { +diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh +index ceb676ff..47de164c 100644 +--- a/ggml/src/ggml-cuda/mmq.cuh ++++ b/ggml/src/ggml-cuda/mmq.cuh +@@ -930,7 +930,9 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( + constexpr int tile_y_stride = GGML_PAD(J*MMQ_TILE_Y_K, nwarps*warp_size); + #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK + constexpr bool async_buffer_y = +- type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0; ++ type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0 || type == GGML_TYPE_PTQ1_0; ++#elif defined(CP_ASYNC_AVAILABLE) && defined(TURING_MMA_AVAILABLE) ++ constexpr bool async_buffer_y = type == GGML_TYPE_PTQ1_0; + #else + constexpr bool async_buffer_y = false; + #endif +@@ -951,6 +953,63 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( + + constexpr int sz = sizeof(block_q8_1_mmq) / sizeof(int); + ++#if !defined(GGML_USE_HIP) && defined(CP_ASYNC_AVAILABLE) && defined(TURING_MMA_AVAILABLE) ++ if constexpr (type == GGML_TYPE_PTQ1_0) { ++ // Software pipeline over the two y buffers: the y half for the next vec_dot is copied with cp.async while the ++ // current one is computed, and the next x tile is fetched into registers before the first vec_dot and only ++ // decoded into shared memory after the second, so no global load latency is exposed between k-iterations. ++ const int tid = threadIdx.y*warp_size + threadIdx.x; ++ const auto stage_y = [&](int * dst, const int kb) { ++ const char * src = reinterpret_cast(y + ncols_y * (kb * qk / ne_block) * sz); ++ char * dst_bytes = reinterpret_cast(dst); ++#pragma unroll ++ for (int byte0 = 16*tid; byte0 < J*MMQ_TILE_Y_K*int(sizeof(int)); byte0 += 16*nwarps*warp_size) { ++ cp_async_cg_16<256>(ggml_cuda_cvta_generic_to_shared(dst_bytes + byte0), src + byte0); ++ } ++ cp_async_commit_group(); ++ }; ++ ++ ggml_cuda_mmq_ptq1_0_x_stage x_stage; ++ stage_y(tile_y, kb0_start); ++ x_stage.fetch(x, offset_x + kb0_start, tile_x_max_i, stride_row_x); ++ x_stage.store(tile_x, tile_x_max_i); ++ ++ for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { ++ const bool next = kb0 + blocks_per_iter < kb0_stop; ++ ++ stage_y(tile_y_next, kb0 + blocks_per_iter/2); ++ cp_async_wait_group<1>(); ++ __syncthreads(); ++ ++ if (next) { ++ x_stage.fetch(x, offset_x + kb0 + blocks_per_iter, tile_x_max_i, stride_row_x); ++ } ++ vec_dot(tile_x, tile_y, sum, 0); ++ ++ cp_async_wait_group<0>(); ++ __syncthreads(); ++ ++ if (next) { ++ stage_y(tile_y, kb0 + blocks_per_iter); ++ } ++ vec_dot(tile_x, tile_y_next, sum, MMQ_TILE_NE_K); ++ ++ __syncthreads(); ++ ++ if (next) { ++ x_stage.store(tile_x, tile_x_max_i); ++ } ++ } ++ ++ if (fixup) { ++ write_back(sum, ids_dst, tmp_fixup + blockIdx.x*(J*I), y_scale, I, I, J); ++ } else { ++ write_back(sum, ids_dst, dst, y_scale, stride_col_dst, tile_x_max_i, tile_y_max_j); ++ } ++ return; ++ } ++#endif // !defined(GGML_USE_HIP) && defined(CP_ASYNC_AVAILABLE) && defined(TURING_MMA_AVAILABLE) ++ + for (int kb0 = kb0_start; kb0 < kb0_stop; kb0 += blocks_per_iter) { + if constexpr (async_buffer_y) { + const char * by0 = reinterpret_cast( +@@ -1488,8 +1547,9 @@ static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const i + const size_t nbs_x = ggml_cuda_mmq_get_nbytes_shared_x(config, cc); + const size_t nbs_y = config.J * (sizeof(block_q8_1_mmq)); + const size_t nbs_y_padded = GGML_PAD(nbs_y, config.nthreads*sizeof(int)); +- const bool async_buffer_y = cc == GGML_CUDA_CC_DGX_SPARK && +- (config.type == GGML_TYPE_Q1_0 || config.type == GGML_TYPE_Q2_0 || config.type == GGML_TYPE_PQ2_0); ++ const bool async_buffer_y = (cc == GGML_CUDA_CC_DGX_SPARK && ++ (config.type == GGML_TYPE_Q1_0 || config.type == GGML_TYPE_Q2_0 || config.type == GGML_TYPE_PQ2_0)) || ++ (config.type == GGML_TYPE_PTQ1_0 && cp_async_available(cc) && turing_mma_available(cc)); + const int y_buffers = async_buffer_y ? 2 : 1; + return nbs_ids + nbs_x + y_buffers*nbs_y_padded; + } diff --git a/pkgs-set/pkgs/llama-cpp-bonsai-mmq-y128.patch b/pkgs-set/pkgs/llama-cpp-bonsai-mmq-y128.patch new file mode 100644 index 0000000..c3086cd --- /dev/null +++ b/pkgs-set/pkgs/llama-cpp-bonsai-mmq-y128.patch @@ -0,0 +1,177 @@ +diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh +index be091d27..65ca3f3d 100644 +--- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh ++++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh +@@ -179,7 +179,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma( + + float dB; + const int j = j0 + tile_C::get_j(0); +- if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4) { ++ if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4 || ds_layout == MMQ_Q8_1_DS_LAYOUT_D1) { + dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + } else { + dB = __low2float(y_ds[j*MMQ_TILE_Y_K + k01/QI8_1]); +@@ -244,6 +244,39 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma( + } + } + ++ if constexpr (ds_layout == MMQ_Q8_1_DS_LAYOUT_D1) { ++ // x and y both have one scale per MMQ_TILE_NE_K ints (128 values), so the k01 mmas accumulate in int32 ++ // and the result is rescaled once. ++#pragma unroll ++ for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) { ++ tile_C C[ntx]; ++#pragma unroll ++ for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_0) { ++ tile_B B; ++ load_generic(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); ++#pragma unroll ++ for (int n = 0; n < ntx; ++n) { ++ mma(C[n], A[n][k01/QI8_0], B); ++ } ++ } ++ ++ float dB[tile_C::ne/2]; ++#pragma unroll ++ for (int l = 0; l < tile_C::ne/2; ++l) { ++ dB[l] = y_df[(j0 + tile_C::get_j(l))*MMQ_TILE_Y_K]; ++ } ++ ++#pragma unroll ++ for (int n = 0; n < ntx; ++n) { ++#pragma unroll ++ for (int l = 0; l < tile_C::ne; ++l) { ++ sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*dA[n][l/2][0]*dB[l%2]; ++ } ++ } ++ } ++ return; ++ } ++ + #pragma unroll + for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) { + #pragma unroll +@@ -257,7 +290,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma( + for (int l = 0; l < tile_C::ne/2; ++l) { + const int j = j0 + tile_C::get_j(l); + +- if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4) { ++ if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4 || ds_layout == MMQ_Q8_1_DS_LAYOUT_D1) { + dB[l] = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + } else { + dB[l] = __low2float(y_ds[j*MMQ_TILE_Y_K + k01/QI8_1]); +diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh +index 47de164c..99b05d86 100644 +--- a/ggml/src/ggml-cuda/mmq.cuh ++++ b/ggml/src/ggml-cuda/mmq.cuh +@@ -8,6 +8,13 @@ + + #define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. + #define MMQ_PTQ1_0_MAX_BATCH_SIZE (1 << 30) // MMQ at every batch unless GGML_CUDA_PTQ1_0_MMQ_MAX_BATCH lowers it ++ ++// PTQ1_0 MMQ quantizes src1 with one int8 scale per 128 values (one PTQ1_0 block) instead of one per 32, so the ++// tensor cores accumulate a whole block in int32 and the float rescale runs once per 128 instead of once per 32. ++// Changes results versus per-32 activation scales; set to 0 for the per-32 layout. ++#ifndef GGML_CUDA_PTQ1_0_MMQ_Y128 ++#define GGML_CUDA_PTQ1_0_MMQ_Y128 1 ++#endif + #define MMQ_ITER_K 256 + #define MMQ_ITER_K_FP4 512 + #define MMQ_NWARPS 8 +@@ -19,6 +26,7 @@ typedef void (*ggml_cuda_mmq_write_back_t)(const float * __restrict__ sum, const + + enum mmq_q8_1_ds_layout { + MMQ_Q8_1_DS_LAYOUT_D4, ++ MMQ_Q8_1_DS_LAYOUT_D1, // like D4, but one scale per 128 values, written to all four d4 slots + MMQ_Q8_1_DS_LAYOUT_DS4, + MMQ_Q8_1_DS_LAYOUT_D2S6, + }; +@@ -64,10 +72,11 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { + case GGML_TYPE_Q1_0: + case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: ++ return MMQ_Q8_1_DS_LAYOUT_D4; + #if !defined(GGML_USE_HIP) + case GGML_TYPE_PTQ1_0: ++ return GGML_CUDA_PTQ1_0_MMQ_Y128 ? MMQ_Q8_1_DS_LAYOUT_D1 : MMQ_Q8_1_DS_LAYOUT_D4; + #endif +- return MMQ_Q8_1_DS_LAYOUT_D4; + case GGML_TYPE_Q4_0: + case GGML_TYPE_Q4_1: + return MMQ_Q8_1_DS_LAYOUT_DS4; +@@ -756,7 +765,8 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func + case GGML_TYPE_PTQ1_0: + return ggml_cuda_mmq_util_funcs( + -1, ggml_cuda_mmq_load_tiles_ptq1_0, +- ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, ++ ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, + ggml_cuda_mmq_write_back_mma); + #endif + case GGML_TYPE_Q4_0: +diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu +index 8cbee011..ad4306b7 100644 +--- a/ggml/src/ggml-cuda/quantize.cu ++++ b/ggml/src/ggml-cuda/quantize.cu +@@ -497,7 +497,7 @@ static __global__ void quantize_mmq_q8_1( + const float * __restrict__ norm_weight = nullptr, + const float * __restrict__ norm_scale = nullptr) { + +- constexpr int vals_per_scale = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 64 : 32; ++ constexpr int vals_per_scale = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 64 : ds_layout == MMQ_Q8_1_DS_LAYOUT_D1 ? 128 : 32; + constexpr int vals_per_sum = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 16 : 32; + + const int64_t i0 = ((int64_t)blockDim.x*blockIdx.y + threadIdx.x)*4; +@@ -559,7 +559,7 @@ static __global__ void quantize_mmq_q8_1( + } + + float sum; +- if (ds_layout != MMQ_Q8_1_DS_LAYOUT_D4) { ++ if (ds_layout != MMQ_Q8_1_DS_LAYOUT_D4 && ds_layout != MMQ_Q8_1_DS_LAYOUT_D1) { + sum = xi.x + xi.y + xi.z + xi.w; + + // Calculate sums across vals_per_sum/4 threads. +@@ -627,6 +627,10 @@ void quantize_mmq_q8_1_swiglu_cuda( + quantize_mmq_q8_1<<>>( + x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate); + break; ++ case MMQ_Q8_1_DS_LAYOUT_D1: ++ quantize_mmq_q8_1<<>>( ++ x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate); ++ break; + case MMQ_Q8_1_DS_LAYOUT_DS4: + quantize_mmq_q8_1<<>>( + x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate); +@@ -655,6 +659,10 @@ void quantize_mmq_q8_1_rms_cuda( + quantize_mmq_q8_1<<>>( + x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, nullptr, weight, scale); + break; ++ case MMQ_Q8_1_DS_LAYOUT_D1: ++ quantize_mmq_q8_1<<>>( ++ x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, nullptr, weight, scale); ++ break; + case MMQ_Q8_1_DS_LAYOUT_DS4: + quantize_mmq_q8_1<<>>( + x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, nullptr, weight, scale); +@@ -719,6 +727,10 @@ void quantize_mmq_q8_1_cuda( + quantize_mmq_q8_1 + <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); + break; ++ case MMQ_Q8_1_DS_LAYOUT_D1: ++ quantize_mmq_q8_1 ++ <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); ++ break; + case MMQ_Q8_1_DS_LAYOUT_DS4: + quantize_mmq_q8_1 + <<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0); +@@ -749,6 +761,10 @@ void quantize_scatter_mmq_q8_1_cuda( + quantize_mmq_q8_1<<>>( + x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used); + break; ++ case MMQ_Q8_1_DS_LAYOUT_D1: ++ quantize_mmq_q8_1<<>>( ++ x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used); ++ break; + case MMQ_Q8_1_DS_LAYOUT_DS4: + quantize_mmq_q8_1<<>>( + x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used); diff --git a/pkgs-set/pkgs/llama-cpp-bonsai-prefill-fusions.patch b/pkgs-set/pkgs/llama-cpp-bonsai-prefill-fusions.patch new file mode 100644 index 0000000..f0a8b29 --- /dev/null +++ b/pkgs-set/pkgs/llama-cpp-bonsai-prefill-fusions.patch @@ -0,0 +1,322 @@ +diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu +index 467a84a9..281a3d97 100644 +--- a/ggml/src/ggml-cuda/fwht.cu ++++ b/ggml/src/ggml-cuda/fwht.cu +@@ -1,5 +1,6 @@ + #include "common.cuh" + #include "fwht.cuh" ++#include "unary.cuh" + + #include + +@@ -13,10 +14,19 @@ __device__ __forceinline__ float fwht_load(const half value) { + return __half2float(value); + } + +-template ++// the transform's input: src, or silu(src) * src2 for a SWIGLU folded into the load (same arithmetic as the GLU kernel) ++template ++__device__ __forceinline__ float fwht_value(const T * src, const T * src2, const int i) { ++ if constexpr (glu) { ++ return ggml_cuda_op_silu_single((float) src[i]) * (float) src2[i]; ++ } ++ return fwht_load(src[i]); ++} ++ ++template + __launch_bounds__(4*ggml_cuda_get_physical_warp_size(), 1) + __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, const float scale, +- const float * signs, const int n_blk) { ++ const float * signs, const int n_blk, const T * src2) { + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + + const int64_t r = (int64_t) blockIdx.x * blockDim.y + threadIdx.y; +@@ -27,6 +37,9 @@ __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, cons + + src += r * N; + dst += r * N; ++ if (glu) { ++ src2 += r * N; ++ } + + static constexpr int el_w = N / warp_size; + float reg[el_w]; +@@ -36,7 +49,7 @@ __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, cons + const float * signs_row = has_signs ? signs + (r % n_blk) * N : nullptr; + #pragma unroll + for (int i = 0; i < el_w; ++i) { +- reg[i] = fwht_load(src[i * warp_size + lane]) * scale; ++ reg[i] = fwht_value(src, src2, i * warp_size + lane) * scale; + if (has_signs) { + reg[i] *= signs_row[i * warp_size + lane]; + } +@@ -83,10 +96,10 @@ __global__ void fwht_cuda(const T * src, float * dst, const int64_t n_rows, cons + // row than the register path, so it is used only where that path cannot go. + #define FWHT_SMEM_THREADS 256 + +-template ++template + __launch_bounds__(FWHT_SMEM_THREADS, 1) + __global__ void fwht_cuda_smem(const T * src, float * dst, const int64_t n_rows, const float scale, +- const float * signs, const int n_blk) { ++ const float * signs, const int n_blk, const T * src2) { + __shared__ float s[N]; + + const int64_t r = blockIdx.x; +@@ -96,12 +109,15 @@ __global__ void fwht_cuda_smem(const T * src, float * dst, const int64_t n_rows, + + src += r * N; + dst += r * N; ++ if (glu) { ++ src2 += r * N; ++ } + + const float * signs_row = has_signs ? signs + (r % n_blk) * N : nullptr; + + ggml_cuda_pdl_sync(); + for (int i = threadIdx.x; i < N; i += FWHT_SMEM_THREADS) { +- float v = fwht_load(src[i]) * scale; ++ float v = fwht_value(src, src2, i) * scale; + if (has_signs) { + v *= signs_row[i]; + } +@@ -134,10 +150,10 @@ __global__ void fwht_cuda_smem(const T * src, float * dst, const int64_t n_rows, + // Both leave most of the GPU idle at these shapes. + #define FWHT_BLOCK_THREADS 256 + +-template ++template + __launch_bounds__(NT, 1) + __global__ void fwht_cuda_block(const T * src, float * dst, const int64_t n_rows, const float scale, +- const float * signs, const int n_blk) { ++ const float * signs, const int n_blk, const T * src2) { + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr int NE = N / NT; + static_assert(NE >= 1 && N % NT == 0 && NT % warp_size == 0, "bad FWHT block shape"); +@@ -151,6 +167,9 @@ __global__ void fwht_cuda_block(const T * src, float * dst, const int64_t n_rows + + src += r * N; + dst += r * N; ++ if (glu) { ++ src2 += r * N; ++ } + + const int tid = threadIdx.x; + const int lane = tid % warp_size; +@@ -161,7 +180,7 @@ __global__ void fwht_cuda_block(const T * src, float * dst, const int64_t n_rows + float reg[NE]; + #pragma unroll + for (int i = 0; i < NE; ++i) { +- reg[i] = fwht_load(src[i * NT + tid]) * scale; ++ reg[i] = fwht_value(src, src2, i * NT + tid) * scale; + if (has_signs) { + reg[i] *= signs_row[i * NT + tid]; + } +@@ -220,7 +239,7 @@ __global__ void fwht_cuda_block(const T * src, float * dst, const int64_t n_rows + template + static bool fwht_launch(ggml_backend_cuda_context & ctx, const T * src_d, float * dst_d, + const int n, const int64_t rows, const float scale, +- const float * signs, const int n_blk) { ++ const float * signs, const int n_blk, const T * src2_d = nullptr) { + const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size; + const int rows_per_block = 4; + const int64_t num_blocks = (rows + rows_per_block - 1) / rows_per_block; +@@ -233,10 +252,12 @@ static bool fwht_launch(ggml_backend_cuda_context & ctx, const T * src_d, float + switch (n) { + #define FWHT_CASE(NN) \ + case NN: \ +- if (signs) { \ +- ggml_cuda_kernel_launch(fwht_cuda, launch_params, src_d, dst_d, rows, scale, signs, n_blk); \ ++ if (src2_d) { \ ++ ggml_cuda_kernel_launch(fwht_cuda, launch_params, src_d, dst_d, rows, scale, signs, n_blk, src2_d); \ ++ } else if (signs) { \ ++ ggml_cuda_kernel_launch(fwht_cuda, launch_params, src_d, dst_d, rows, scale, signs, n_blk, nullptr); \ + } else { \ +- ggml_cuda_kernel_launch(fwht_cuda, launch_params, src_d, dst_d, rows, scale, nullptr, 1); \ ++ ggml_cuda_kernel_launch(fwht_cuda, launch_params, src_d, dst_d, rows, scale, nullptr, 1, nullptr); \ + } \ + return true; + FWHT_CASE(64) +@@ -252,10 +273,12 @@ static bool fwht_launch(ggml_backend_cuda_context & ctx, const T * src_d, float + case NN: { \ + const dim3 g((unsigned) rows, 1, 1), b(FWHT_SMEM_THREADS, 1, 1); \ + const ggml_cuda_kernel_launch_params lp = ggml_cuda_kernel_launch_params(g, b, 0, stream); \ +- if (signs) { \ +- ggml_cuda_kernel_launch(fwht_cuda_smem, lp, src_d, dst_d, rows, scale, signs, n_blk); \ ++ if (src2_d) { \ ++ ggml_cuda_kernel_launch(fwht_cuda_smem, lp, src_d, dst_d, rows, scale, signs, n_blk, src2_d); \ ++ } else if (signs) { \ ++ ggml_cuda_kernel_launch(fwht_cuda_smem, lp, src_d, dst_d, rows, scale, signs, n_blk, nullptr); \ + } else { \ +- ggml_cuda_kernel_launch(fwht_cuda_smem, lp, src_d, dst_d, rows, scale, nullptr, 1); \ ++ ggml_cuda_kernel_launch(fwht_cuda_smem, lp, src_d, dst_d, rows, scale, nullptr, 1, nullptr); \ + } \ + return true; \ + } +@@ -263,10 +286,12 @@ static bool fwht_launch(ggml_backend_cuda_context & ctx, const T * src_d, float + case NN: { \ + const dim3 g((unsigned) rows, 1, 1), b(FWHT_BLOCK_THREADS, 1, 1); \ + const ggml_cuda_kernel_launch_params lp = ggml_cuda_kernel_launch_params(g, b, 0, stream); \ +- if (signs) { \ +- ggml_cuda_kernel_launch(fwht_cuda_block, lp, src_d, dst_d, rows, scale, signs, n_blk); \ ++ if (src2_d) { \ ++ ggml_cuda_kernel_launch(fwht_cuda_block, lp, src_d, dst_d, rows, scale, signs, n_blk, src2_d); \ ++ } else if (signs) { \ ++ ggml_cuda_kernel_launch(fwht_cuda_block, lp, src_d, dst_d, rows, scale, signs, n_blk, nullptr); \ + } else { \ +- ggml_cuda_kernel_launch(fwht_cuda_block, lp, src_d, dst_d, rows, scale, nullptr, 1); \ ++ ggml_cuda_kernel_launch(fwht_cuda_block, lp, src_d, dst_d, rows, scale, nullptr, 1, nullptr); \ + } \ + return true; \ + } +@@ -337,3 +362,19 @@ bool ggml_cuda_op_fwht_signed(ggml_backend_cuda_context & ctx, const ggml_tensor + const ggml_tensor * signs, ggml_tensor * dst) { + return fwht_dispatch(ctx, src, dst, signs); + } ++ ++bool ggml_cuda_op_fwht_swiglu_signed(ggml_backend_cuda_context & ctx, const ggml_tensor * gate, const ggml_tensor * up, ++ const ggml_tensor * signs_t, ggml_tensor * dst) { ++ // F32 split SWIGLU only, both halves laid out like dst ++ if (gate->type != GGML_TYPE_F32 || up->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32 || ++ !ggml_is_contiguous(gate) || !ggml_is_contiguous(up) || !ggml_is_contiguous(dst) || ++ ggml_nelements(gate) != ggml_nelements(dst) || ggml_nelements(up) != ggml_nelements(dst)) { ++ return false; ++ } ++ const int n = dst->ne[0]; ++ if (signs_t->type != GGML_TYPE_F32 || !ggml_is_contiguous(signs_t) || signs_t->ne[0] % n != 0) { ++ return false; ++ } ++ return fwht_launch(ctx, (const float *) gate->data, (float *) dst->data, n, ggml_nelements(dst) / n, ++ 1 / sqrtf(n), (const float *) signs_t->data, signs_t->ne[0] / n, (const float *) up->data); ++} +diff --git a/ggml/src/ggml-cuda/fwht.cuh b/ggml/src/ggml-cuda/fwht.cuh +index 62b2f288..faf3e88b 100644 +--- a/ggml/src/ggml-cuda/fwht.cuh ++++ b/ggml/src/ggml-cuda/fwht.cuh +@@ -4,3 +4,6 @@ + bool ggml_cuda_op_fwht(ggml_backend_cuda_context & ctx, const ggml_tensor * src, ggml_tensor * dst); + bool ggml_cuda_op_fwht_signed(ggml_backend_cuda_context & ctx, const ggml_tensor * src, + const ggml_tensor * signs, ggml_tensor * dst); ++// the same, with a split F32 SWIGLU (silu(gate) * up) folded into the transform's load ++bool ggml_cuda_op_fwht_swiglu_signed(ggml_backend_cuda_context & ctx, const ggml_tensor * gate, const ggml_tensor * up, ++ const ggml_tensor * signs, ggml_tensor * dst); +diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu +index 7536a5d2..69b2652b 100644 +--- a/ggml/src/ggml-cuda/ggml-cuda.cu ++++ b/ggml/src/ggml-cuda/ggml-cuda.cu +@@ -2971,7 +2971,18 @@ static bool ggml_cuda_check_fusion_memory_ranges(const ggml_cgraph * cgraph, + const int node_count, + const int * out_nodes, + const int out_count, +- const bool is_topk_moe = false) { ++ const bool is_topk_moe = false, ++ const bool allow_exact_alias = false) { ++ // An exact alias is a destination that shares the base pointer, type and ++ // contiguous layout of a source, i.e. element k of the destination is ++ // element k of the source. Fusions whose kernels are elementwise within the ++ // unit they parallelize over can opt into accepting those; a partial ++ // overlap is still rejected for everyone. ++ auto is_exact_alias = [&](const ggml_tensor * a, const ggml_tensor * b) { ++ return allow_exact_alias && a->data == b->data && a->type == b->type && ++ ggml_are_same_shape(a, b) && ggml_is_contiguous(a) && ggml_is_contiguous(b); ++ }; ++ + auto nodes_overlap = [&](const ggml_tensor * a, const ggml_tensor * b) { + const int64_t a_start = (int64_t) a->data; + const int64_t a_end = a_start + ggml_backend_buft_get_alloc_size(a->buffer->buft, a); +@@ -3007,7 +3018,7 @@ static bool ggml_cuda_check_fusion_memory_ranges(const ggml_cgraph * cgraph, + continue; + } + +- if (nodes_overlap(dst, src)) { ++ if (nodes_overlap(dst, src) && !is_exact_alias(dst, src)) { + bool found = false; + + for (int k = node_idx; k < j; ++k) { +@@ -3359,13 +3370,20 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph + + // Preserve the residual sum for later graph consumers while normalizing it + // in the same launch. This removes a full read of the residual tensor. ++ // ++ // add_rms_norm_f32 uses no architecture-specific instructions, so this runs ++ // on every NVIDIA device rather than GB10 only. It is elementwise within a ++ // row -- each thread reads a[col] and b[col] before writing sum[col], and ++ // block_reduce synchronizes the block before dst[col] is written -- so the ++ // in-place residual add and the norm output reusing an input buffer are ++ // both safe, hence allow_exact_alias below. + if (node->op == GGML_OP_ADD && i + 2 < cgraph->n_nodes) { + ggml_tensor * rms_norm = cgraph->nodes[i + 1]; + ggml_tensor * mul = cgraph->nodes[i + 2]; + const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + const ggml_op ops[] = { GGML_OP_ADD, GGML_OP_RMS_NORM, GGML_OP_MUL }; + const int out_nodes[] = { i, i + 2 }; +- if (cc == GGML_CUDA_CC_DGX_SPARK && rms_norm->op == GGML_OP_RMS_NORM && ++ if (GGML_CUDA_CC_IS_NVIDIA(cc) && rms_norm->op == GGML_OP_RMS_NORM && + mul->op == GGML_OP_MUL && (mul->src[0] == rms_norm || mul->src[1] == rms_norm) && + rms_norm->src[0] == node && node->src[0] && node->src[1] && + node->type == GGML_TYPE_F32 && node->src[0]->type == GGML_TYPE_F32 && +@@ -3375,10 +3393,12 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph + ggml_is_contiguous(node->src[1]) && ggml_is_contiguous(node) && + ggml_is_contiguous(rms_norm) && ggml_is_contiguous(mul) && + ggml_can_fuse_subgraph(cgraph, i, 3, ops, out_nodes, 2) && +- ggml_cuda_check_fusion_memory_ranges(cgraph, i, 3, out_nodes, 2)) { ++ ggml_cuda_check_fusion_memory_ranges(cgraph, i, 3, out_nodes, 2, false, true)) { + const ggml_tensor * weight = mul->src[0] == rms_norm ? mul->src[1] : mul->src[0]; + if (weight && weight->type == GGML_TYPE_F32 && ggml_is_contiguous(weight) && +- weight->ne[0] == node->ne[0] && ggml_nrows(weight) == 1) { ++ weight->ne[0] == node->ne[0] && ggml_nrows(weight) == 1 && ++ weight->data != node->data && weight->data != rms_norm->data && ++ weight->data != mul->data) { + ggml_cuda_op_add_rms_norm_fused(*cuda_ctx, node, rms_norm, mul); + return 2; + } +@@ -3475,6 +3495,49 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph + } + } + ++ // SWIGLU [+ reshapes] + Hadamard sign flip + reshape + FWHT-hint matmul (the FFN down projection's input, and the ++ // DeltaNet gated norm's output before ssm_out): the transform computes silu(a) * b as it loads, so the SWIGLU output ++ // never goes through memory ++ static const bool swiglu_fwht = getenv("GGML_CUDA_NO_SWIGLU_FWHT") == nullptr; ++ if (swiglu_fwht && node->op == GGML_OP_GLU && ggml_get_glu_op(node) == GGML_GLU_OP_SWIGLU) { ++ int k = i + 1; ++ while (k < cgraph->n_nodes && k < i + 4 && cgraph->nodes[k]->op == GGML_OP_RESHAPE) { ++ ++k; ++ } ++ const int n_ops = k - i + 3; ++ if (k + 2 < cgraph->n_nodes && n_ops <= 6) { ++ ggml_op ops[6]; ++ for (int j = 0; j < n_ops; ++j) { ++ ops[j] = cgraph->nodes[i + j]->op; ++ } ++ const int out_nodes[] = { k + 2 }; ++ const ggml_tensor * mul = cgraph->nodes[k]; ++ const ggml_tensor * reshape = cgraph->nodes[k + 1]; ++ ggml_tensor * mm = cgraph->nodes[k + 2]; ++ ++ bool chain_ok = true; ++ for (int j = i + 1; j <= k; ++j) { ++ chain_ok = chain_ok && cgraph->nodes[j]->src[0] == cgraph->nodes[j - 1]; ++ } ++ const ggml_tensor * signs = mul->src[1]; ++ ++ const bool pattern_ok = chain_ok && mul->op == GGML_OP_MUL && reshape->op == GGML_OP_RESHAPE && ++ mm->op == GGML_OP_MUL_MAT && node->src[1] != nullptr && ++ ggml_get_op_params_i32(node, 1) == 0 && node->type == GGML_TYPE_F32 && mul->type == GGML_TYPE_F32 && ++ ggml_get_op_params_i32(mm, 1) == GGML_HINT_SRC0_IS_HADAMARD && ++ mm->src[1] == reshape && reshape->src[0] == mul && signs && ++ signs->ne[1] == 1 && signs->ne[2] == 1 && signs->ne[3] == 1 && ++ ggml_are_same_shape(node->src[0], node) && ggml_are_same_shape(node->src[1], node) && ++ ggml_nelements(mul) == ggml_nelements(node) && ggml_is_contiguous(node) && ++ signs->ne[0] == mul->ne[0] && signs->ne[0] % mm->src[0]->ne[0] == 0 && ++ ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1); ++ ++ if (pattern_ok && ggml_cuda_op_fwht_swiglu_signed(*cuda_ctx, node->src[0], node->src[1], signs, mm)) { ++ return n_ops - 1; ++ } ++ } ++ } ++ + // Hadamard sign flip + reshape + FWHT-hint matmul: multiply the sign + // vector during the transform's load instead of a separate full pass + if (ggml_can_fuse_subgraph(cgraph, i, { GGML_OP_MUL, GGML_OP_RESHAPE, GGML_OP_MUL_MAT }, { i + 2 })) { diff --git a/pkgs-set/pkgs/llama-cpp-bonsai.nix b/pkgs-set/pkgs/llama-cpp-bonsai.nix index 167aaf1..ed6fb4a 100644 --- a/pkgs-set/pkgs/llama-cpp-bonsai.nix +++ b/pkgs-set/pkgs/llama-cpp-bonsai.nix @@ -59,6 +59,18 @@ cudaPackages.backendStdenv.mkDerivation (finalAttrs: { # slots hold at shutdown, go to files that a later prefix match reads back (an 8K-token conversation # with its DeltaNet checkpoints is ~1 GB and restores in ~0.6 s instead of a ~14 s prefill). ./llama-cpp-bonsai-prompt-cache-disk.patch + # Prefill fusions, bit-identical: ADD+RMS_NORM+MUL on every NVIDIA GPU (PrismML-Eng/llama.cpp#209 by + # @cklxx, from professorpalmer's 0040), and the SWIGLU feeding a Hadamard-rotated matmul computed inside + # the transform's load. ~+2% prefill. + ./llama-cpp-bonsai-prefill-fusions.patch + # PTQ1_0 MMQ: activations double-buffered with cp.async and the next weight tile fetched into registers + # during compute. Bit-identical, +3.3% on the matmuls. + ./llama-cpp-bonsai-mmq-pipeline.patch + # PTQ1_0 MMQ quantizes activations with one scale per 128 values and accumulates a whole weight block in + # int32, so the float epilogue runs once per block instead of per 32: +15-18% on the matmuls, ~+11% + # prefill. Not bit-identical: PPL +0.12%, top token unchanged at 99.42% of positions. Decode and + # speculative verify (mat-vec) are untouched. + ./llama-cpp-bonsai-mmq-y128.patch ]; nativeBuildInputs = [