wmantly/llama.cpp-xrip
1
0 Fork 0
Code Issues Pull Requests Actions 32 Packages Projects Releases Wiki Activity
Files
master
llama.cpp-xrip/sm75-prefill-patch/sm75-prefill-code-only-af084752.patch
T
xrip 60c62ce9f2 Merge upstream master; take upstream DFlash2, MMQ fusion, and top-k 
Assisted-by: GLM-5.3-Flash
2026-11-01 22:29:01 +03:00

375 lines
15 KiB
Diff
Raw Permalink Blame History

diff --git a/ggml/src/ggml-cuda/CMakeLists.txt b/ggml/src/ggml-cuda/CMakeLists.txt
--- a/ggml/src/ggml-cuda/CMakeLists.txt
+++ b/ggml/src/ggml-cuda/CMakeLists.txt
@@ -1,6 +1,8 @@
cmake_minimum_required(VERSION 3.18) # for CMAKE_CUDA_ARCHITECTURES
+
+option(GGML_CUDA_SM75_Q8_0_FAST_HALF "SM75: experimental packed Q8_0 to FP16 conversion in FlashAttention" OFF)
find_package(CUDAToolkit)
if (CUDAToolkit_FOUND)
message(STATUS "CUDA Toolkit found")
@@ -118,10 +120,14 @@
ggml_add_backend_library(ggml-cuda
${GGML_HEADERS_CUDA}
${GGML_SOURCES_CUDA}
)
+ if (GGML_CUDA_SM75_Q8_0_FAST_HALF)
+ target_compile_definitions(ggml-cuda PRIVATE GGML_CUDA_SM75_Q8_0_FAST_HALF)
+ endif()
+
if (GGML_CUDA_GRAPHS)
add_compile_definitions(GGML_CUDA_USE_GRAPHS)
endif()
if (GGML_CUDA_FORCE_MMQ)
diff --git a/ggml/src/ggml-cuda/fattn-common.cuh b/ggml/src/ggml-cuda/fattn-common.cuh
--- a/ggml/src/ggml-cuda/fattn-common.cuh
+++ b/ggml/src/ggml-cuda/fattn-common.cuh
@@ -1,10 +1,13 @@
#pragma once
#include "common.cuh"
#include "convert.cuh"
#include "vecdotq.cuh"
+#ifdef GGML_CUDA_SM75_Q8_0_FAST_HALF
+#include "sm75-q8.cuh"
+#endif
#include <cstdint>
#define FATTN_KQ_STRIDE 256
#define HALF_MAX_HALF __float2half(65504.0f/2) // Use neg. of this instead of -INFINITY to initialize KQ max vals to avoid NaN upon subtraction.
@@ -599,11 +602,15 @@
if constexpr (std::is_same<T, half>::value) {
const half2 d = __half2half2(x[ib].d);
#pragma unroll
for (int l0 = 0; l0 < ne; l0 += 2) {
+#if defined(GGML_CUDA_SM75_Q8_0_FAST_HALF) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ == 750 && !defined(GGML_USE_HIP)
+ ((half2 *) dst)[l0/2] = d * ggml_cuda_sm75_q8_pair_to_half2(qs[l0 + 0], qs[l0 + 1]);
+#else
((half2 *) dst)[l0/2] = d * make_half2(qs[l0 + 0], qs[l0 + 1]);
+#endif
}
} else
#endif // FP16_AVAILABLE
if constexpr (std::is_same<T, float>::value) {
const float d = x[ib].d;
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -29,10 +29,11 @@
#include "ggml-cuda/fwht.cuh"
#include "ggml-cuda/getrows.cuh"
#include "ggml-cuda/im2col.cuh"
#include "ggml-cuda/mmf.cuh"
#include "ggml-cuda/mmq.cuh"
+#include "ggml-cuda/sm75-tuning.cuh"
#include "ggml-cuda/mmvf.cuh"
#include "ggml-cuda/mmvq.cuh"
#include "ggml-cuda/moe-weighted-reduction.cuh"
#include "ggml-cuda/norm.cuh"
#include "ggml-cuda/opt-step-adamw.cuh"
@@ -1929,20 +1930,29 @@
if (ggml_cuda_should_use_mmvq(src0->type, cc, ne11)) {
ggml_cuda_mul_mat_vec_q(ctx, src0, src1, nullptr, dst);
return;
}
#ifndef GGML_CUDA_FORCE_MMQ
- // Bound FP16 weight workspace and keep small batches on MMQ.
- const bool sm75_ffn = (ne00 == 4096 && ne01 == 12288) || (ne00 == 12288 && ne01 == 4096)
- || (ne00 == 5120 && ne01 == 17408) || (ne00 == 17408 && ne01 == 5120);
- if (cc == GGML_CUDA_CC_TURING && sm75_ffn && ne11 >= 1024 && ne11 <= 2048
- && (src0->type == GGML_TYPE_Q4_K || src0->type == GGML_TYPE_IQ4_XS)
+ // Preserve the existing FFN override in auto mode. Experimental modes are
+ // restricted to dense prefill; MMVQ/decode and MUL_MAT_ID are untouched.
+ const bool sm75_legacy_type = src0->type == GGML_TYPE_Q4_K || src0->type == GGML_TYPE_IQ4_XS;
+ const bool sm75_test_type = sm75_legacy_type || src0->type == GGML_TYPE_Q5_K || src0->type == GGML_TYPE_Q6_K;
+ if (cc == GGML_CUDA_CC_TURING && sm75_test_type && ne11 >= 512
&& ne02 == 1 && ne03 == 1 && ne12 == 1 && ne13 == 1
&& ggml_is_contiguous(src0) && ggml_is_contiguous(src1)
&& !ggml_backend_buffer_is_host(src0->buffer)) {
- ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
- return;
+ const auto & tuning = ggml_cuda_sm75::get_options();
+ const bool use_cublas = ggml_cuda_sm75::use_cublas(tuning, sm75_legacy_type, sm75_test_type, ne00, ne01, ne11);
+ if (tuning.trace) {
+ GGML_LOG_INFO("SM75 prefill: device=%d type=%s M=%lld K=%lld N=%lld mode=%s override=%s\n",
+ ctx.device, ggml_type_name(src0->type), (long long) ne01, (long long) ne00, (long long) ne11,
+ ggml_cuda_sm75::mode_name(tuning.mode), use_cublas ? "cublas" : "none");
+ }
+ if (use_cublas) {
+ ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
+ return;
+ }
}
#endif
if (ggml_cuda_should_use_mmq(src0->type, cc, ne11, /*n_experts =*/ 0)) {
ggml_cuda_mul_mat_q(ctx, src0, src1, nullptr, dst);
return;
diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
--- a/ggml/src/ggml-cuda/mmq.cuh
+++ b/ggml/src/ggml-cuda/mmq.cuh
@@ -1,8 +1,9 @@
#pragma once
#include "common.cuh"
+#include "sm75-tuning.cuh"
#include <climits>
#include <cstdint>
#define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available.
@@ -1478,29 +1479,46 @@
void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) {
const int id = ggml_cuda_get_device();
const int cc = ggml_cuda_info().devices[id].cc;
const size_t smpbo = ggml_cuda_info().devices[id].smpbo;
- int J_best = 0;
- int ntiles_J_best = INT_MAX;
+ const bool sm75_dense = cc == GGML_CUDA_CC_TURING && !fallback &&
+ (type == GGML_TYPE_Q4_K || type == GGML_TYPE_IQ4_XS) &&
+ !args.ids_dst && !args.expert_bounds && args.nchannels_x == 1 && args.nchannels_y == 1 &&
+ args.nsamples_x == 1 && args.nsamples_y == 1 && args.ncols_y >= 512 && args.ncols_y % 128 == 0;
+ const auto * tuning = sm75_dense ? &ggml_cuda_sm75::get_options() : nullptr;
+ const int J_limit = tuning ? ggml_cuda_sm75::mmq_j_limit(*tuning, true, args.ncols_y) : 128;
- for (int J = 8; J <= 128 && ntiles_J_best > 1; J += 8) {
- const ggml_cuda_mmq_config config = ggml_cuda_mmq_get_config(type, J, fallback, cc);
- if (config.type == GGML_TYPE_COUNT) {
- continue;
+ const auto select_J = [&](const int limit) {
+ int best = 0;
+ int ntiles_best = INT_MAX;
+ for (int J = 8; J <= limit && ntiles_best > 1; J += 8) {
+ const ggml_cuda_mmq_config config = ggml_cuda_mmq_get_config(type, J, fallback, cc);
+ if (config.type == GGML_TYPE_COUNT || mmq_get_nbytes_shared(config, cc) > smpbo) {
+ continue;
+ }
+ const int ntiles_x = (args.ncols_opt + config.J - 1) / config.J;
+ if (ntiles_x < ntiles_best) {
+ best = J;
+ ntiles_best = ntiles_x;
+ }
}
+ return best;
+ };
- if (mmq_get_nbytes_shared(config, cc) > smpbo) {
- continue;
- }
-
- const int ntiles_x = (args.ncols_opt + config.J - 1) / config.J;
-
- if (ntiles_x < ntiles_J_best) {
- J_best = J;
- ntiles_J_best = ntiles_x;
- }
+ int J_best = select_J(J_limit);
+ if (J_best == 0 && J_limit < 128) {
+ // A cap is a tuning hint, not permission to launch an invalid kernel.
+ J_best = select_J(128);
+ }
+ // Keep ggml_cuda_mmq_get_J_max and all padding allocations unchanged: a
+ // narrower launch must not shrink buffers needed by another dispatch path.
+ if (tuning && tuning->trace && J_best != 0) {
+ GGML_LOG_INFO("SM75 MMQ: device=%d type=%s M=%lld K=%lld N=%lld J=%d cap=%d shared=%zu\n",
+ id, ggml_type_name(type), (long long) args.nrows_x, (long long) args.ncols_x,
+ (long long) args.ncols_y, J_best, J_limit,
+ mmq_get_nbytes_shared(ggml_cuda_mmq_get_config(type, J_best, fallback, cc), cc));
}
switch (J_best) {
case 8:
launch_mul_mat_q<type, 8, fallback>(ctx, args, stream);
diff --git a/ggml/src/ggml-cuda/sm75-q8.cuh b/ggml/src/ggml-cuda/sm75-q8.cuh
new file mode 100644
index 0000000..7ce29f0
--- /dev/null
+++ b/ggml/src/ggml-cuda/sm75-q8.cuh
@@ -0,0 +1,30 @@
+#pragma once
+
+#include <cstdint>
+
+#if defined(__CUDACC__) && !defined(GGML_USE_HIP)
+#include <cuda_fp16.h>
+#define GGML_SM75_Q8_HD __host__ __device__
+#else
+#define GGML_SM75_Q8_HD
+#endif
+
+// An unsigned byte u in the mantissa of 0x6400 encodes 1024 + u exactly.
+// Flipping the INT8 sign bit gives u = q + 128. Subtracting 1152 then recovers
+// every q in [-128, 127], including -128. No float conversion or wide load.
+GGML_SM75_Q8_HD static constexpr uint32_t ggml_cuda_sm75_q8_pair_bits(uint8_t a, uint8_t b) {
+ return 0x64006400u | uint32_t(a ^ 0x80u) | (uint32_t(b ^ 0x80u) << 16);
+}
+
+#undef GGML_SM75_Q8_HD
+
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == 750 && !defined(GGML_USE_HIP)
+static __device__ __forceinline__ half2 ggml_cuda_sm75_q8_pair_to_half2(int8_t a, int8_t b) {
+ const uint32_t bits = ggml_cuda_sm75_q8_pair_bits(uint8_t(a), uint8_t(b));
+ const half2 encoded = __halves2half2(__ushort_as_half(uint16_t(bits)),
+ __ushort_as_half(uint16_t(bits >> 16)));
+ // The explicit rounding intrinsic prevents a following scale multiply from
+ // contracting with the subtraction. Keep the same multiply as the old path.
+ return __hsub2_rn(encoded, __float2half2_rn(1152.0f));
+}
+#endif
diff --git a/ggml/src/ggml-cuda/sm75-tuning.cuh b/ggml/src/ggml-cuda/sm75-tuning.cuh
new file mode 100644
index 0000000..33ee38e
--- /dev/null
+++ b/ggml/src/ggml-cuda/sm75-tuning.cuh
@@ -0,0 +1,144 @@
+#pragma once
+
+// Opt-in host policy. Keep environment variables fixed while the process runs.
+// CUDA Graphs can retain the selected kernels.
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+namespace ggml_cuda_sm75 {
+
+enum class prefill_mode { automatic, mmq, cublas };
+
+struct options {
+ prefill_mode mode = prefill_mode::automatic;
+ int64_t min_batch = 1024;
+ uint64_t max_scratch_mib = 512;
+ int mmq_max_j = 0;
+ bool trace = false;
+};
+
+// Reject overflow, signs, whitespace, and trailing text.
+inline bool parse_uint(const char * text, uint64_t lo, uint64_t hi, uint64_t & value) {
+ if (!text || !*text || lo > hi) {
+ return false;
+ }
+ uint64_t n = 0;
+ for (const char * p = text; *p; ++p) {
+ if (*p < '0' || *p > '9') {
+ return false;
+ }
+ const uint64_t digit = uint64_t(*p - '0');
+ if (digit > hi || n > (hi - digit) / 10) {
+ return false;
+ }
+ n = 10*n + digit;
+ }
+ if (n < lo) {
+ return false;
+ }
+ value = n;
+ return true;
+}
+
+inline void warn_env(const char * name, const char * value) {
+ std::fprintf(stderr, "ggml_cuda_sm75: invalid %s='%s'; using its default\n", name, value ? value : "(unset)");
+}
+
+inline uint64_t env_uint(const char * name, uint64_t fallback, uint64_t lo, uint64_t hi) {
+ const char * text = std::getenv(name);
+ if (!text) {
+ return fallback;
+ }
+ uint64_t value;
+ if (!parse_uint(text, lo, hi, value)) {
+ warn_env(name, text);
+ return fallback;
+ }
+ return value;
+}
+
+inline options read_options() {
+ options o;
+ const char * mode = std::getenv("GGML_CUDA_SM75_PREFILL_MODE");
+ if (mode) {
+ if (std::strcmp(mode, "mmq") == 0) {
+ o.mode = prefill_mode::mmq;
+ } else if (std::strcmp(mode, "cublas") == 0) {
+ o.mode = prefill_mode::cublas;
+ } else if (std::strcmp(mode, "auto") != 0) {
+ warn_env("GGML_CUDA_SM75_PREFILL_MODE", mode);
+ }
+ }
+ o.min_batch = int64_t(env_uint("GGML_CUDA_SM75_PREFILL_MIN_BATCH", 1024, 512, 1048576));
+ o.max_scratch_mib = env_uint("GGML_CUDA_SM75_PREFILL_MAX_SCRATCH_MIB", 512, 1, 65536);
+ o.mmq_max_j = int(env_uint("GGML_CUDA_SM75_MMQ_MAX_J", 0, 0, 128));
+ if (o.mmq_max_j != 0 && o.mmq_max_j != 32 && o.mmq_max_j != 64 &&
+ o.mmq_max_j != 96 && o.mmq_max_j != 128) {
+ warn_env("GGML_CUDA_SM75_MMQ_MAX_J", std::getenv("GGML_CUDA_SM75_MMQ_MAX_J"));
+ o.mmq_max_j = 0;
+ }
+ o.trace = env_uint("GGML_CUDA_SM75_PREFILL_TRACE", 0, 0, 1) != 0;
+ return o;
+}
+
+inline const options & get_options() {
+ static const options o = read_options();
+ return o;
+}
+
+inline const char * mode_name(prefill_mode mode) {
+ switch (mode) {
+ case prefill_mode::automatic: return "auto";
+ case prefill_mode::mmq: return "mmq";
+ case prefill_mode::cublas: return "cublas";
+ }
+ return "auto";
+}
+
+inline bool legacy_ffn_shape(int64_t k, int64_t m) {
+ return (k == 4096 && m == 12288) || (k == 12288 && m == 4096) ||
+ (k == 5120 && m == 17408) || (k == 17408 && m == 5120);
+}
+
+// Bound conversion scratch by 4*(M*K + K*N + M*N), also covering FP32 overrides.
+// Excludes cuBLAS workspace and cached pool memory. Check products before multiplication.
+inline bool scratch_fits(int64_t k, int64_t m, int64_t n, uint64_t limit_mib) {
+ if (k <= 0 || m <= 0 || n <= 0 || limit_mib == 0 || limit_mib > 65536) {
+ return false;
+ }
+ uint64_t budget = limit_mib * (1024u * 1024u / 4u);
+ const uint64_t a[] = { uint64_t(m), uint64_t(k), uint64_t(m) };
+ const uint64_t b[] = { uint64_t(k), uint64_t(n), uint64_t(n) };
+ for (int i = 0; i < 3; ++i) {
+ if (a[i] > budget / b[i]) {
+ return false;
+ }
+ budget -= a[i]*b[i];
+ }
+ return true;
+}
+
+// Caller checks SM75, dense shapes, F32 I/O, contiguous inputs, and device weights.
+// legacy_type: Q4_K/IQ4_XS; experiment_type: also Q5_K/Q6_K.
+inline bool use_cublas(const options & o, bool legacy_type, bool experiment_type,
+ int64_t k, int64_t m, int64_t n) {
+ if (o.mode == prefill_mode::automatic) {
+ // Exactly the pre-existing af084752 policy, including its tail behavior.
+ return legacy_type && legacy_ffn_shape(k, m) && n >= 1024 && n <= 2048;
+ }
+ if (o.mode == prefill_mode::mmq) {
+ return false;
+ }
+ return experiment_type && n >= o.min_batch && n >= 512 && n % 128 == 0 &&
+ k >= 256 && m >= 128 && k % 256 == 0 && m % 128 == 0 &&
+ scratch_fits(k, m, n, o.max_scratch_mib);
+}
+
+// Select existing specializations only; do not change layouts or buffer sizes.
+inline int mmq_j_limit(const options & o, bool eligible, int64_t n) {
+ return eligible && n >= 512 && n % 128 == 0 && o.mmq_max_j != 0 ? o.mmq_max_j : 128;
+}
+
+} // namespace ggml_cuda_sm75
Reference in new issue View Git Blame Copy Permalink