Compare commits

..
Author SHA1 Message Date
Xuan Son Nguyen 431e72841b server: add --no-chat-input-special 2026-11-10 12:18:01 +01:00
366 changed files with 6985 additions and 12167 deletions
+12 -12
View File
@@ -1,4 +1,4 @@
ARG ONEAPI_VERSION=2026.2.0-devel-ubuntu24.04
ARG ONEAPI_VERSION=2026.1.1-devel-ubuntu24.04
ARG BUILD_DATE=N/A
ARG APP_VERSION=N/A
ARG APP_REVISION=N/A
@@ -22,20 +22,20 @@ RUN LLAMA_BUILD_NUMBER="$APP_VERSION" npm run build
FROM docker.io/intel/deep-learning-essentials:$ONEAPI_VERSION AS build
RUN echo "Download oneDNN installation package" && \
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/46452a4a-7b7f-49f2-be4c-602a900180bb/intel-onednn-2026.1.0.625_offline.sh -O intel-onednn_offline.sh && \
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/94c3dbac-0852-45be-a57d-21c204cada3e/intel-onednn-2026.0.2.46_offline.sh -O intel-onednn_offline.sh && \
echo "Install oneDNN components" && \
bash intel-onednn_offline.sh -s -a --silent --eula accept
ARG GGML_SYCL_F16=ON
ARG LEVEL_ZERO_VERSION=1.34.0
ARG LEVEL_ZERO_VERSION=1.28.2
ARG LEVEL_ZERO_UBUNTU_VERSION=u24.04
RUN apt-get update && \
apt-get install -y git libssl-dev wget ca-certificates && \
cd /tmp && \
wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/libze1_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O libze1.deb && \
wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/libze-dev_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O libze-dev.deb && \
apt-get -o Dpkg::Options::="--force-overwrite" install -y ./libze1.deb ./libze-dev.deb && \
rm -f /tmp/libze1.deb /tmp/libze-dev.deb
wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/level-zero_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O level-zero.deb && \
wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/level-zero-devel_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O level-zero-devel.deb && \
apt-get -o Dpkg::Options::="--force-overwrite" install -y ./level-zero.deb ./level-zero-devel.deb && \
rm -f /tmp/level-zero.deb /tmp/level-zero-devel.deb
WORKDIR /app
@@ -91,11 +91,11 @@ COPY --from=build /opt/intel/oneapi/dnnl /opt/intel/oneapi/dnnl
#ARG IGDGMM_VERSION=22.8.2
ARG IGC_VERSION=v2.42.6
ARG IGC_VERSION_FULL=2_2.42.6+23070
ARG COMPUTE_RUNTIME_VERSION=26.40.40196.8
ARG COMPUTE_RUNTIME_VERSION_FULL=26.40.40196.8-0
ARG IGDGMM_VERSION=22.10.2
ARG IGC_VERSION=v2.34.4
ARG IGC_VERSION_FULL=2_2.34.4+21428
ARG COMPUTE_RUNTIME_VERSION=26.18.38308.1
ARG COMPUTE_RUNTIME_VERSION_FULL=26.18.38308.1-0
ARG IGDGMM_VERSION=22.10.0
RUN mkdir /tmp/neo/ && cd /tmp/neo/ \
&& wget https://github.com/intel/intel-graphics-compiler/releases/download/$IGC_VERSION/intel-igc-core-${IGC_VERSION_FULL}_amd64.deb \
&& wget https://github.com/intel/intel-graphics-compiler/releases/download/$IGC_VERSION/intel-igc-opencl-${IGC_VERSION_FULL}_amd64.deb \
+1 -1
View File
@@ -37,7 +37,7 @@ env:
LLAMA_ARG_LOG_PREFIX: 1
LLAMA_ARG_LOG_TIMESTAMPS: 1
# TODO: fix failing tests on OpenVINO backend
CTEST_EXCLUDE: "test-llama-archs|test-llama-context|test-fp8-kv"
CTEST_EXCLUDE: "test-llama-archs|test-llama-context"
jobs:
ubuntu-24-openvino:
+8 -8
View File
@@ -52,8 +52,8 @@ jobs:
env:
ONEAPI_ROOT: /opt/intel/oneapi/
ONEAPI_INSTALLER_VERSION: "2026.2"
LEVEL_ZERO_VERSION: "1.34.0"
ONEAPI_INSTALLER_VERSION: "2026.1"
LEVEL_ZERO_VERSION: "1.33.1"
LEVEL_ZERO_UBUNTU_VERSION: "u24.04"
continue-on-error: true
@@ -68,12 +68,12 @@ jobs:
run: |
cd /tmp
echo "Download oneAPI Deep Learning Essentials installation package"
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/aef232d2-aac5-4cd6-bceb-911019bbb6a8/intel-deep-learning-essentials-2026.2.0.640_offline.sh -O intel-deep-learning-essentials_offline.sh
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/b5c100a7-8052-4e88-90e2-63d9916a998b/intel-deep-learning-essentials-2026.1.4.17_offline.sh -O intel-deep-learning-essentials_offline.sh
echo "Install oneAPI Deep Learning Essentials components"
sudo bash intel-deep-learning-essentials_offline.sh -s -a --silent --eula accept
echo "Download oneDNN installation package"
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/46452a4a-7b7f-49f2-be4c-602a900180bb/intel-onednn-2026.1.0.625_offline.sh -O intel-onednn_offline.sh
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/94c3dbac-0852-45be-a57d-21c204cada3e/intel-onednn-2026.0.2.46_offline.sh -O intel-onednn_offline.sh
echo "Install oneDNN components"
sudo bash intel-onednn_offline.sh -s -a --silent --eula accept
@@ -150,13 +150,13 @@ jobs:
shell: bash
env:
WINDOWS_DLE_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/b177a006-5cd9-4cb2-807f-8199cc73d53c/intel-deep-learning-essentials-2026.2.0.638_offline.exe
WINDOWS_ONEDNN_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/8f905666-c85c-43fb-9bd0-2618c7ecd33b/intel-onednn-2026.1.0.622_offline.exe
WINDOWS_DLE_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/a51104eb-6f4a-4a0e-9dae-1b81b0c81ae1/intel-deep-learning-essentials-2026.1.4.16_offline.exe
WINDOWS_ONEDNN_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/34e28083-d04b-46a3-826c-8c1d00bf1e1c/intel-onednn-2026.0.2.47_offline.exe
WINDOWS_DPCPP_MKL: intel.oneapi.win.cpp-dpcpp-common:intel.oneapi.win.mkl.devel:intel.oneapi.win.tbb.devel
WINDOWS_DNNL: intel.oneapi.win.dnnl
LEVEL_ZERO_SDK_URL: https://github.com/oneapi-src/level-zero/releases/download/v1.34.0/level-zero-win-sdk-1.34.0.zip
LEVEL_ZERO_SDK_URL: https://github.com/oneapi-src/level-zero/releases/download/v1.33.1/level-zero-win-sdk-1.33.1.zip
ONEAPI_ROOT: "C:/Program Files (x86)/Intel/oneAPI"
ONEAPI_INSTALLER_VERSION: "2026.2"
ONEAPI_INSTALLER_VERSION: "2026.1"
steps:
- name: Clone
id: checkout
+1 -1
View File
@@ -35,7 +35,7 @@ jobs:
uses: actions/setup-python@v6
with:
python-version: "3.11"
pip-install: -r requirements/requirements-all.txt ty==0.0.100
pip-install: -r requirements/requirements-all.txt ty==0.0.88
# - name: Type-check with Pyright
# uses: jakebailey/pyright-action@v2
# with:
+10 -9
View File
@@ -1307,14 +1307,14 @@ jobs:
shell: bash
env:
WINDOWS_DLE_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/b177a006-5cd9-4cb2-807f-8199cc73d53c/intel-deep-learning-essentials-2026.2.0.638_offline.exe
WINDOWS_ONEDNN_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/8f905666-c85c-43fb-9bd0-2618c7ecd33b/intel-onednn-2026.1.0.622_offline.exe
WINDOWS_DLE_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/a51104eb-6f4a-4a0e-9dae-1b81b0c81ae1/intel-deep-learning-essentials-2026.1.4.16_offline.exe
WINDOWS_ONEDNN_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/34e28083-d04b-46a3-826c-8c1d00bf1e1c/intel-onednn-2026.0.2.47_offline.exe
WINDOWS_DPCPP_MKL: intel.oneapi.win.cpp-dpcpp-common:intel.oneapi.win.mkl.devel:intel.oneapi.win.tbb.devel
WINDOWS_DNNL: intel.oneapi.win.dnnl
LEVEL_ZERO_SDK_URL: https://github.com/oneapi-src/level-zero/releases/download/v1.34.0/level-zero-win-sdk-1.34.0.zip
LEVEL_ZERO_SDK_URL: https://github.com/oneapi-src/level-zero/releases/download/v1.33.1/level-zero-win-sdk-1.33.1.zip
ONEAPI_ROOT: "C:/Program Files (x86)/Intel/oneAPI"
ONEAPI_INSTALLER_VERSION: "2026.2"
ONEAPI_INSTALLER_VERSION: "2026.1"
steps:
- name: Clone
@@ -1366,7 +1366,8 @@ jobs:
cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_adapter_level_zero.dll" ./build/bin
cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_adapter_level_zero_v2.dll" ./build/bin
cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_loader_shared.dll" ./build/bin
cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_adapter_opencl.dll" ./build/bin
cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_loader.dll" ./build/bin
cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_win_proxy_loader.dll" ./build/bin
ZE_LOADER_DLL=$(find "${{ env.ONEAPI_ROOT }}" "$LEVEL_ZERO_V1_SDK_PATH" -iname ze_loader.dll -print -quit 2>/dev/null || true)
if [ -n "$ZE_LOADER_DLL" ]; then
@@ -1420,8 +1421,8 @@ jobs:
env:
ONEAPI_ROOT: /opt/intel/oneapi/
ONEAPI_INSTALLER_VERSION: "2026.2"
LEVEL_ZERO_VERSION: "1.34.0"
ONEAPI_INSTALLER_VERSION: "2026.1"
LEVEL_ZERO_VERSION: "1.33.1"
LEVEL_ZERO_UBUNTU_VERSION: "u24.04"
steps:
@@ -1436,12 +1437,12 @@ jobs:
run: |
cd /tmp
echo "Download oneAPI Deep Learning Essentials installation package"
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/aef232d2-aac5-4cd6-bceb-911019bbb6a8/intel-deep-learning-essentials-2026.2.0.640_offline.sh -O intel-deep-learning-essentials_offline.sh
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/b5c100a7-8052-4e88-90e2-63d9916a998b/intel-deep-learning-essentials-2026.1.4.17_offline.sh -O intel-deep-learning-essentials_offline.sh
echo "Install oneAPI Deep Learning Essentials components"
sudo bash intel-deep-learning-essentials_offline.sh -s -a --silent --eula accept
echo "Download oneDNN installation package"
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/46452a4a-7b7f-49f2-be4c-602a900180bb/intel-onednn-2026.1.0.625_offline.sh -O intel-onednn_offline.sh
wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/94c3dbac-0852-45be-a57d-21c204cada3e/intel-onednn-2026.0.2.46_offline.sh -O intel-onednn_offline.sh
echo "Install oneDNN components"
sudo bash intel-onednn_offline.sh -s -a --silent --eula accept
+1 -1
View File
@@ -1,6 +1,6 @@
# llama.cpp
![llama.cpp](https://raw.githubusercontent.com/ggml-org/llama.brand/refs/heads/master/logo/svg/cover/llama-cpp/cover-llama-cpp-dark.svg)
![llama](https://raw.githubusercontent.com/ggml-org/llama.brand/refs/heads/master/cover/llama-cpp/cover-llama-cpp-dark.svg)
<div align="center">
+1 -1
View File
@@ -180,7 +180,7 @@ if [ ! -z ${GG_BUILD_OPENVINO} ]; then
CMAKE_EXTRA="${CMAKE_EXTRA} -DGGML_OPENVINO=ON"
# TODO: fix failing tests on OpenVINO backend
CTEST_EXTRA="-E test-llama-archs|test-llama-context|test-fp8-kv"
CTEST_EXTRA="-E test-llama-archs|test-llama-context"
fi
## helpers
+11 -3
View File
@@ -282,7 +282,6 @@ const std::vector<ggml_type> kv_cache_types = {
GGML_TYPE_F32,
GGML_TYPE_F16,
GGML_TYPE_BF16,
GGML_TYPE_F8_E4M3,
GGML_TYPE_Q8_0,
GGML_TYPE_Q4_0,
GGML_TYPE_Q4_1,
@@ -3829,6 +3828,15 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.prefill_assistant = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_PREFILL_ASSISTANT"));
add_opt(common_arg(
{"--chat-input-special"},
{"--no-chat-input-special"},
string_format("experimental: whether to parse special tokens from user input messages in chat completions, does not affect the raw /completions endpoint (default: %s)\n"
"note: this is not a security feature, do not rely on it to prevent prompt injection", params.chat_input_special ? "enabled" : "disabled"),
[](common_params & params, bool value) {
params.chat_input_special = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_CHAT_INPUT_SPECIAL"));
add_opt(common_arg(
{"-sps", "--slot-prompt-similarity"}, "SIMILARITY",
string_format("how much the prompt of a request must match the prompt of a slot in order to use that slot (default: %.2f, 0.0 = disabled)\n", params.slot_prompt_similarity),
@@ -4492,7 +4500,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params, const std::string & value) {
params.tts_lang = value;
}
).set_examples({LLAMA_EXAMPLE_TTS, LLAMA_EXAMPLE_SERVER}));
).set_examples({LLAMA_EXAMPLE_TTS}));
add_opt(common_arg(
{"--tts-voice"}, "VOICE",
"voice for audio generation (model-dependent)\n"
@@ -4500,7 +4508,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params, const std::string & value) {
params.tts_voice = value;
}
).set_examples({LLAMA_EXAMPLE_TTS, LLAMA_EXAMPLE_SERVER}));
).set_examples({LLAMA_EXAMPLE_TTS}));
add_opt(common_arg(
{"--tts-speaker-file"}, "FNAME",
"speaker file path for audio generation",
-6
View File
@@ -131,12 +131,6 @@ static std::string normalize_quotes_to_json(const std::string & input) {
continue;
}
// Models can emit raw control characters inside string literals; escape them to keep the JSON valid.
if ((in_single_quoted || in_double_quoted) && static_cast<unsigned char>(c) < 0x20) {
result += escape_json_string_inner(std::string(1, c));
continue;
}
// Handle quote characters
if (c == '"') {
if (in_single_quoted) {
+6 -35
View File
@@ -819,29 +819,21 @@ static bool is_lfm3_template(const std::string & src) {
return src.find("<additional_tools>") != std::string::npos;
}
static std::string string_lower(std::string s) {
std::transform(s.begin(), s.end(), s.begin(), [](unsigned char c) { return (char) std::tolower(c); });
return s;
}
// LFM3-Audio: language code -> name
static std::string lfm3_language_name(const std::string & language) {
// LFM3-Audio: a known language code maps to its name, no language lets the model detect it
static std::string lfm3_asr_system_prompt(const std::string & language) {
static const std::map<std::string, std::string> names = {
{ "ar", "Arabic" }, { "da", "Danish" }, { "de", "German" }, { "el", "Greek" }, { "en", "English" },
{ "es", "Spanish" }, { "fi", "Finnish" }, { "fr", "French" }, { "he", "Hebrew" }, { "it", "Italian" },
{ "ja", "Japanese" }, { "ko", "Korean" }, { "nl", "Dutch" }, { "pl", "Polish" }, { "pt", "Portuguese" },
{ "sv", "Swedish" }, { "tr", "Turkish" }, { "zh", "Chinese" },
};
const auto it = names.find(string_lower(language));
return it != names.end() ? it->second : language;
}
// LFM3-Audio: no language = auto detect
static std::string lfm3_asr_system_prompt(const std::string & language) {
if (language.empty()) {
return "Perform ASR.";
}
return "Perform ASR in " + lfm3_language_name(language) + ".";
std::string code = language;
std::transform(code.begin(), code.end(), code.begin(), [](unsigned char c) { return (char) std::tolower(c); });
const auto it = names.find(code);
return "Perform ASR in " + (it != names.end() ? it->second : language) + ".";
}
common_chat_prompt_preset common_chat_get_asr_prompt(const common_chat_templates * chat_templates, const std::string & language) {
@@ -876,27 +868,6 @@ std::string common_chat_get_tts_prompt(const common_chat_templates * chat_templa
});
}
std::string common_chat_get_speech_system_prompt(const common_chat_templates * chat_templates,
const std::string & language, const std::string & voice) {
if (!chat_templates || !chat_templates->template_default || !is_lfm3_template(chat_templates->template_default->source())) {
return "";
}
// same voices as the "tts" template
const std::string v = string_lower(voice);
const bool has_voice = v == "female" || v == "male";
if (language.empty() && !has_voice) {
return "";
}
std::string out = "Speak";
if (!language.empty()) {
out += " in " + lfm3_language_name(language);
}
if (has_voice) {
out += " using a " + v + " voice";
}
return out + ".";
}
std::string common_chat_templates_source(const struct common_chat_templates * tmpls, const std::string & variant) {
if (!variant.empty()) {
if (variant == "tool_use") {
-4
View File
@@ -500,7 +500,3 @@ common_chat_prompt_preset common_chat_get_asr_prompt(const common_chat_templates
// language and voice are optional, the template defines their defaults
std::string common_chat_get_tts_prompt(const common_chat_templates * chat_templates, const std::string & text,
const std::string & language = "", const std::string & voice = "");
// system prompt line for the speech language and voice (e.g. LFM3-Audio), empty if not used
std::string common_chat_get_speech_system_prompt(const common_chat_templates * chat_templates,
const std::string & language, const std::string & voice);
+1
View File
@@ -679,6 +679,7 @@ struct common_params {
common_reasoning_format reasoning_format = COMMON_REASONING_FORMAT_DEEPSEEK;
int enable_reasoning = -1; // -1 = auto, 0 = disable, 1 = enable
bool prefill_assistant = true; // if true, any trailing assistant message will be prefilled into the response
bool chat_input_special = true; // if false, special tokens in user input are tokenized as plain text
int sleep_idle_seconds = -1; // if >0, server will sleep after this many seconds of idle time
std::vector<std::string> api_keys;
-37
View File
@@ -860,7 +860,6 @@ public:
std::unordered_set<std::string> required;
std::vector<std::pair<std::string, const common_chat_schema *>> properties;
std::map<std::string, size_t> enum_values;
std::vector<const common_chat_schema *> leaves;
std::function<void(const common_chat_schema &, bool, int)> add_component = [&](const common_chat_schema & comp, bool is_required, int depth) {
if (depth >= COMMON_RECURSIVE_LIMIT) {
throw std::invalid_argument("JSON schema is nested more than " + std::to_string(COMMON_RECURSIVE_LIMIT) + " levels deep");
@@ -880,10 +879,6 @@ public:
for (const auto & v : as<common_chat_schema_enum>(comp).values) {
enum_values[_generate_constant_rule(v)] += 1;
}
} else if (is_required && comp.kind() != common_chat_schema::KIND_ANY) {
// Direct non-structural children are used for merging. anyOf alternatives are not merged.
// Empty schemas are skipped.
leaves.push_back(&comp);
}
};
for (const auto & child : schema.children) {
@@ -906,38 +901,6 @@ public:
return _add_rule(rule_name, "(" + string_join(enum_intersection, " | ") + ")");
}
}
if (properties.empty() && !leaves.empty()) {
if (leaves.size() == 1) {
return visit(*leaves[0], name);
}
bool all_numeric = true;
bool any_integer = false;
int64_t lo = std::numeric_limits<int64_t>::min();
int64_t hi = std::numeric_limits<int64_t>::max();
for (const auto * leaf : leaves) {
if (leaf->kind() == common_chat_schema::KIND_INTEGER) {
const auto & i = as<common_chat_schema_integer>(*leaf);
any_integer = true;
lo = std::max(lo, i.minimum);
hi = std::min(hi, i.maximum);
} else if (leaf->kind() != common_chat_schema::KIND_NUMBER) {
all_numeric = false;
}
}
if (all_numeric) {
if (lo > hi) {
throw std::invalid_argument("allOf has an empty numeric range");
}
if (any_integer) {
common_chat_schema_integer merged;
merged.minimum = lo;
merged.maximum = hi;
return visit(merged, name);
}
return visit(*leaves[0], name);
}
}
return _add_rule(rule_name, _build_object_rule(properties, required, name, nullptr));
}
+1 -1
View File
@@ -300,7 +300,7 @@ common_chat_params common_chat_params_init_deepseek_v3_2(const common_chat_templ
});
if (include_grammar) {
data.grammar_lazy = !has_response_format && has_tools && !require_tools;
data.grammar_lazy = has_tools && !require_tools;
data.grammar = build_grammar([&](const common_grammar_builder & builder) {
data.parser.build_grammar(builder, data.grammar_lazy);
});
+1 -1
View File
@@ -109,7 +109,7 @@ common_chat_params common_chat_params_init_ministral_3(const common_chat_templat
});
if (include_grammar) {
data.grammar_lazy = !has_response_format && has_tools && inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_AUTO;
data.grammar_lazy = has_tools && inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_AUTO;
data.grammar = build_grammar([&](const common_grammar_builder & builder) {
data.parser.build_grammar(builder, data.grammar_lazy);
+4 -14
View File
@@ -46,10 +46,9 @@ common_chat_params common_chat_params_init_muse_spark(const common_chat_template
auto extract_reasoning = inputs.reasoning_format != COMMON_REASONING_FORMAT_NONE;
auto has_tools = inputs.tools.is_array() && !inputs.tools.empty();
auto has_response_format = !inputs.json_schema.is_null() && inputs.json_schema.is_object();
// Constrained grammar whenever tools are offered or a response format is requested.
auto include_grammar = has_response_format || (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE);
auto has_tools = inputs.tools.is_array() && !inputs.tools.empty();
// Constrained grammar whenever tools are offered.
auto include_grammar = has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE;
data.parser = build_chat_peg_parser([&](common_chat_peg_builder & p) {
auto start = p.rule("start", p.literal("<|start|>assistant"));
@@ -76,15 +75,6 @@ common_chat_params common_chat_params_init_muse_spark(const common_chat_template
auto final_msg = p.rule("final", recipient + p.literal("<|message|>") +
p.content(p.until_one_of({ "<|eot|>", "<|eom|>" })));
if (has_response_format) {
auto response_json = p.content(p.schema(p.json(), "response-format-schema", inputs.json_schema));
auto response_format = p.rule("response-format",
recipient + p.literal("<|message|>") +
((p.literal("```json") + p.space() + response_json + p.space() + p.literal("```")) | response_json));
return p.zero_or_more(start + analysis) + start + response_format;
}
if (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE) {
auto string_value = p.ac(
p.tool_arg_string_value(p.until("</atem:parameter>")) + p.tool_arg_close(p.literal("</atem:parameter>")),
@@ -142,7 +132,7 @@ common_chat_params common_chat_params_init_muse_spark(const common_chat_template
});
if (include_grammar) {
data.grammar_lazy = !(has_response_format || (has_tools && inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_REQUIRED));
data.grammar_lazy = inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_REQUIRED;
data.grammar = build_grammar([&](const common_grammar_builder & builder) {
data.parser.build_grammar(builder, data.grammar_lazy);
});
+1 -1
View File
@@ -191,7 +191,7 @@ common_chat_params common_chat_params_init_qwen3_coder(const common_chat_templat
});
if (include_grammar) {
data.grammar_lazy = !has_response_format && has_tools && inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_AUTO;
data.grammar_lazy = has_tools && inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_AUTO;
data.grammar = build_grammar([&](const common_grammar_builder & builder) {
data.parser.build_grammar(builder, data.grammar_lazy);
-1
View File
@@ -371,7 +371,6 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
"InternVisionModel": "internvl",
"JanusForConditionalGeneration": "januspro",
"KimiK25ForConditionalGeneration": "kimivl",
"KimiK31ForConditionalGeneration": "kimivl",
"KimiVLForConditionalGeneration": "kimivl",
"Lfm2AudioForConditionalGeneration": "lfm2",
"D1OmniModel": "lfm2",
-13
View File
@@ -1,6 +1,5 @@
from __future__ import annotations
import re
from typing import Callable, Iterable, TYPE_CHECKING
import torch
@@ -199,18 +198,6 @@ class SarvamMoEModel(BailingMoeV2Model):
# - full rotary (no partial_rotary_factor)
# - expert bias is zero-mean normalized at load time
def set_vocab(self):
# BPE with byte fallback: mark <0xXX> tokens as BYTE so they decode to raw bytes
tokens, toktypes, tokpre = self.get_vocab_base()
toktypes = [gguf.TokenType.BYTE if re.fullmatch(r"<0x[0-9A-F]{2}>", tok) else typ for tok, typ in zip(tokens, toktypes)]
self.gguf_writer.add_tokenizer_model("gpt2")
self.gguf_writer.add_tokenizer_pre(tokpre)
self.gguf_writer.add_token_list(tokens)
self.gguf_writer.add_token_types(toktypes)
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
special_vocab.add_to_gguf(self.gguf_writer)
def set_gguf_parameters(self):
super().set_gguf_parameters()
hparams = self.hparams
+32 -71
View File
@@ -7,7 +7,6 @@ import ast
import logging
import contextlib
import json
import math
import os
import re
import sys
@@ -32,10 +31,10 @@ import gguf
from gguf.vocab import MistralTokenizerType, MistralVocab
try:
from mistral_common.tokens.tokenizers.base import TokenizerVersion # ty: ignore[import-not-found]
from mistral_common.tokens.tokenizers.multimodal import DATASET_MEAN as _MISTRAL_COMMON_DATASET_MEAN, DATASET_STD as _MISTRAL_COMMON_DATASET_STD # ty: ignore[import-not-found]
from mistral_common.tokens.tokenizers.tekken import Tekkenizer # ty: ignore[import-not-found]
from mistral_common.tokens.tokenizers.sentencepiece import ( # ty: ignore[import-not-found]
from mistral_common.tokens.tokenizers.base import TokenizerVersion # type: ignore[import-not-found, ty:unresolved-import]
from mistral_common.tokens.tokenizers.multimodal import DATASET_MEAN as _MISTRAL_COMMON_DATASET_MEAN, DATASET_STD as _MISTRAL_COMMON_DATASET_STD # type: ignore[import-not-found, ty:unresolved-import]
from mistral_common.tokens.tokenizers.tekken import Tekkenizer # type: ignore[import-not-found, ty:unresolved-import]
from mistral_common.tokens.tokenizers.sentencepiece import ( # type: ignore[import-not-found, ty:unresolved-import]
SentencePieceTokenizer,
)
@@ -273,7 +272,7 @@ class ModelBase:
if weight_map is None or not isinstance(weight_map, dict):
raise ValueError(f"Can't load 'weight_map' from {index_name!r}")
tensor_names_from_index.update(weight_map.keys())
part_dict: dict[str, None] = dict.fromkeys(weight_map.values(), None)
part_dict: dict[str, None] = dict.fromkeys(weight_map.values(), None) # ty: ignore[invalid-assignment]
part_names = sorted(part_dict.keys())
else:
weight_map = {}
@@ -344,36 +343,6 @@ class ModelBase:
logger.info(f" + {scale_name} (per-expert scale, shape [{len(scales)}])")
self.gguf_writer.add_tensor(scale_name, scale_vals)
def _prepare_kv_cache_scales(self):
for name in list(self.model_tensors.keys()):
if not name.endswith((".k_scale", ".v_scale")):
continue
mapped = self.tensor_map.get_type_and_name(key=name, try_suffixes=(".k_scale", ".v_scale"))
if mapped is None:
continue
tensor_type, new_name = mapped
if tensor_type not in (gguf.MODEL_TENSOR.ATTN_K, gguf.MODEL_TENSOR.ATTN_V, gguf.MODEL_TENSOR.ATTN_QKV):
continue
bid = int(new_name.split(".")[1])
scale_type = gguf.MODEL_TENSOR.ATTN_K if name.endswith(".k_scale") else gguf.MODEL_TENSOR.ATTN_V
suffix = ".k_scale" if name.endswith(".k_scale") else ".v_scale"
new_name = gguf.TENSOR_NAMES[scale_type].format(bid=bid) + suffix
scale = LazyTorchTensor.to_eager(self.model_tensors.pop(name)())
if scale.dtype not in (torch.bfloat16, torch.float32) or scale.numel() != 1:
raise ValueError(f"KV cache scale {name!r} must be a scalar BF16 or FP32 tensor")
scale = scale.float()
value = float(scale.item())
if not math.isfinite(value) or value == 0.0:
raise ValueError(f"KV cache scale {name!r} must be finite and nonzero")
logger.info(f" + {new_name} (KV cache scale, shape [1])")
self.gguf_writer.add_tensor(new_name, scale.flatten().numpy())
def dequant_model(self):
# If all quantized tensors were already handled (e.g. pure NVFP4), skip
if self._is_nvfp4 and not any(k.endswith((".weight_scale", ".weight_scale_inv", ".input_scale", ".activation_scale", "_activation_scale", ".k_scale", ".v_scale")) for k in self.model_tensors):
@@ -1486,8 +1455,6 @@ class ModelBase:
unmatched = layer_algo.keys() - matched
assert not unmatched, f"quantized_layers entries not in the model tensors: {sorted(unmatched)}"
self._prepare_kv_cache_scales()
# NVFP4 weights are repacked and written directly to gguf_writer.
# This must run before dequant_model so NVFP4 tensors are removed
# from model_tensors, leaving only non-NVFP4 (e.g. FP8) for dequant.
@@ -2055,16 +2022,15 @@ class TextModel(ModelBase):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
vocab_size = self.hparams.get("vocab_size", len(tokenizer.vocab))
assert max(tokenizer.vocab.values()) < vocab_size
vocab_size = self.hparams.get("vocab_size", len(tokenizer.vocab)) # ty: ignore[unresolved-attribute]
assert max(tokenizer.vocab.values()) < vocab_size # ty: ignore[unresolved-attribute]
tokpre = self.get_vocab_base_pre(tokenizer)
reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()}
added_vocab = tokenizer.get_added_vocab()
reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()} # ty: ignore[unresolved-attribute]
added_vocab = tokenizer.get_added_vocab() # ty: ignore[unresolved-attribute]
added_tokens_decoder = tokenizer.added_tokens_decoder
added_tokens_decoder = tokenizer.added_tokens_decoder # ty: ignore[unresolved-attribute]
for i in range(vocab_size):
if i not in reverse_vocab:
@@ -2077,7 +2043,7 @@ class TextModel(ModelBase):
# To avoid unexpected issues - we make sure to normalize non-normalized tokens
if not added_tokens_decoder[i].normalized:
previous_token = token
token = tokenizer.decode(tokenizer.encode(token, add_special_tokens=False)) # ty: ignore[invalid-assignment]
token = tokenizer.decode(tokenizer.encode(token, add_special_tokens=False)) # ty: ignore[unresolved-attribute, invalid-assignment]
if previous_token != token:
logger.info(f"{repr(previous_token)} is encoded and decoded back to {repr(token)} using AutoTokenizer")
@@ -2487,7 +2453,6 @@ class TextModel(ModelBase):
def _set_vocab_hybriddna(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
vocab = tokenizer.get_vocab()
vocab_size = self.hparams.get("vocab_size", len(vocab))
assert max(vocab.values()) < vocab_size
@@ -2501,11 +2466,11 @@ class TextModel(ModelBase):
if tokenizer.k != 6 or tail not in ("pad_a", "pad_n"):
raise NotImplementedError(f"HybridDNATokenizer with k={tokenizer.k}, tail={tail!r}")
for kmer in [*tokenizer.kmers, *getattr(tokenizer, "npad_tokens", [])]:
reverse_vocab[tokenizer.dna_token_to_id[kmer]] = kmer + "\ue000"
reverse_vocab[tokenizer.dna_token_to_id[kmer]] = kmer + "\ue000" # ty: ignore[unresolved-attribute]
# older HybridDNATokenizer versions (Carbon-500M/3B/8B) do not register the base tokenizer's added tokens (<|im_start|>, <think>, ...) as their own
# but they are still encoded as special tokens: take their types from the base tokenizer
base_tokenizer = getattr(tokenizer, "_base_tokenizer", tokenizer)
added_tokens_decoder = {**base_tokenizer.added_tokens_decoder, **tokenizer.added_tokens_decoder}
added_tokens_decoder = {**base_tokenizer.added_tokens_decoder, **tokenizer.added_tokens_decoder} # ty: ignore[unresolved-attribute]
tokens: list[str] = []
toktypes: list[int] = []
@@ -2552,25 +2517,24 @@ class TextModel(ModelBase):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(dir_model, trust_remote_code=True)
assert tokenizer is not None
vocab_size = hparams["vocab_size"]
assert max(tokenizer.get_vocab().values()) < vocab_size
assert max(tokenizer.get_vocab().values()) < vocab_size # ty: ignore[unresolved-attribute]
tokpre = self.get_vocab_base_pre(tokenizer)
merges = []
vocab = {}
mergeable_ranks = tokenizer.mergeable_ranks
mergeable_ranks = tokenizer.mergeable_ranks # ty: ignore[unresolved-attribute]
for token, rank in mergeable_ranks.items():
vocab[QwenModel.token_bytes_to_string(token)] = rank
if len(token) == 1:
continue
merged = QwenModel.bpe(mergeable_ranks, token, max_rank=rank)
assert len(merged) == 2
merges.append(str(' '.join(map(QwenModel.token_bytes_to_string, merged))))
merges.append(' '.join(map(QwenModel.token_bytes_to_string, merged)))
# for this kind of tokenizer, added_vocab is not a subset of vocab, so they need to be combined
added_vocab = tokenizer.special_tokens
added_vocab = tokenizer.special_tokens # ty: ignore[unresolved-attribute]
reverse_vocab = {id_ : encoded_tok for encoded_tok, id_ in {**vocab, **added_vocab}.items()}
for i in range(vocab_size):
@@ -2593,10 +2557,10 @@ class TextModel(ModelBase):
special_vocab.merges = merges
# only add special tokens when they were not already loaded from config.json
if len(special_vocab.special_token_ids) == 0:
special_vocab._set_special_token("bos", tokenizer.special_tokens["<|endoftext|>"])
special_vocab._set_special_token("eos", tokenizer.special_tokens["<|endoftext|>"])
special_vocab._set_special_token("bos", tokenizer.special_tokens["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("eos", tokenizer.special_tokens["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
# this one is usually not in config.json anyway
special_vocab._set_special_token("unk", tokenizer.special_tokens["<|endoftext|>"])
special_vocab._set_special_token("unk", tokenizer.special_tokens["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab.add_to_gguf(self.gguf_writer)
def _set_vocab_sentencepiece(self, add_to_gguf=True):
@@ -2870,23 +2834,21 @@ class TextModel(ModelBase):
def _set_vocab_glmedge(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
tokens, toktypes, tokpre = self.get_vocab_base()
self.gguf_writer.add_tokenizer_model("gpt2")
self.gguf_writer.add_tokenizer_pre(tokpre)
self.gguf_writer.add_token_list(tokens)
self.gguf_writer.add_token_types(toktypes)
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"])
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab.add_to_gguf(self.gguf_writer)
def _set_vocab_glm(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
tokens, toktypes, tokpre = self.get_vocab_base()
self.gguf_writer.add_tokenizer_model("gpt2")
@@ -2895,10 +2857,10 @@ class TextModel(ModelBase):
self.gguf_writer.add_token_types(toktypes)
# Special tokens
# Note: Using <|endoftext|> (151329) for eot causes endless generation
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["[gMASK]"]) # 151331
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"]) # 151336
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"]) # 151329
special_vocab._set_special_token("eom", tokenizer.get_added_vocab()["<|observation|>"]) # 151338
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["[gMASK]"]) # ty: ignore[unresolved-attribute] # 151331
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"]) # ty: ignore[unresolved-attribute] # 151336
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute] # 151329
special_vocab._set_special_token("eom", tokenizer.get_added_vocab()["<|observation|>"]) # ty: ignore[unresolved-attribute] # 151338
special_vocab.add_to_gguf(self.gguf_writer)
def _set_vocab_interns1(self):
@@ -2907,17 +2869,16 @@ class TextModel(ModelBase):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
vocab = getattr(tokenizer, 'vocab', tokenizer.get_vocab())
vocab = getattr(tokenizer, 'vocab', tokenizer.get_vocab()) # ty: ignore[unresolved-attribute]
vocab_size = self.hparams.get("vocab_size", len(vocab))
assert max(vocab.values()) < vocab_size
tokpre = self.get_vocab_base_pre(tokenizer)
reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in vocab.items()}
added_vocab = tokenizer.get_added_vocab()
added_vocab = tokenizer.get_added_vocab() # ty: ignore[unresolved-attribute]
added_tokens_decoder = tokenizer.added_tokens_decoder
added_tokens_decoder = tokenizer.added_tokens_decoder # ty: ignore[unresolved-attribute]
for i in range(vocab_size):
if i not in reverse_vocab:
@@ -2930,7 +2891,7 @@ class TextModel(ModelBase):
# To avoid unexpected issues - we make sure to normalize non-normalized tokens
if not added_tokens_decoder[i].normalized:
previous_token = token
token = tokenizer.decode(tokenizer.encode(token, add_special_tokens=False)) # ty: ignore[invalid-assignment]
token = tokenizer.decode(tokenizer.encode(token, add_special_tokens=False)) # ty: ignore[unresolved-attribute, invalid-assignment]
if previous_token != token:
logger.info(f"{repr(previous_token)} is encoded and decoded back to {repr(token)} using AutoTokenizer")
+1 -2
View File
@@ -60,8 +60,7 @@ class BeamModel(TextModel):
# otherwise EOT is guessed from a hash-ordered scan
from transformers import AutoTokenizer
tok = AutoTokenizer.from_pretrained(self.dir_model)
assert tok is not None
ret_id = tok.convert_tokens_to_ids("<|return|>")
ret_id = tok.convert_tokens_to_ids("<|return|>") # ty: ignore[unresolved-attribute]
if isinstance(ret_id, int) and ret_id >= 0:
self.gguf_writer.add_eot_token_id(ret_id)
+10 -11
View File
@@ -132,7 +132,6 @@ class BertModel(TextModel):
from base64 import b64decode
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
with open(tokenizer_path, "r", encoding="utf-8") as fp:
tokenizer_json = json.load(fp)
@@ -141,13 +140,13 @@ class BertModel(TextModel):
with open(tokenizer_config_path, "r", encoding="utf-8") as fp:
tokenizer_config_json = json.load(fp)
add_prefix = tokenizer.add_prefix_space
remove_whitespaces = tokenizer.clean_up_tokenization_spaces
add_prefix = tokenizer.add_prefix_space # ty: ignore[unresolved-attribute]
remove_whitespaces = tokenizer.clean_up_tokenization_spaces # ty: ignore[unresolved-attribute]
precompiled_charsmap = b64decode(tokenizer_json["normalizer"]["precompiled_charsmap"])
vocab_size = max(self.hparams.get("vocab_size", 0), tokenizer.vocab_size)
vocab_size = max(self.hparams.get("vocab_size", 0), tokenizer.vocab_size) # ty: ignore[unresolved-attribute]
else:
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
sentencepiece_model.ParseFromString(open(tokenizer_path, "rb").read())
assert sentencepiece_model.trainer_spec.model_type == 1 # UNIGRAM
@@ -184,20 +183,20 @@ class BertModel(TextModel):
scores[token_id] = score
toktypes[token_id] = toktype
else:
added_vocab = tokenizer.get_added_vocab()
added_vocab = tokenizer.get_added_vocab() # ty: ignore[unresolved-attribute]
unk_token = tokenizer_config_json.get("unk_token")
unk_token_id = added_vocab.get(unk_token, tokenizer_json["model"].get("unk_id", 3))
unk_token_id = added_vocab.get(unk_token, tokenizer_json["model"].get("unk_id", 3)) # ty: ignore[no-matching-overload]
for token_id in range(tokenizer.vocab_size):
piece = tokenizer._convert_id_to_token(token_id)
if (piece := tokenizer._convert_id_to_token(token_id)) is not None:
for token_id in range(tokenizer.vocab_size): # ty: ignore[unresolved-attribute]
piece = tokenizer._convert_id_to_token(token_id) # ty: ignore[unresolved-attribute]
if (piece := tokenizer._convert_id_to_token(token_id)) is not None: # ty: ignore[unresolved-attribute]
text = piece.encode("utf-8")
score = tokenizer_json["model"]["vocab"][token_id][1]
toktype = SentencePieceTokenTypes.NORMAL
if token_id == unk_token_id:
toktype = SentencePieceTokenTypes.UNKNOWN
elif token_id in tokenizer.all_special_ids:
elif token_id in tokenizer.all_special_ids: # ty: ignore[unresolved-attribute]
toktype = SentencePieceTokenTypes.CONTROL
elif token_id in added_vocab.values():
toktype = SentencePieceTokenTypes.USER_DEFINED
+16 -19
View File
@@ -22,13 +22,12 @@ class ChatGLMModel(TextModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(dir_model, trust_remote_code=True)
assert tokenizer is not None
vocab_size = hparams.get("padded_vocab_size", len(tokenizer.get_vocab()))
assert max(tokenizer.get_vocab().values()) < vocab_size
vocab_size = hparams.get("padded_vocab_size", len(tokenizer.get_vocab())) # ty: ignore[unresolved-attribute]
assert max(tokenizer.get_vocab().values()) < vocab_size # ty: ignore[unresolved-attribute]
role_special_tokens = ["<|system|>", "<|user|>", "<|assistant|>", "<|observation|>"]
special_tokens = ["[MASK]", "[gMASK]", "[sMASK]", "sop", "eop"] + role_special_tokens
for token_id in range(vocab_size):
piece = tokenizer._convert_id_to_token(token_id)
piece = tokenizer._convert_id_to_token(token_id) # ty: ignore[unresolved-attribute]
if token_id == 0:
piece = "<unk>"
elif token_id == 1:
@@ -36,18 +35,17 @@ class ChatGLMModel(TextModel):
elif token_id == 2:
piece = "<eos>"
assert isinstance(piece, str)
text = piece.encode("utf-8")
text = piece.encode("utf-8") # ty: ignore[unresolved-attribute]
score = 0.0
# Referencing the tokenizer Python implementation(https://huggingface.co/THUDM/chatglm3-6b/blob/main/tokenization_chatglm.py),
# it is only valid if it is less than tokenizer.tokenizer.sp_model.vocab_size()
if len(piece) != 0 and token_id < tokenizer.tokenizer.sp_model.vocab_size():
score = tokenizer.tokenizer.sp_model.get_score(token_id)
if len(piece) != 0 and token_id < tokenizer.tokenizer.sp_model.vocab_size(): # ty: ignore[unresolved-attribute, invalid-argument-type]
score = tokenizer.tokenizer.sp_model.get_score(token_id) # ty: ignore[unresolved-attribute]
if token_id >= tokenizer.tokenizer.sp_model.vocab_size():
if token_id >= tokenizer.tokenizer.sp_model.vocab_size(): # ty: ignore[unresolved-attribute]
if piece in special_tokens:
toktype = SentencePieceTokenTypes.CONTROL
elif len(piece) == 0:
elif len(piece) == 0: # ty: ignore[invalid-argument-type]
text = f"[PAD{token_id}]".encode("utf-8")
toktype = SentencePieceTokenTypes.UNUSED
else:
@@ -58,13 +56,13 @@ class ChatGLMModel(TextModel):
continue
toktype = SentencePieceTokenTypes.NORMAL
if tokenizer.tokenizer.sp_model.is_unknown(token_id):
if tokenizer.tokenizer.sp_model.is_unknown(token_id): # ty: ignore[unresolved-attribute]
toktype = SentencePieceTokenTypes.UNKNOWN
elif tokenizer.tokenizer.sp_model.is_control(token_id):
elif tokenizer.tokenizer.sp_model.is_control(token_id): # ty: ignore[unresolved-attribute]
toktype = SentencePieceTokenTypes.CONTROL
elif tokenizer.tokenizer.sp_model.is_unused(token_id):
elif tokenizer.tokenizer.sp_model.is_unused(token_id): # ty: ignore[unresolved-attribute]
toktype = SentencePieceTokenTypes.UNUSED
elif tokenizer.tokenizer.sp_model.is_byte(token_id):
elif tokenizer.tokenizer.sp_model.is_byte(token_id): # ty: ignore[unresolved-attribute]
toktype = SentencePieceTokenTypes.BYTE
tokens.append(text)
@@ -117,9 +115,8 @@ class ChatGLMModel(TextModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(dir_model, trust_remote_code=True)
assert tokenizer is not None
vocab_size = hparams.get("padded_vocab_size",hparams["vocab_size"])
assert max(tokenizer.get_vocab().values()) < vocab_size
assert max(tokenizer.get_vocab().values()) < vocab_size # ty: ignore[unresolved-attribute]
tokens, toktypes, tokpre = self.get_vocab_base()
self.gguf_writer.add_tokenizer_model("gpt2")
@@ -128,10 +125,10 @@ class ChatGLMModel(TextModel):
self.gguf_writer.add_token_types(toktypes)
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
# only add special tokens when they were not already loaded from config.json
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"])
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"]) # ty: ignore[unresolved-attribute]
# this one is usually not in config.json anyway
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab.add_to_gguf(self.gguf_writer)
def set_gguf_parameters(self):
+2 -3
View File
@@ -277,14 +277,13 @@ class DeepseekV2Model(TextModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
tokpre = self.get_vocab_base_pre(tokenizer)
if tokpre == "kimi-k2":
# Build merges list using the approach similar to HunYuanMoE
merges = []
vocab = {}
mergeable_ranks = tokenizer.model._mergeable_ranks # ty: ignore[missing-union-attribute]
mergeable_ranks = tokenizer.model._mergeable_ranks # ty: ignore[unresolved-attribute]
for token, rank in mergeable_ranks.items():
vocab[QwenModel.token_bytes_to_string(token)] = rank
if len(token) == 1:
@@ -295,7 +294,7 @@ class DeepseekV2Model(TextModel):
# Build token list
vocab_size = self.hparams["vocab_size"]
special_tokens = tokenizer.special_tokens
special_tokens = tokenizer.special_tokens # ty: ignore[unresolved-attribute]
reverse_vocab = {id_ : encoded_tok for encoded_tok, id_ in {**vocab, **special_tokens}.items()}
tokens: list[str] = []
toktypes: list[int] = []
-2
View File
@@ -62,7 +62,6 @@ def engram_compressed_token_map(tokenizer_dir: Path) -> tuple[list[int], int]:
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(tokenizer_dir)
assert tokenizer is not None
sentinel = "\ue000"
normalizer = normalizers.Sequence([
normalizers.NFKC(),
@@ -81,7 +80,6 @@ def engram_compressed_token_map(tokenizer_dir: Path) -> tuple[list[int], int]:
text = backend.decode([token_id], skip_special_tokens=False)
if "\ufffd" in text:
key = backend.id_to_token(token_id)
assert key is not None
else:
normalized = normalizer.normalize_str(text)
key = normalized if normalized else text
+1 -2
View File
@@ -75,14 +75,13 @@ class Dots3NoteModel(DeepseekV2Model):
def set_vocab(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
tokens, toktypes, tokpre = self.get_vocab_base()
self.gguf_writer.add_tokenizer_model("gpt2")
self.gguf_writer.add_tokenizer_pre(tokpre)
self.gguf_writer.add_token_list(tokens)
self.gguf_writer.add_token_types(toktypes)
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|endofassistant|>"])
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|endofassistant|>"]) # ty: ignore[unresolved-attribute]
special_vocab.add_to_gguf(self.gguf_writer)
@classmethod
+2 -3
View File
@@ -19,16 +19,15 @@ class DreamModel(TextModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
vocab_dict = tokenizer.get_vocab()
vocab_dict = tokenizer.get_vocab() # ty: ignore[unresolved-attribute]
vocab_size = self.hparams.get("vocab_size", len(vocab_dict))
assert max(vocab_dict.values()) < vocab_size
tokpre = self.get_vocab_base_pre(tokenizer)
reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in vocab_dict.items()}
added_vocab = tokenizer.get_added_vocab()
added_vocab = tokenizer.get_added_vocab() # ty: ignore[unresolved-attribute]
for i in range(vocab_size):
if i not in reverse_vocab:
+8 -10
View File
@@ -31,7 +31,6 @@ class Glm4Model(TextModel):
def set_vocab(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
tokens, toktypes, tokpre = self.get_vocab_base()
self.gguf_writer.add_tokenizer_model("gpt2")
@@ -39,10 +38,10 @@ class Glm4Model(TextModel):
self.gguf_writer.add_token_list(tokens)
self.gguf_writer.add_token_types(toktypes)
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"])
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|user|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab.add_to_gguf(self.gguf_writer)
def set_gguf_parameters(self):
@@ -392,17 +391,16 @@ class SolarOpenModel(Glm4MoeModel):
def set_vocab(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True)
tokens, toktypes, tokpre = self.get_vocab_base()
self.gguf_writer.add_tokenizer_model("gpt2")
self.gguf_writer.add_tokenizer_pre(tokpre)
self.gguf_writer.add_token_list(tokens)
self.gguf_writer.add_token_types(toktypes)
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|endoftext|>"])
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<unk>"])
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["<|startoftext|>"])
special_vocab._set_special_token("eos", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|endoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("unk", tokenizer.get_added_vocab()["<unk>"]) # ty: ignore[unresolved-attribute]
special_vocab._set_special_token("bos", tokenizer.get_added_vocab()["<|startoftext|>"]) # ty: ignore[unresolved-attribute]
special_vocab.add_to_gguf(self.gguf_writer)
-10
View File
@@ -85,16 +85,6 @@ class GptOssModel(TextModel):
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
name, gen = item
# Normalize erroneous GPT-OSS MXFP4 tensor naming.
# https://github.com/p-e-w/heretic/pull/211#issuecomment-5094370390
if "mlp.experts." in name:
for proj in ("down_proj", "gate_up_proj"):
name = name.replace(
f"{proj}_blocks_blocks", f"{proj}_blocks"
).replace(
f"{proj}_blocks_scales", f"{proj}_scales"
)
if "sinks" in name:
name += ".weight"
+6 -8
View File
@@ -24,7 +24,6 @@ class HunYuanMoEModel(TextModel):
def set_vocab(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
# 1. Get the pre-tokenizer identifier hash
tokpre = self.get_vocab_base_pre(tokenizer)
@@ -32,7 +31,7 @@ class HunYuanMoEModel(TextModel):
# 2. Reverse-engineer the merges list from mergeable_ranks
merges = []
vocab = {}
mergeable_ranks = tokenizer.mergeable_ranks
mergeable_ranks = tokenizer.mergeable_ranks # ty: ignore[unresolved-attribute]
for token, rank in mergeable_ranks.items():
vocab[QwenModel.token_bytes_to_string(token)] = rank
if len(token) == 1:
@@ -43,8 +42,8 @@ class HunYuanMoEModel(TextModel):
# 3. Generate the tokens and toktypes lists
vocab_size = self.hparams["vocab_size"]
assert tokenizer.vocab_size == vocab_size
special_tokens = tokenizer.special_tokens
assert tokenizer.vocab_size == vocab_size # ty: ignore[unresolved-attribute]
special_tokens = tokenizer.special_tokens # ty: ignore[unresolved-attribute]
reverse_vocab = {id_ : encoded_tok for encoded_tok, id_ in {**vocab, **special_tokens}.items()}
tokens: list[str] = []
toktypes: list[int] = []
@@ -185,7 +184,6 @@ class HunYuanModel(TextModel):
else:
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
# 1. Get the pre-tokenizer identifier hash
tokpre = self.get_vocab_base_pre(tokenizer)
@@ -193,7 +191,7 @@ class HunYuanModel(TextModel):
# 2. Reverse-engineer the merges list from mergeable_ranks
merges = []
vocab = {}
mergeable_ranks = tokenizer.mergeable_ranks
mergeable_ranks = tokenizer.mergeable_ranks # ty: ignore[unresolved-attribute]
for token, rank in mergeable_ranks.items():
vocab[QwenModel.token_bytes_to_string(token)] = rank
if len(token) == 1:
@@ -204,8 +202,8 @@ class HunYuanModel(TextModel):
# 3. Generate the tokens and toktypes lists
vocab_size = self.hparams["vocab_size"]
assert tokenizer.vocab_size == vocab_size
special_tokens = tokenizer.special_tokens
assert tokenizer.vocab_size == vocab_size # ty: ignore[unresolved-attribute]
special_tokens = tokenizer.special_tokens # ty: ignore[unresolved-attribute]
reverse_vocab = {id_ : encoded_tok for encoded_tok, id_ in {**vocab, **special_tokens}.items()}
tokens: list[str] = []
toktypes: list[int] = []
+1 -1
View File
@@ -36,7 +36,7 @@ class InternLM2Model(TextModel):
logger.error(f'Error: Missing {tokenizer_path}')
sys.exit(1)
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
sentencepiece_model.ParseFromString(open(tokenizer_path, "rb").read())
add_prefix = sentencepiece_model.normalizer_spec.add_dummy_prefix
+3 -4
View File
@@ -29,14 +29,13 @@ class KimiLinearModel(TextModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
tokpre = self.get_vocab_base_pre(tokenizer)
if tokpre == "kimi-k2":
# Build merges list using the approach similar to HunYuanMoE
merges = []
vocab = {}
mergeable_ranks = tokenizer.model._mergeable_ranks # ty: ignore[missing-union-attribute]
mergeable_ranks = tokenizer.model._mergeable_ranks # ty: ignore[unresolved-attribute]
for token, rank in mergeable_ranks.items():
vocab[QwenModel.token_bytes_to_string(token)] = rank
if len(token) == 1:
@@ -46,7 +45,7 @@ class KimiLinearModel(TextModel):
merges.append(' '.join(map(QwenModel.token_bytes_to_string, merged)))
# Build token list
vocab_size = self.hparams["vocab_size"]
special_tokens = tokenizer.special_tokens
special_tokens = tokenizer.special_tokens # ty: ignore[unresolved-attribute]
reverse_vocab = {id_ : encoded_tok for encoded_tok, id_ in {**vocab, **special_tokens}.items()}
tokens: list[str] = []
toktypes: list[int] = []
@@ -72,7 +71,7 @@ class KimiLinearModel(TextModel):
special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=False)
special_vocab.add_to_gguf(self.gguf_writer)
# override eos id in config.json with tiktoken eos id
self.gguf_writer.add_eos_token_id(tokenizer.eos_id)
self.gguf_writer.add_eos_token_id(tokenizer.eos_id) # ty: ignore[unresolved-attribute]
else:
raise NotImplementedError(f"Deepseek pre-tokenizer {tokpre!r} is not supported yet!")
+3 -25
View File
@@ -57,8 +57,6 @@ class KimiVLModel(MmprojModel):
class KimiK25Model(MmprojModel):
"""Kimi-K2.5 with MoonViT3d vision encoder"""
projector_type = gguf.VisionProjectorType.KIMIK25
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
@@ -80,7 +78,7 @@ class KimiK25Model(MmprojModel):
super().set_gguf_parameters()
assert self.hparams_vision is not None
self.gguf_writer.add_clip_projector_type(self.projector_type)
self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.KIMIK25)
# Position embedding parameters (for interpolation)
self.gguf_writer.add_uint32("vision.pos_emb_height", self.hparams_vision.get("init_pos_emb_height", 64))
@@ -89,11 +87,7 @@ class KimiK25Model(MmprojModel):
# Projector parameters
self.gguf_writer.add_vision_use_gelu(self.hparams_vision.get("projector_hidden_act", "gelu") == "gelu")
if self.hparams_vision.get("norm_type", "layernorm") == "rmsnorm":
# nn.RMSNorm without eps uses finfo(float32).eps, also for bf16 inputs
self.gguf_writer.add_vision_attention_layernorm_eps(torch.finfo(torch.float32).eps)
else:
self.gguf_writer.add_vision_attention_layernorm_eps(self.hparams_vision.get("projector_ln_eps", 1e-5))
self.gguf_writer.add_vision_attention_layernorm_eps(self.hparams_vision.get("projector_ln_eps", 1e-5))
self.gguf_writer.add_vision_projector_scale_factor(self.merge_kernel_size[0])
# Image size limits
@@ -126,7 +120,7 @@ class KimiK25Model(MmprojModel):
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
assert self.hparams_vision is not None
n_head = self.find_vparam(["num_attention_heads", "vt_num_attention_heads"], optional=True) or 16
n_head = self.hparams_vision.get("num_attention_heads", 16)
# Permute Q/K weights/biases from interleaved to split RoPE format
# This allows using build_rope_2d at runtime without post-permutation.
@@ -177,19 +171,3 @@ class Glm5vModel(KimiK25Model):
name = name.replace("mm_projector.linear_", "mm_projector.proj.linear_", 1)
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("KimiK31ForConditionalGeneration")
@ModelBase.example("moonshotai/Kimi-K3.1-Pro")
class KimiK31VisionModel(KimiK25Model):
"""Kimi-K2.5 MoonViT3d with RMSNorm, a wider QKV and the patchmergerv2 projector"""
projector_type = gguf.VisionProjectorType.KIMIK31
def set_gguf_parameters(self):
super().set_gguf_parameters()
assert self.hparams_vision is not None
assert self.hparams_vision["norm_type"] == "rmsnorm"
assert self.hparams_vision["mm_projector_type"] == "patchmergerv2"
n_head = self.find_vparam(["vt_num_attention_heads"])
self.gguf_writer.add_vision_head_dim(self.hparams_vision["qkv_hidden_size"] // n_head)
+2 -2
View File
@@ -77,10 +77,10 @@ class _DecisionLoraMixin:
delta = self.lora_scale * (lora["B"] @ lora["A"])
data_torch = data_torch.float() + LazyTorchTensor.from_eager(delta)
self.lora_merged.add(name[name.index("layers."):])
yield from super().modify_tensors(data_torch, name, bid) # ty: ignore[missing-attribute]
yield from super().modify_tensors(data_torch, name, bid) # ty: ignore[unresolved-attribute]
def prepare_tensors(self):
super().prepare_tensors() # ty: ignore[missing-attribute]
super().prepare_tensors() # ty: ignore[unresolved-attribute]
if len(self.lora_merged) != len(self.lora):
raise ValueError(f"only {len(self.lora_merged)} of {len(self.lora)} LoRA tensors were merged into the base model")
+2 -3
View File
@@ -22,16 +22,15 @@ class LLaDAModel(TextModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
vocab_dict = tokenizer.get_vocab()
vocab_dict = tokenizer.get_vocab() # ty: ignore[unresolved-attribute]
vocab_size = self.hparams.get("vocab_size", len(vocab_dict))
assert max(vocab_dict.values()) < vocab_size
tokpre = self.get_vocab_base_pre(tokenizer)
reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in vocab_dict.items()}
added_vocab = tokenizer.get_added_vocab()
added_vocab = tokenizer.get_added_vocab() # ty: ignore[unresolved-attribute]
for i in range(vocab_size):
if i not in reverse_vocab:
-1
View File
@@ -263,7 +263,6 @@ class MiniCPMV4_7VisionModel(MiniCPMV4_6VisionModel):
# the tile separators are text tokens; clip appends their embeddings so that one chunk holds the whole image
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
ids = []
for text in self.tok_embd_sep:
tok = tokenizer.encode(text, add_special_tokens=False)
+1 -2
View File
@@ -25,8 +25,7 @@ class MiniMaxText01Model(TextModel):
# they get in the way of the token sampling process and must be suppressed
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
tokenizer_vocab_size = tokenizer.vocab_size
tokenizer_vocab_size = tokenizer.vocab_size # ty: ignore[unresolved-attribute]
with open(self.dir_model / "model.safetensors.index.json", "r", encoding="utf-8") as f:
weight_map = json.load(f)["weight_map"]
+3 -3
View File
@@ -12,9 +12,9 @@ from .deepseek import DeepseekV2Model
from .llama import LlamaModel
if _mistral_common_installed:
from mistral_common.tokens.tokenizers.base import TokenizerVersion # ty: ignore[import-not-found]
from mistral_common.tokens.tokenizers.tekken import Tekkenizer # ty: ignore[import-not-found]
from mistral_common.tokens.tokenizers.sentencepiece import SentencePieceTokenizer # ty: ignore[import-not-found]
from mistral_common.tokens.tokenizers.base import TokenizerVersion # type: ignore[import-not-found, ty:unresolved-import]
from mistral_common.tokens.tokenizers.tekken import Tekkenizer # type: ignore[import-not-found, ty:unresolved-import]
from mistral_common.tokens.tokenizers.sentencepiece import SentencePieceTokenizer # type: ignore[import-not-found, ty:unresolved-import]
else:
TokenizerVersion = None # type: ignore[assignment]
Tekkenizer = None # type: ignore[assignment]
+2 -3
View File
@@ -37,8 +37,7 @@ class MuseGlimmerModel(TextModel):
from transformers import AutoTokenizer
tok = AutoTokenizer.from_pretrained(self.dir_model)
assert tok is not None
eot_id = tok.convert_tokens_to_ids("<|eot|>")
eot_id = tok.convert_tokens_to_ids("<|eot|>") # ty: ignore[unresolved-attribute]
if isinstance(eot_id, int) and eot_id >= 0:
self.gguf_writer.add_eot_token_id(eot_id)
@@ -153,7 +152,7 @@ class MuseGlimmerAssistantModel(TextModel):
target_arch = json.load(f)["architectures"][0]
target_cls = get_model_class(target_arch)
if target_cls is not type(self):
target_cls.set_vocab(self) # ty: ignore[missing-attribute]
target_cls.set_vocab(self) # ty: ignore[unresolved-attribute]
else:
super().set_vocab()
+1 -2
View File
@@ -44,8 +44,7 @@ class MuseSparkModel(TextModel):
from transformers import AutoTokenizer
tok = AutoTokenizer.from_pretrained(self.dir_model)
assert tok is not None
eot_id = tok.convert_tokens_to_ids("<|eot|>")
eot_id = tok.convert_tokens_to_ids("<|eot|>") # ty: ignore[unresolved-attribute]
if isinstance(eot_id, int) and eot_id >= 0:
self.gguf_writer.add_eot_token_id(eot_id)
+6 -7
View File
@@ -73,7 +73,7 @@ class LiLiCorrDraftModel(DFlashModel):
try:
if target_arch == "NemotronHForCausalLM":
self.is_moe = "num_experts_per_tok" in target_text_config
target_cls.set_vocab(self) # ty: ignore[missing-attribute]
target_cls.set_vocab(self) # ty: ignore[unresolved-attribute]
finally:
self.dir_model = original_dir
self.gguf_writer.add_mask_token_id(self.hparams["dflash_config"]["mask_token_id"])
@@ -609,7 +609,6 @@ class NemotronHModel(GraniteHybridModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
assert tokenizer is not None
# Pad vocab size (from Mamba2Model/GraniteHybridModel)
self.hparams["pad_vocab_size_multiple"] = 8 # Setting this here since GraniteHybridModel.set_vocab() isn't being invoked now.
@@ -624,14 +623,14 @@ class NemotronHModel(GraniteHybridModel):
vocab_skip = getattr(self, "_vocab_skip", ())
user_defined_tokens = getattr(self, "_user_defined_tokens", None)
assert max(id_ for tok, id_ in tokenizer.vocab.items() if tok not in vocab_skip) < vocab_size
assert max(id_ for tok, id_ in tokenizer.vocab.items() if tok not in vocab_skip) < vocab_size # ty: ignore[unresolved-attribute]
tokpre = self.get_vocab_base_pre(tokenizer)
reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()}
added_vocab = tokenizer.get_added_vocab()
reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()} # ty: ignore[unresolved-attribute]
added_vocab = tokenizer.get_added_vocab() # ty: ignore[unresolved-attribute]
added_tokens_decoder = tokenizer.added_tokens_decoder
added_tokens_decoder = tokenizer.added_tokens_decoder # ty: ignore[unresolved-attribute]
for i in range(vocab_size):
if i not in reverse_vocab:
@@ -642,7 +641,7 @@ class NemotronHModel(GraniteHybridModel):
if token in added_vocab:
if not added_tokens_decoder[i].normalized:
previous_token = token
token = tokenizer.decode(tokenizer.encode(token, add_special_tokens=False)) # ty: ignore[invalid-assignment]
token = tokenizer.decode(tokenizer.encode(token, add_special_tokens=False)) # ty: ignore[unresolved-attribute, invalid-assignment]
if previous_token != token:
logger.info(f"{repr(previous_token)} is encoded and decoded back to {repr(token)} using AutoTokenizer")
+1 -1
View File
@@ -191,7 +191,7 @@ class PocketTTSModel(TextModel):
# unigram segmentation, so use the UGM tokenizer instead
from sentencepiece import sentencepiece_model_pb2 as model
proto = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
proto = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
proto.ParseFromString(open(self.dir_model / "tokenizer.model", "rb").read())
assert proto.trainer_spec.model_type == 1, "expected a unigram tokenizer"
+11 -12
View File
@@ -225,13 +225,12 @@ class Qwen3Model(Qwen2Model):
def _find_rerank_config(self):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
assert tokenizer is not None
self.is_rerank = True
self.is_tied_embeddings = self.hparams.get("tie_word_embeddings", False)
self.token_false_id = tokenizer.convert_tokens_to_ids("no") # ty: ignore[invalid-assignment]
self.token_true_id = tokenizer.convert_tokens_to_ids("yes") # ty: ignore[invalid-assignment]
self.sep_token_id = tokenizer.convert_tokens_to_ids("|")
self.token_false_id = tokenizer.convert_tokens_to_ids("no") # ty: ignore[unresolved-attribute, invalid-assignment]
self.token_true_id = tokenizer.convert_tokens_to_ids("yes") # ty: ignore[unresolved-attribute, invalid-assignment]
self.sep_token_id = tokenizer.convert_tokens_to_ids("|") # ty: ignore[unresolved-attribute]
assert self.token_false_id is not None and self.token_true_id is not None
@@ -330,7 +329,7 @@ class _QwenMtpMixin:
key = next((k for k in ["n_layers", "num_hidden_layers", "n_layer", "num_layers"] if k in hparams), None)
type(self)._original_block_count = hparams.get(key)
type(self).opt_num_mtp_layers = 0
return super().index_tensors(remote_hf_model_id=remote_hf_model_id) # ty: ignore[missing-attribute]
return super().index_tensors(remote_hf_model_id=remote_hf_model_id) # ty: ignore[unresolved-attribute]
@classmethod
def filter_tensors(cls, item):
@@ -367,7 +366,7 @@ class _QwenMtpMixin:
return name, gen
def set_gguf_parameters(self):
super().set_gguf_parameters() # ty: ignore[missing-attribute]
super().set_gguf_parameters() # ty: ignore[unresolved-attribute]
if self.no_mtp:
return
if (n := self.block_count - self.hparams["num_hidden_layers"]) > 0:
@@ -375,15 +374,15 @@ class _QwenMtpMixin:
def prepare_metadata(self, vocab_only: bool):
from_dir = self.fname_out.is_dir()
super().prepare_metadata(vocab_only=vocab_only) # ty: ignore[missing-attribute]
super().prepare_metadata(vocab_only=vocab_only) # ty: ignore[unresolved-attribute]
if not self.mtp_only or not from_dir:
return
output_type: str = self.ftype.name.partition("_")[2] # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
output_type: str = self.ftype.name.partition("_")[2] # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
fname_default: str = gguf.naming_convention(
self.metadata.name, self.metadata.basename, self.metadata.finetune, # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
self.metadata.version, size_label=None, output_type=output_type, model_type=None) # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
self.metadata.name, self.metadata.basename, self.metadata.finetune, # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
self.metadata.version, size_label=None, output_type=output_type, model_type=None) # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
self.fname_out = self.fname_out.parent / f"mtp-{fname_default}.gguf"
@@ -686,7 +685,7 @@ class _Qwen35MRopeMixin:
rope_parameters: dict
def set_gguf_parameters(self):
super().set_gguf_parameters() # ty: ignore[missing-attribute]
super().set_gguf_parameters() # ty: ignore[unresolved-attribute]
if "mrope_section" not in self.rope_parameters:
self.gguf_writer.add_rope_dimension_sections(self._QWEN35_DEFAULT_MROPE_SECTION)
@@ -844,7 +843,7 @@ class DFlashModel(Qwen3Model):
if target_cls is not type(self):
if target_arch == "NemotronHForCausalLM":
setattr(self, "is_moe", "num_experts_per_tok" in target_hparams)
target_cls.set_vocab(self) # ty: ignore[missing-attribute]
target_cls.set_vocab(self) # ty: ignore[unresolved-attribute]
else:
super().set_vocab()
+2 -2
View File
@@ -40,7 +40,7 @@ class T5Model(TextModel):
if not tokenizer_path.is_file():
raise FileNotFoundError(f"File not found: {tokenizer_path}")
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
sentencepiece_model.ParseFromString(open(tokenizer_path, "rb").read())
# some models like Pile-T5 family use BPE tokenizer instead of Unigram
@@ -178,7 +178,7 @@ class T5EncoderModel(TextModel):
if not tokenizer_path.is_file():
raise FileNotFoundError(f"File not found: {tokenizer_path}")
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[missing-attribute]
sentencepiece_model = model.ModelProto() # pyright: ignore[reportAttributeAccessIssue] # ty: ignore[unresolved-attribute]
sentencepiece_model.ParseFromString(open(tokenizer_path, "rb").read())
# some models like Pile-T5 family use BPE tokenizer instead of Unigram
+5 -6
View File
@@ -25,16 +25,15 @@ class XverseModel(TextModel):
from transformers import AutoTokenizer
tokenizer = AutoTokenizer.from_pretrained(dir_model)
assert tokenizer is not None
vocab_size = hparams.get("vocab_size", len(tokenizer.vocab))
vocab_size = hparams.get("vocab_size", len(tokenizer.vocab)) # ty: ignore[unresolved-attribute]
# Since we are checking the maximum index, we need to ensure it's strictly less than vocab_size,
# because vocab_size is the count of items, and indexes start at 0.
max_vocab_index = max(tokenizer.get_vocab().values())
max_vocab_index = max(tokenizer.get_vocab().values()) # ty: ignore[unresolved-attribute]
if max_vocab_index >= vocab_size:
raise ValueError("Vocabulary size exceeds expected maximum size.")
reverse_vocab: dict[int, str] = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()}
added_vocab = tokenizer.get_added_vocab()
reverse_vocab: dict[int, str] = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()} # ty: ignore[unresolved-attribute]
added_vocab = tokenizer.get_added_vocab() # ty: ignore[unresolved-attribute]
for token_id in range(vocab_size):
token_text = reverse_vocab[token_id].encode('utf-8')
@@ -45,7 +44,7 @@ class XverseModel(TextModel):
elif re.fullmatch(br"<0x[0-9A-Fa-f]{2}>", token_text):
toktype = gguf.TokenType.BYTE # special
elif reverse_vocab[token_id] in added_vocab:
if tokenizer.added_tokens_decoder[token_id].special:
if tokenizer.added_tokens_decoder[token_id].special: # ty: ignore[unresolved-attribute]
toktype = gguf.TokenType.CONTROL
else:
toktype = gguf.TokenType.USER_DEFINED
+2 -4
View File
@@ -325,8 +325,7 @@ for model in [*pre_computed_hashes, *all_models]:
except Exception as e:
raise OSError(f"Error loading tokenizer for model {name}.") from e
assert tokenizer is not None
chktok = tokenizer.encode(CHK_TXT)
chktok = tokenizer.encode(CHK_TXT) # ty: ignore[unresolved-attribute]
chkhsh = sha256(str(chktok).encode()).hexdigest()
logger.info(f"model: {name}")
@@ -487,7 +486,6 @@ for model in models:
logger.error(f"Failed to load tokenizer for model {name}. Error: {e}")
continue # Skip this model and continue with the next one in the loop
assert tokenizer is not None
if not os.path.exists(f"models/ggml-vocab-{name}.gguf"):
logger.info(f"Skip vocab files for model {name}, no GGUF file found")
continue
@@ -499,7 +497,7 @@ for model in models:
with open(f"models/ggml-vocab-{name}.gguf.out", "w") as f:
for text in tests:
res = tokenizer.encode(text, add_special_tokens=False)
res = tokenizer.encode(text, add_special_tokens=False) # ty: ignore[unresolved-attribute]
for r in res:
f.write(f" {r}")
f.write("\n")
+1 -2
View File
@@ -443,8 +443,7 @@ if __name__ == '__main__':
# the invocation string includes the "<|start_of_turn|>"
# token, but the adapters themselves were trained to
# activate _after_ that first token, so we drop it here.
assert tokenizer is not None
alora_invocation_tokens = tokenizer(invocation_string)["input_ids"][1:]
alora_invocation_tokens = tokenizer(invocation_string)["input_ids"][1:] # ty: ignore[call-non-callable]
if alora_invocation_tokens:
logger.debug("GGUF KV: %s = %s", gguf.Keys.Adapter.ALORA_INVOCATION_TOKENS, alora_invocation_tokens)
self.gguf_writer.add_key_value(
+1 -6
View File
@@ -52,9 +52,6 @@ The packages for FP32 and FP16 would have different accuracy and performance on
## News
- 2026.11
- Update the CI and release build for oneAPI 2026.2 and Level Zero SDK 1.34.0. oneDNN is updated to the 2026.1 stand-alone package.
- 2026.09
- Update the CI build environment for oneAPI 2026.1 (unified oneAPI Toolkit). oneDNN is removed from the Deep Learning Essentials package in 2026.0, so the CI now uses the oneAPI Toolkit installer which still includes oneDNN.
- oneAPI 2026.1 improves the SYCL build performance: measured with the same code on Arc B570, prompt processing 1331 vs 434 t/s (3.1x) vs the 2025.3-based release build.
@@ -288,7 +285,6 @@ Upon a successful installation, SYCL is enabled for the available Intel devices,
|Verified release|
|-|
|2026.2 |
|2026.1 |
|2025.3.3 |
|2025.2.1|
@@ -829,8 +825,7 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm
| GGML_SYCL_ENABLE_XMX_ESIMD | 0 or 1 (default) | Use the XMX (DPAS) multi-column mat-vec kernels for reordered Q2_K to Q6_K and Q8_0 weights on Xe2 GPUs. Speeds up speculative decoding and short prompt batches. |
| GGML_SYCL_XMX_MAX_COLS | non-negative integer, 80 (default) | Most activation columns (batch size) the `GGML_SYCL_ENABLE_XMX_ESIMD` kernels take; wider batches use the other mat-mul kernels. A larger value keeps XMX for bigger batches (as repeated launches), `0` never uses it. |
| GGML_SYCL_XMX_GLU_MAX_COLS | 0 to 16, 16 (default) | Most columns the fused gate and up XMX kernel takes; wider batches run gate and up as separate mat-vecs. 16 is the most the kernel is built for, larger values are clamped to 16 and `0` turns the fused kernel off. |
| GGML_SYCL_XMX_GATHER_TYPES | decimal bitmask, all bits set (default) | Weight formats that may use the XMX dequant-GEMM paths, which dequantize weights straight into the XMX tiles. This speeds up prompt processing of MoE models on GPUs with XMX units (Arc A- and B-series, Arc Pro, Data Center GPU Max), for example pp512 of Qwen3-30B-A3B UD-IQ3_XXS by about 50% on an Arc Pro B60. Bits:<br>* 1: IQ4_NL, 2: IQ3_S, 4: IQ4_XS, 8: IQ3_XXS, 16: IQ2_XXS, 32: IQ2_XS, 64: IQ2_S, 128: IQ1_S, 256: IQ1_M<br>* 512: Q8_0, 1024: Q4_K, 2048: Q5_K, 4096: Q6_K, 8192: Q2_K, 16384: Q3_K (MoE `MUL_MAT_ID` only)<br>Add values to combine them, for example `3` for IQ4_NL and IQ3_S; `0` disables the paths. A set bit does not force the path: batches of more than 64 tokens per expert use the library GEMM unless the device schedule width heuristic is overridden with `GGML_SYCL_MMID_SCHED=3`. Row lengths must be a multiple of 256 (32 for IQ4_NL and Q8_0). |
| GGML_SYCL_MMID_SCHED | decimal bitmask, 1 (default) | Build the grouped MoE routing and tile schedule on the GPU, avoiding the expert-ID host readback and CPU sort.<br>* 1: enable the device schedule.<br>* 2: override the fallback for models with few routed experts by disabling the performance heuristic that rejects batches averaging more than 64 routed rows per expert. This may reduce performance; type, shape, precision and device capability checks still apply.<br>Use `3` to enable the device schedule and override this heuristic, or `0` to use the host schedule. Bit 2 alone does not enable the device schedule. Higher bits are ignored. |
| GGML_SYCL_XMX_GATHER_TYPES | decimal bitmask, all bits set (default) | Weight formats that may use the XMX dequant-GEMM paths, which dequantize weights straight into the XMX tiles. This speeds up prompt processing of MoE models on GPUs with XMX units (Arc A- and B-series, Arc Pro, Data Center GPU Max), for example pp512 of Qwen3-30B-A3B UD-IQ3_XXS by about 50% on an Arc Pro B60. Bits:<br>* 1: IQ4_NL, 2: IQ3_S, 4: IQ4_XS, 8: IQ3_XXS, 16: IQ2_XXS, 32: IQ2_XS, 64: IQ2_S, 128: IQ1_S, 256: IQ1_M<br>* 512: Q8_0, 1024: Q4_K, 2048: Q5_K, 4096: Q6_K (MoE `MUL_MAT_ID` only)<br>Add values to combine them, for example `3` for IQ4_NL and IQ3_S; `0` disables the paths. A set bit does not force the path: batches of more than 64 tokens per expert or row lengths that are not a multiple of 256 (32 for IQ4_NL and Q8_0) use the library GEMM. |
| GGML_SYCL_XMX_GATHER_SHAPES | decimal bitmask, 255 (default) | XMX `joint_matrix` combinations the paths of `GGML_SYCL_XMX_GATHER_TYPES` may use; the operand type comes from `GGML_SYCL_DYNAMIC_PRECISION` and the best supported combination is picked automatically (logged as `fg_pick_combo`). Bits:<br>* Xe2, Xe3, Xe-HPC: 1: f16 8x16x16, 2: f16 16x16x16, 4: f16 32x64x16, 8: f16 32x64x32, 32: tf32 8x16x8, 64: bf16 8x16x16<br>* Xe-HPG (Arc A770, ARL-H): 16: f16 8x8x16, 128: bf16 8x8x16<br>Clear a bit to exclude a combination, or set a single bit to force one for testing. |
| GGML_SYCL_DYNAMIC_PRECISION | `F16` (default with `GGML_SYCL_F16=ON`), `BF16`, `TF32` or `F32` (default otherwise) | Operand type of the XMX dequant-GEMM paths (`GGML_SYCL_XMX_GATHER_TYPES`); accumulation is always f32. `F16` is the fastest, but activations above 65504 overflow. `BF16` keeps the f32 range at a 7-bit mantissa, `TF32` keeps the range and the f16 mantissa but is about 30% slower and needs Xe2, Xe3 or Xe-HPC, and `F32` turns the XMX paths off. Ops that request a higher src1 precision ([TAG_GGML_PREC]) get it regardless of this setting. |
| GGML_SYCL_DYNAMIC_REQUIRED_PRECISION | `F32` (default), `TF32`, `BF16` or `F16` | Lowest type the XMX paths may use for an op that requests an F32 src1, such as Mistral 4 `ffn_down_exps`. The default runs such ops on the library f32 GEMM; `TF32` or `BF16` trade mantissa for speed while keeping the f32 range. `F16` ignores the request and can overflow; it is meant for testing only. |
+1 -1
View File
@@ -86,7 +86,7 @@ Legend:
| OPT_STEP_SGD | ❌ | ❌ | ✅ | ✅ | ❌ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| OUT_PROD | 🟡 | 🟡 | 🟡 | 🟡 | ❌ | ❌ | ❌ | ❌ | 🟡 | 🟡 | ❌ | ❌ | ❌ |
| PAD | ❌ | 🟡 | ✅ | 🟡 | ❌ | ✅ | 🟡 | 🟡 | 🟡 | ✅ | ✅ | ❌ | ❌ |
| PAD_REFLECT_1D | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ✅ | ❌ | ✅ | ❌ | ✅ | ❌ | ❌ |
| PAD_REFLECT_1D | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ | ✅ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
| POOL_1D | ❌ | ❌ | ✅ | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| POOL_2D | ❌ | 🟡 | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
| REGLU | ❌ | ✅ | ✅ | ✅ | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
+2 -7
View File
@@ -12336,13 +12336,8 @@
"WebGPU: WebGPU","PAD","type=f32,ne_a=[100,1,1,1],pad_0=100,pad_1=0,circular=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD","type=f32,ne_a=[100,1,1,1],pad_0=0,pad_1=100,circular=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD","type=f32,ne_a=[100,100,1,1],pad_0=50,pad_1=50,circular=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[512,34,2,1],pad_0=10,pad_1=9,view=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[1,3,2,2],pad_0=0,pad_1=0,view=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[5,3,2,2],pad_0=4,pad_1=1,view=1","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[281,128,1,1],pad_0=2,pad_1=2,view=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[257,7,1,1],pad_0=0,pad_1=256,view=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[2,16777217,1,1],pad_0=1,pad_1=1,view=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[3000,384,4,1],pad_0=10,pad_1=9,view=0","support","1","yes","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[512,34,2,1],pad_0=10,pad_1=9","support","0","no","WebGPU"
"WebGPU: WebGPU","PAD_REFLECT_1D","type=f32,ne_a=[3000,384,4,1],pad_0=10,pad_1=9","support","0","no","WebGPU"
"WebGPU: WebGPU","ROLL","shift0=3,shift1=-2,shift3=1,shift4=-1","support","0","no","WebGPU"
"WebGPU: WebGPU","ARANGE","type=f32,start=0.000000,stop=10.000000,step=1.000000","support","1","yes","WebGPU"
"WebGPU: WebGPU","ARANGE","type=f32,start=3.000000,stop=3.500000,step=1.000000","support","1","yes","WebGPU"
Can't render this file because it is too large.
@@ -7,7 +7,7 @@ import os
# Add utils directory to path for direct script execution
sys.path.insert(0, str(Path(__file__).parent.parent / "utils"))
from common import get_model_name_from_env_path, compare_tokens, exit_with_warning # ty: ignore[import-not-found]
from common import get_model_name_from_env_path, compare_tokens, exit_with_warning # type: ignore[import-not-found, ty:unresolved-import]
def quick_logits_check(pytorch_file, llamacpp_file):
"""Lightweight sanity check before NMSE"""
@@ -33,7 +33,6 @@ print("EOS token id: ", config.eos_token_id)
print("Loading model and tokenizer using AutoTokenizer:", model_path)
tokenizer = AutoTokenizer.from_pretrained(model_path)
assert tokenizer is not None
if unreleased_model_name:
model_name_lower = unreleased_model_name.lower()
@@ -57,11 +56,11 @@ model_name = os.path.basename(model_path)
print(f"Model name: {model_name}")
prompt = "Hello world today"
input_ids = tokenizer(prompt, return_tensors="pt").input_ids
input_ids = tokenizer(prompt, return_tensors="pt").input_ids # ty: ignore[call-non-callable]
token_ids = input_ids[0].cpu().tolist()
print(f"Input tokens: {input_ids}")
print(f"Input text: {repr(prompt)}")
print(f"Tokenized: {tokenizer.convert_ids_to_tokens(input_ids[0])}")
print(f"Tokenized: {tokenizer.convert_ids_to_tokens(input_ids[0])}") # ty: ignore[unresolved-attribute]
with torch.no_grad():
outputs = model(input_ids, output_hidden_states=True)
@@ -84,7 +83,7 @@ with torch.no_grad():
# Print embeddings per token in the requested format
print("\nToken embeddings:")
tokens = tokenizer.convert_ids_to_tokens(input_ids[0])
tokens = tokenizer.convert_ids_to_tokens(input_ids[0]) # ty: ignore[unresolved-attribute]
for i, embedding in enumerate(token_embeddings):
# Format: show first few values, ..., then last few values
if len(embedding) > 10:
@@ -64,7 +64,7 @@ def load_model_and_tokenizer(model_path, use_sentence_transformers=False, device
print("Using SentenceTransformer to apply all numbered layers")
model = SentenceTransformer(model_path)
tokenizer = model.tokenizer
config = model[0].auto_model.config # ty: ignore[missing-union-attribute]
config = model[0].auto_model.config # ty: ignore[unresolved-attribute]
else:
tokenizer = AutoTokenizer.from_pretrained(model_path)
config = AutoConfig.from_pretrained(model_path, trust_remote_code=True)
@@ -5,7 +5,7 @@ import sys
import os
import argparse
from pathlib import Path
from common import get_model_name_from_env_path # ty: ignore[import-not-found]
from common import get_model_name_from_env_path # type: ignore[import-not-found, ty:unresolved-import]
def calculate_nmse(reference, test):
mse = np.mean((test - reference) ** 2)
@@ -2,7 +2,7 @@
import argparse
import sys
from common import compare_tokens # ty: ignore[import-not-found]
from common import compare_tokens # type: ignore[import-not-found, ty:unresolved-import]
def parse_arguments():
@@ -7,7 +7,7 @@ import importlib
from pathlib import Path
from transformers import AutoTokenizer, AutoConfig, AutoModelForCausalLM, AutoModel
from common import compare_tokens, exit_with_warning # ty: ignore[import-not-found]
from common import compare_tokens, exit_with_warning # type: ignore[import-not-found, ty:unresolved-import]
unreleased_model_name = os.getenv('UNRELEASED_MODEL_NAME')
@@ -183,7 +183,6 @@ def main():
# Load the python model to get configuration information and also to load the tokenizer.
print("Loading model and tokenizer using AutoTokenizer:", args.model_path)
tokenizer = AutoTokenizer.from_pretrained(args.model_path)
assert tokenizer is not None
config = AutoConfig.from_pretrained(args.model_path, trust_remote_code=True)
if unreleased_model_name:
@@ -208,8 +207,8 @@ def main():
else:
model = AutoModel.from_pretrained(args.model_path, trust_remote_code=True)
encoded = tokenizer(prompt, return_tensors="pt")
tokens = tokenizer.convert_ids_to_tokens(encoded['input_ids'][0])
encoded = tokenizer(prompt, return_tensors="pt") # ty: ignore[call-non-callable]
tokens = tokenizer.convert_ids_to_tokens(encoded['input_ids'][0]) # ty: ignore[unresolved-attribute]
n_tokens = len(tokens)
print(f"n_tokens: {n_tokens}");
print(f"hidden_size: {model.config.hidden_size}")
+1 -1
View File
@@ -1285,7 +1285,7 @@ def convert_dictionary_to_pydantic_model(dictionary: dict[str, Any], model_name:
if items != {}:
array = {"properties": items}
array_type = convert_dictionary_to_pydantic_model(array, f"{model_name}_{field_name}_items")
fields[field_name] = (list[array_type], ...)
fields[field_name] = (list[array_type], ...) # ty: ignore[invalid-type-form]
else:
fields[field_name] = (list, ...)
elif field_type == "object":
+1 -1
View File
@@ -205,7 +205,7 @@ option(GGML_CUDA_NO_PEER_COPY "ggml: do not use peer to peer copie
option(GGML_CUDA_NO_VMM "ggml: do not try to use CUDA VMM" OFF)
option(GGML_CUDA_FA "ggml: compile ggml FlashAttention CUDA kernels" ON)
option(GGML_CUDA_FA_ALL_QUANTS "ggml: compile all quants for FlashAttention" OFF)
set (GGML_CUDA_FA_QUANTS "q4_0-q4_0;q8_0-q8_0;f16-f16;bf16-bf16;f8_e4m3-f8_e4m3" CACHE STRING
set (GGML_CUDA_FA_QUANTS "q4_0-q4_0;q8_0-q8_0;f16-f16;bf16-bf16" CACHE STRING
"ggml: FlashAttention K-V type combinations to compile, \"all\" or a list such as \"q8_0-q8_0;q8_0-q4_0\"")
option(GGML_CUDA_GRAPHS "ggml: use CUDA graphs (llama.cpp only)" ${GGML_CUDA_GRAPHS_DEFAULT})
option(GGML_CUDA_NCCL "ggml: use NVIDIA Collective Comm. Library" ON)
+1 -1
View File
@@ -58,7 +58,7 @@ endfunction()
# Determines which FlashAttention vector kernel template instances to compile, returns them in OUT_SRCS.
function(ggml_cuda_fattn_vec_instances DIR OUT_SRCS)
set(FA_TYPES q4_0 q4_1 q5_0 q5_1 q8_0 bf16 f16 f8_e4m3)
set(FA_TYPES q4_0 q4_1 q5_0 q5_1 q8_0 bf16 f16)
string(TOLOWER "${GGML_CUDA_FA_QUANTS}" FA_QUANTS)
string(STRIP "${FA_QUANTS}" FA_QUANTS)
+9 -3
View File
@@ -131,9 +131,15 @@ endif()
# POSIX conformance
#
# POSIX.1-2001 with XSI extensions provides clock_gettime, CLOCK_MONOTONIC, posix_memalign and M_PI.
# OpenBSD needs POSIX.1-2008 for locale_t. FreeBSD needs it for F_DUPFD_CLOEXEC and strnlen.
if (CMAKE_SYSTEM_NAME MATCHES "OpenBSD|FreeBSD")
# clock_gettime came in POSIX.1b (1993)
# CLOCK_MONOTONIC came in POSIX.1-2001 / SUSv3 as optional
# posix_memalign came in POSIX.1-2001 / SUSv3
# M_PI is an XSI extension since POSIX.1-2001 / SUSv3, came in XPG1 (1985)
# Somehow in OpenBSD whenever POSIX conformance is specified
# some string functions rely on locale_t availability,
# which was introduced in POSIX.1-2008, forcing us to go higher
if (CMAKE_SYSTEM_NAME MATCHES "OpenBSD")
add_compile_definitions(_XOPEN_SOURCE=700)
elseif (CMAKE_SYSTEM_NAME MATCHES "AIX")
# Don't define _XOPEN_SOURCE. We need _ALL_SOURCE, which is the default,
+28 -33
View File
@@ -1398,35 +1398,34 @@ static void ggml_backend_cann_buffer_get_tensor(ggml_backend_buffer_t buffer,
static bool ggml_backend_cann_buffer_cpy_tensor(ggml_backend_buffer_t buffer,
const ggml_tensor * src,
ggml_tensor * dst) {
if (!ggml_backend_buft_is_cann(src->buffer->buft)) {
return false;
}
if (ggml_backend_buft_is_cann(src->buffer->buft)) {
ggml_backend_cann_buffer_context * src_ctx = (ggml_backend_cann_buffer_context *) src->buffer->context;
ggml_backend_cann_buffer_context * dst_ctx = (ggml_backend_cann_buffer_context *) buffer->context;
ggml_backend_cann_buffer_context * src_ctx = (ggml_backend_cann_buffer_context *) src->buffer->context;
ggml_backend_cann_buffer_context * dst_ctx = (ggml_backend_cann_buffer_context *) buffer->context;
size_t memcpy_size = ggml_nbytes(src);
// Same device.
if (src_ctx->device == dst_ctx->device) {
ACL_CHECK(aclrtMemcpy((char *) dst->data, memcpy_size, (const char *) src->data, memcpy_size,
ACL_MEMCPY_DEVICE_TO_DEVICE));
return true;
} else {
#ifdef ASCEND_310P
// TODO: Support 310p P2P copy
return false;
#endif
// Different device but can access by peer.
int32_t canAccessPeer = 0;
ACL_CHECK(aclrtDeviceCanAccessPeer(&canAccessPeer, src_ctx->device, dst_ctx->device));
if (canAccessPeer) {
ggml_cann_set_device(src_ctx->device);
ACL_CHECK(aclrtDeviceEnablePeerAccess(dst_ctx->device, 0));
size_t memcpy_size = ggml_nbytes(src);
// Same device.
if (src_ctx->device == dst_ctx->device) {
ACL_CHECK(aclrtMemcpy((char *) dst->data, memcpy_size, (const char *) src->data, memcpy_size,
ACL_MEMCPY_DEVICE_TO_DEVICE));
ACL_MEMCPY_DEVICE_TO_DEVICE));
return true;
} else {
#ifdef ASCEND_310P
// TODO: Support 310p P2P copy
return false;
#endif
// Different device but can access by peer.
int32_t canAccessPeer = 0;
ACL_CHECK(aclrtDeviceCanAccessPeer(&canAccessPeer, src_ctx->device, dst_ctx->device));
if (canAccessPeer) {
ggml_cann_set_device(src_ctx->device);
ACL_CHECK(aclrtDeviceEnablePeerAccess(dst_ctx->device, 0));
ACL_CHECK(aclrtMemcpy((char *) dst->data, memcpy_size, (const char *) src->data, memcpy_size,
ACL_MEMCPY_DEVICE_TO_DEVICE));
return true;
}
}
}
return false;
}
/**
@@ -2447,9 +2446,6 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
break;
case GGML_OP_MUL_MAT:
{
if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
return false;
}
switch (op->src[0]->type) {
#ifndef ASCEND_310P
case GGML_TYPE_BF16:
@@ -2868,13 +2864,12 @@ static ggml_backend_t ggml_backend_cann_device_init(ggml_backend_dev_t dev, cons
* otherwise false.
*/
static bool ggml_backend_cann_supports_buft(ggml_backend_dev_t dev, ggml_backend_buffer_type_t buft) {
if (!ggml_backend_buft_is_cann(buft)) {
return false;
if (ggml_backend_buft_is_cann(buft)) {
ggml_backend_cann_device_context * dev_ctx = (ggml_backend_cann_device_context *) dev->context;
ggml_backend_cann_buffer_type_context * buft_ctx = (ggml_backend_cann_buffer_type_context *) buft->context;
return buft_ctx->device == dev_ctx->device;
}
ggml_backend_cann_device_context * dev_ctx = (ggml_backend_cann_device_context *) dev->context;
ggml_backend_cann_buffer_type_context * buft_ctx = (ggml_backend_cann_buffer_type_context *) buft->context;
return buft_ctx->device == dev_ctx->device;
return false;
}
static ggml_backend_buffer_type_t ggml_backend_cann_device_get_buffer_type(ggml_backend_dev_t dev) {
+48 -22
View File
@@ -8911,9 +8911,14 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
float * VKQ32 = (float *) params->wdata + ith*(1*DK + 2*DV + CACHE_LINE_SIZE_F32); // FP32 VKQ accumulator
float * V32 = (VKQ32 + 1*DV); // (temporary) FP32 V buffer
ggml_fp16_t * VKQ16 = (ggml_fp16_t *) (VKQ32 + 1*DV); // (temporary) FP16 VKQ accumulator
ggml_fp16_t * Q_q = (ggml_fp16_t *) (VKQ32 + 2*DV); // (temporary) buffer for Q converted to quantized/FP16
memset(VKQ32, 0, DV*sizeof(float));
if (v->type == GGML_TYPE_F16) {
memset(VKQ16, 0, DV*sizeof(ggml_fp16_t));
} else {
memset(VKQ32, 0, DV*sizeof(float));
}
const ggml_fp16_t * mp = mask ? (ggml_fp16_t *)((char *) mask->data + iq1*mask->nb[1] + (iq2%mask->ne[2])*mask->nb[2] + (iq3%mask->ne[3])*mask->nb[3]) : NULL;
@@ -8965,32 +8970,53 @@ static void ggml_compute_forward_flash_attn_ext_f16_one_chunk(
const char * v_data = ((const char *) v->data + (ik1*nbv1 + iv2*nbv2 + iv3*nbv3));
if (s > M) {
// s is new maximum, ms < 1.0f, vs == expf(s - s) == 1.0f
M = s;
ms = expf(Mold - M);
// V = V*expf(Mold - M)
ggml_vec_scale_f32(DV, VKQ32, ms);
} else {
// no new maximum, ms == 1.0f, vs != 1.0f
vs = expf(s - M);
}
// V += v*expf(s - M)
if (v->type == GGML_TYPE_F16) {
ggml_vec_mad_f16_f32(DV, VKQ32, (const ggml_fp16_t *) v_data, vs);
} else if (v_to_float) {
v_to_float(v_data, V32, DV);
ggml_vec_mad_f32(DV, VKQ32, V32, vs);
if (s > M) {
// s is new maximum, ms < 1.0f, vs == expf(s - s) == 1.0f
M = s;
ms = expf(Mold - M);
// V = V*expf(Mold - M)
ggml_vec_scale_f16(DV, VKQ16, ms);
} else {
// no new maximum, ms == 1.0f, vs != 1.0f
vs = expf(s - M);
}
// V += v*expf(s - M)
ggml_vec_mad_f16(DV, VKQ16, (const ggml_fp16_t *) v_data, vs);
} else {
// V is F32
ggml_vec_mad_f32(DV, VKQ32, (const float *) v_data, vs);
if (s > M) {
// s is new maximum, ms < 1.0f, vs == expf(s - s) == 1.0f
M = s;
ms = expf(Mold - M);
// V = V*expf(Mold - M)
ggml_vec_scale_f32(DV, VKQ32, ms);
} else {
// no new maximum, ms == 1.0f, vs != 1.0f
vs = expf(s - M);
}
// V += v*expf(s - M)
if (v_to_float) {
v_to_float(v_data, V32, DV);
ggml_vec_mad_f32(DV, VKQ32, V32, vs);
} else {
// V is F32
ggml_vec_mad_f32(DV, VKQ32, (const float *) v_data, vs);
}
}
S = S*ms + vs; // scale and increment sum with partial sum
}
if (v->type == GGML_TYPE_F16) {
for (int64_t d = 0; d < DV; ++d) {
VKQ32[d] = GGML_CPU_FP16_TO_FP32(VKQ16[d]);
}
}
// sinks - apply only on the first kv-chunk
if (sinks && ic_start == 0) {
const float s = ((float *)((char *) sinks->data))[h];
@@ -9543,8 +9569,8 @@ static void ggml_compute_forward_flash_attn_ext_f16(
void ggml_compute_forward_flash_attn_ext(
const ggml_compute_params * params,
ggml_tensor * dst) {
switch (ggml_prec_get_acc(dst)) {
case GGML_PREC_UNDEFINED:
switch (dst->op_params[3]) {
case GGML_PREC_DEFAULT:
case GGML_PREC_F32:
case GGML_PREC_BF16:
{
+1 -1
View File
@@ -1180,7 +1180,7 @@ class tensor_traits_common : public tensor_traits_base {
const int64_t DK = nek0;
const int64_t DV = nev0;
const bool supported_prec = (ggml_prec_get_acc(dst) == GGML_PREC_F32 || ggml_prec_get_acc(dst) == GGML_PREC_BF16 || ggml_prec_get_acc(dst) == GGML_PREC_UNDEFINED);
const bool supported_prec = (dst->op_params[3] == GGML_PREC_F32 || dst->op_params[3] == GGML_PREC_BF16 || dst->op_params[3] == GGML_PREC_DEFAULT);
const bool supported_types = (q->type == GGML_TYPE_F32 && k->type == GGML_TYPE_F16 && v->type == GGML_TYPE_F16);
const bool supported_shape = (DK > 0 && DK <= 128 && DV > 0 && DV <= 128);
const bool supported_vlen = (__riscv_vlenb() == 128);
-35
View File
@@ -436,41 +436,6 @@ inline static void ggml_vec_mad_f32(const int n, float * GGML_RESTRICT y, const
#endif
}
// y (F32) += x (F16)*v
// on AVX512, AVX with F16C, NEON and RVV, x is widened to F32 in-register and fused into the multiply-add
inline static void ggml_vec_mad_f16_f32(const int n, float * GGML_RESTRICT y, const ggml_fp16_t * GGML_RESTRICT x, const float v) {
int i = 0;
#if defined(__AVX512F__)
const __m512 vx = _mm512_set1_ps(v);
for (; i + 16 <= n; i += 16) {
const __m512 ax = _mm512_cvtph_ps(_mm256_loadu_si256((const __m256i *)(x + i)));
_mm512_storeu_ps(y + i, _mm512_fmadd_ps(ax, vx, _mm512_loadu_ps(y + i)));
}
#elif defined(__AVX__) && defined(__F16C__)
const __m256 vx = _mm256_set1_ps(v);
for (; i + 8 <= n; i += 8) {
const __m256 ax = _mm256_cvtph_ps(_mm_loadu_si128((const __m128i *)(x + i)));
_mm256_storeu_ps(y + i, GGML_F32x8_FMA(_mm256_loadu_ps(y + i), ax, vx));
}
#elif defined(__ARM_NEON) && defined(__ARM_FEATURE_FMA) && defined(__ARM_FP16_FORMAT_IEEE)
const float32x4_t vx = vdupq_n_f32(v);
for (; i + 8 <= n; i += 8) {
const float16x8_t ax = vld1q_f16((const __fp16 *)(x + i));
vst1q_f32(y + i + 0, vfmaq_f32(vld1q_f32(y + i + 0), vcvt_f32_f16(vget_low_f16 (ax)), vx));
vst1q_f32(y + i + 4, vfmaq_f32(vld1q_f32(y + i + 4), vcvt_f32_f16(vget_high_f16(ax)), vx));
}
#elif defined(__riscv_v_intrinsic) && defined(__riscv_zvfhmin)
for (int vl; i < n; i += vl) {
vl = __riscv_vsetvl_e16m2(n - i);
const vfloat32m4_t ax = __riscv_vfwcvt_f_f_v_f32m4(__riscv_vle16_v_f16m2((const _Float16 *)x + i, vl), vl);
__riscv_vse32_v_f32m4(y + i, __riscv_vfmacc_vf_f32m4(__riscv_vle32_v_f32m4(y + i, vl), v, ax, vl), vl);
}
#endif
for (; i < n; ++i) {
y[i] += GGML_CPU_FP16_TO_FP32(x[i])*v;
}
}
inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * GGML_RESTRICT y, const ggml_fp16_t * GGML_RESTRICT x, const float v) {
#if defined(GGML_SIMD) && defined(__ARM_FEATURE_SVE)
const int sve_register_length = svcntb() * 8;
+2 -2
View File
@@ -163,7 +163,7 @@ if (CUDAToolkit_FOUND)
if (GGML_STATIC)
if (WIN32)
# As of 12.3.1 CUDA Toolkit for Windows does not offer a static cublas library
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart_static CUDA::cublas CUDA::cublasLt)
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart_static CUDA::cublas)
else ()
if (CUDAToolkit_VERSION VERSION_GREATER_EQUAL "10.1")
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart_static CUDA::cublas_static CUDA::cublasLt_static)
@@ -172,7 +172,7 @@ if (CUDAToolkit_FOUND)
endif()
endif()
else()
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart CUDA::cublas CUDA::cublasLt)
target_link_libraries(ggml-cuda PRIVATE CUDA::cudart CUDA::cublas)
endif()
if (GGML_CUDA_NO_VMM)
+290 -91
View File
@@ -7,6 +7,7 @@
#include <cstdint>
#include <cstdlib>
#include <memory>
#include <mutex>
#if defined(GGML_USE_HIP)
#define GGML_COMMON_DECL_HIP
@@ -388,10 +389,6 @@ static bool blackwell_mma_available(const int cc) {
ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_RUBIN;
}
static bool fp8_mma_hardware_available(const int cc) {
return GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_ADA_LOVELACE && cc < GGML_CUDA_CC_RUBIN;
}
// Checks whether the tensor's base data pointer and higher-dimensional strides are byte-aligned to `alignment` bytes.
static bool ggml_cuda_is_aligned(const ggml_tensor * tensor, const size_t alignment) {
GGML_ASSERT(tensor != nullptr);
@@ -887,67 +884,6 @@ static __device__ __forceinline__ float ggml_cuda_ue4m3_to_fp32(uint8_t x) {
#endif // defined(GGML_USE_HIP) && defined(CDNA3) && defined(FP8_AVAILABLE) && HIP_VERSION >= 60200000
}
static __host__ __device__ __forceinline__ float ggml_cuda_f8_e4m3_to_fp32(uint8_t x) {
#if defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
__nv_fp8_e4m3 xf;
xf.__x = x;
return static_cast<float>(xf);
#else
const uint8_t magnitude = x & 0x7F;
if (magnitude == 0x7F) {
return NAN;
}
const int exp = (magnitude >> 3) & 0x0F;
const int man = magnitude & 0x07;
float value;
if (exp == 0) {
value = ldexpf((float) man, -9);
} else {
value = ldexpf(1.0f + (float) man / 8.0f, exp - 7);
}
return x & 0x80 ? -value : value;
#endif // defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
}
static __host__ __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_f8_e4m3(float x) {
#if defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
return __nv_cvt_float_to_fp8(x, __NV_SATFINITE, __NV_E4M3);
#elif !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) && !defined(__MUSA_ARCH__)
// ggml_cuda_cast also runs on the host, where __float2int_rn is not available.
return ggml_fp32_to_f8_e4m3(x);
#else
const uint8_t sign = signbit(x) ? 0x80 : 0;
x = fabsf(x);
if (isnan(x)) {
return sign | 0x7F;
}
if (x == 0.0f) {
return sign;
}
if (isinf(x) || x >= 448.0f) {
return sign | 0x7E;
}
if (x < 0.015625f) {
return sign | (uint8_t) __float2int_rn(x * 512.0f);
}
int exp;
const float mantissa = frexpf(x, &exp) * 2.0f;
int encoded_exp = exp + 6;
int encoded_man = __float2int_rn((mantissa - 1.0f) * 8.0f);
if (encoded_man == 8) {
encoded_man = 0;
encoded_exp++;
}
if (encoded_exp > 15 || (encoded_exp == 15 && encoded_man > 6)) {
return sign | 0x7E;
}
return sign | (uint8_t) (encoded_exp << 3) | (uint8_t) encoded_man;
#endif // defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
}
static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_ue4m3(float x) {
#if defined(BLACKWELL_MMA_AVAILABLE) // This is used for NVFP4 subblock scale quantizations only
if (!(x > 0.0f)) {
@@ -1335,9 +1271,186 @@ struct ggml_tensor_extra_gpu {
#define USE_CUDA_GRAPH
#endif
// defined in graph.cuh
struct ggml_cuda_graph;
struct ggml_cuda_stream_context;
struct ggml_cuda_graph {
#ifdef USE_CUDA_GRAPH
~ggml_cuda_graph() {
for (cudaGraphExec_t instance : instances) {
CUDA_CHECK(cudaGraphExecDestroy(instance));
}
}
// the ggml graph is captured as a series of CUDA graphs, so the GPU can run one chunk while the next is enqueued
std::vector<cudaGraphExec_t> instances;
size_t num_nodes = 0;
std::vector<cudaGraphNode_t> nodes;
bool disable_due_to_gpu_arch = false;
bool warmup_complete = false;
uint64_t uid = 0;
int64_t last_used_time = 0;
struct node_properties {
ggml_tensor node;
void * node_src_data_ptrs[GGML_MAX_SRC];
int64_t node_src_ne[GGML_MAX_SRC][GGML_MAX_DIMS];
size_t node_src_nb[GGML_MAX_SRC][GGML_MAX_DIMS];
};
std::vector<node_properties> node_props;
bool is_enabled() const {
static const bool disable_cuda_graphs_due_to_env = (getenv("GGML_CUDA_DISABLE_GRAPHS") != nullptr);
return !(disable_due_to_gpu_arch || disable_cuda_graphs_due_to_env);
}
#endif
};
struct ggml_cuda_concurrent_event {
std::vector<cudaEvent_t> join_events;
cudaEvent_t fork_event = nullptr;
int n_streams = 0;
std::unordered_map<const ggml_tensor *, int> stream_mapping;
// Original order of nodes in this concurrent region (before interleaving)
// Used to restore grouping for fusion within streams
std::vector<const ggml_tensor *> original_order;
const ggml_tensor * join_node;
ggml_cuda_concurrent_event() = default;
ggml_cuda_concurrent_event(const ggml_cuda_concurrent_event &) = delete;
ggml_cuda_concurrent_event & operator=(const ggml_cuda_concurrent_event &) = delete;
explicit ggml_cuda_concurrent_event(int n_streams) : n_streams(n_streams) {
join_events.resize(n_streams);
for (size_t i = 0; i < join_events.size(); ++i) {
CUDA_CHECK(cudaEventCreateWithFlags(&join_events[i], cudaEventDisableTiming));
}
CUDA_CHECK(cudaEventCreateWithFlags(&fork_event, cudaEventDisableTiming));
}
ggml_cuda_concurrent_event(ggml_cuda_concurrent_event && other) noexcept
: join_events(std::move(other.join_events))
, fork_event(other.fork_event)
, n_streams(other.n_streams)
, stream_mapping(std::move(other.stream_mapping))
, original_order(std::move(other.original_order))
, join_node(other.join_node) {
other.fork_event = nullptr;
}
// 1. check if any branches write to overlapping memory ranges (except the join node)
// 2. check whether all srcs are either within the branch or outside the nodes covered by ggml_cuda_concurrent_event
// we assume all nodes have the same buffer
bool is_valid() const {
std::vector<std::vector<std::pair<int64_t, int64_t>>> write_ranges;
write_ranges.resize(n_streams);
// get join_node's memory range to exclude from overlap checking.
// multiple nodes can use join_node's buffer; we synchronize on the join node.
const ggml_tensor * join_t = join_node->view_src ? join_node->view_src : join_node;
const int64_t join_start = (int64_t) join_t->data;
const int64_t join_end = join_start + ggml_nbytes(join_t);
for (const auto & [tensor, stream] : stream_mapping) {
const ggml_tensor * t = tensor->view_src ? tensor->view_src : tensor;
const int64_t t_start = (int64_t) t->data;
const int64_t t_end = t_start + ggml_nbytes(t);
// skip tensors that overlap with join_node's buffer.
if ((t_start <= join_start && join_start < t_end) || (join_start <= t_start && t_start < join_end)) {
continue;
}
// concurrent streams begin from 1
write_ranges[stream - 1].emplace_back(t_start, t_end);
}
for (int i = 0; i < n_streams; ++i) {
// sorts first by start then by end of write range
std::sort(write_ranges[i].begin(), write_ranges[i].end());
}
bool writes_overlap = false;
bool dependent_srcs = false;
for (const auto & [tensor, stream] : stream_mapping) {
const ggml_tensor * t = tensor->view_src ? tensor->view_src : tensor;
const int64_t t_start = (int64_t) t->data;
const int64_t t_end = t_start + ggml_nbytes(t);
// skip tensors that overlap with join_node's buffer
if ((t_start <= join_start && join_start < t_end) || (join_start <= t_start && t_start < join_end)) {
continue;
}
// check if this buffer's write data overlaps with another stream's
std::pair<int64_t, int64_t> data_range = std::make_pair(t_start, t_end);
for (int i = 0; i < n_streams; ++i) {
if (i == stream - 1) {
continue;
}
auto it = std::lower_bound(write_ranges[i].begin(), write_ranges[i].end(), data_range);
if (it != write_ranges[i].end()) {
const std::pair<int64_t, int64_t> & other = *it;
// std::lower_bound returns the first element where other >= data_range (lexicographically).
// This guarantees other.first >= data_range.first.
// Therefore, overlap occurs iff other.first < data_range.second
// (i.e., the other range starts before this range ends).
if (other.first < data_range.second) {
GGML_LOG_DEBUG("Writes overlap for %s", tensor->name);
writes_overlap = true;
break;
}
}
}
//check if all srcs are either in branch or don't have a branch
for (int i = 0; i < GGML_MAX_SRC; ++i) {
if (!tensor->src[i]) {
continue;
}
auto it = stream_mapping.find(tensor->src[i]);
if (it == stream_mapping.end()) {
continue;
}
if (it->second != stream) {
dependent_srcs = true;
break;
}
}
if (dependent_srcs || writes_overlap) {
break;
}
}
return !writes_overlap && !dependent_srcs;
}
~ggml_cuda_concurrent_event() {
if (fork_event != nullptr) {
CUDA_CHECK(cudaEventDestroy(fork_event));
}
for (cudaEvent_t e : join_events) {
if (e != nullptr) {
CUDA_CHECK(cudaEventDestroy(e));
}
}
}
};
struct ggml_cuda_stream_context {
std::unordered_map<const ggml_tensor *, ggml_cuda_concurrent_event> concurrent_events;
void reset() {
concurrent_events.clear();
}
};
struct ggml_backend_cuda_context {
int device;
@@ -1348,9 +1461,6 @@ struct ggml_backend_cuda_context {
cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {};
void * cublas_workspaces[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {};
size_t cublas_workspace_sizes[GGML_CUDA_MAX_DEVICES] = {};
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
cublasLtHandle_t cublaslt_handles[GGML_CUDA_MAX_DEVICES] = {};
#endif
int curr_stream_no = 0;
@@ -1362,20 +1472,61 @@ struct ggml_backend_cuda_context {
int64_t next_graph_eviction_time = 0;
int64_t last_graph_use_time = 0;
// defined in ggml-cuda.cu, which includes graph.cuh
ggml_cuda_graph * cuda_graph(const void * first_node_ptr, int64_t time_now = ggml_time_us());
ggml_cuda_graph * cuda_graph(const void * first_node_ptr, int64_t time_now = ggml_time_us()) {
// Give all cached graphs 10s to be reused after an idle period, without updating each graph.
if (!cuda_graphs.empty() && time_now - last_graph_use_time >= 10'000'000) {
next_graph_eviction_time = time_now + 10'000'000;
}
last_graph_use_time = time_now;
// Sweep every 5s during active use, evicting graphs unused for at least 10s.
if (time_now >= next_graph_eviction_time) {
next_graph_eviction_time = time_now + 5'000'000;
for (auto it = cuda_graphs.begin(); it != cuda_graphs.end(); ) {
if (it->first != first_node_ptr && time_now - it->second->last_used_time >= 10'000'000) {
it = cuda_graphs.erase(it);
} else {
++it;
}
}
}
auto it = cuda_graphs.find(first_node_ptr);
if (it == cuda_graphs.end()) {
it = cuda_graphs.emplace(first_node_ptr, std::make_unique<ggml_cuda_graph>()).first;
}
it->second->last_used_time = time_now;
return it->second.get();
}
// Check if any CUDA graph is enabled for this context (used by kernels that need to know
// if graphs are in use without having access to the specific graph key)
bool any_cuda_graph_enabled() const;
bool any_cuda_graph_enabled() const {
for (const auto & [key, graph] : cuda_graphs) {
if (graph && graph->is_enabled()) {
return true;
}
}
return false;
}
// Check if any CUDA graph has an instance for this context
bool any_cuda_graph_has_instance() const;
bool any_cuda_graph_has_instance() const {
for (const auto & [key, graph] : cuda_graphs) {
if (graph && !graph->instances.empty()) {
return true;
}
}
return false;
}
#endif // USE_CUDA_GRAPH
explicit ggml_backend_cuda_context(int device);
explicit ggml_backend_cuda_context(int device) :
device(device),
name(GGML_CUDA_NAME + std::to_string(device)) {
}
std::unique_ptr<ggml_cuda_stream_context> concurrent_stream_context;
ggml_cuda_stream_context concurrent_stream_context;
~ggml_backend_cuda_context();
@@ -1389,7 +1540,7 @@ struct ggml_backend_cuda_context {
cudaStream_t stream() { return stream(device, curr_stream_no); }
ggml_cuda_stream_context & stream_context() { return *concurrent_stream_context; }
ggml_cuda_stream_context & stream_context() { return concurrent_stream_context; }
cublasHandle_t cublas_handle() {
if (cublas_handles[device][curr_stream_no] == nullptr) {
@@ -1409,16 +1560,6 @@ struct ggml_backend_cuda_context {
return cublas_handles[device][curr_stream_no];
}
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
cublasLtHandle_t cublaslt_handle() {
if (cublaslt_handles[device] == nullptr) {
ggml_cuda_set_device(device);
CUBLAS_CHECK(cublasLtCreate(&cublaslt_handles[device]));
}
return cublaslt_handles[device];
}
#endif
// pool
std::unique_ptr<ggml_cuda_pool> pools[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS];
@@ -1503,8 +1644,60 @@ struct ggml_cuda_pdl_config {
};
// whether kernel can be launched with PDL on the current device (also honors GGML_CUDA_PDL=0), defined in ggml-cuda.cu
bool ggml_cuda_kernel_can_use_pdl(const void * kernel);
static bool ggml_cuda_kernel_can_use_pdl(const void * kernel) {
const int device = ggml_cuda_get_device();
struct cache_key {
int device;
const void * kernel;
bool operator==(const cache_key & other) const { return device == other.device && kernel == other.kernel; }
};
struct cache_key_hash {
// MurmurHash3 mixing function for better hash distribution (vs. just std::hash which in some implementations simply returns the identity)
static size_t hash_mix(size_t x) {
std::uint64_t y = x;
const std::uint64_t m = 0xe9846af9b1a615d;
y ^= y >> 32;
y *= m;
y ^= y >> 32;
y *= m;
y ^= y >> 28;
return static_cast<size_t>(y);
}
size_t operator()(const cache_key & key) const {
// Use a nonzero seed to avoid mapping all-zero keys to zero
size_t h = 42;
h = hash_mix(h + key.device);
h = hash_mix(h + reinterpret_cast<size_t>(key.kernel));
return h;
}
};
static std::mutex cache_mutex;
static std::unordered_map<cache_key, bool, cache_key_hash> cache;
const cache_key key = { device, kernel };
std::lock_guard<std::mutex> lock(cache_mutex);
const auto it = cache.find(key);
if (it != cache.end()) {
return it->second;
}
cudaFuncAttributes attr = {};
CUDA_CHECK(cudaFuncGetAttributes(&attr, kernel));
// PDL device-side primitives are emitted only for PTX versions >= 90.
// We have to guard on a loaded kernel's PTX version so a kernel forward-JIT'ed
// from pre-Hopper PTX to a Hopper-or-newer GPU does not opt into PDL.
const bool can_use_pdl = attr.ptxVersion >= 90;
cache.emplace(key, can_use_pdl);
return can_use_pdl;
}
#endif //defined(GGML_CUDA_USE_PDL)
@@ -1518,7 +1711,13 @@ bool ggml_cuda_kernel_can_use_pdl(const void * kernel);
template<typename Kernel, typename... Args>
static __inline__ void ggml_cuda_kernel_launch(Kernel kernel, const ggml_cuda_kernel_launch_params & launch_params, Args&&... args) {
#if defined(GGML_CUDA_USE_PDL)
if (ggml_cuda_kernel_can_use_pdl(reinterpret_cast<const void *>(kernel))) {
static const bool env_pdl_enabled = []() {
const char * env = getenv("GGML_CUDA_PDL");
return env == nullptr || std::atoi(env) != 0;
}();
if (env_pdl_enabled && ggml_cuda_kernel_can_use_pdl(reinterpret_cast<const void *>(kernel))) {
auto pdl_cfg = ggml_cuda_pdl_config(launch_params);
CUDA_CHECK(cudaLaunchKernelEx(&pdl_cfg.cfg, kernel, std::forward<Args>(args)... ));
-12
View File
@@ -541,8 +541,6 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
return dequantize_row_mxfp4_cuda;
case GGML_TYPE_NVFP4:
return dequantize_row_nvfp4_cuda;
case GGML_TYPE_F8_E4M3:
return convert_unary_cont_cuda<ggml_fp8_e4m3_t>;
case GGML_TYPE_F32:
return convert_unary_cont_cuda<float>;
case GGML_TYPE_F16:
@@ -603,8 +601,6 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
return dequantize_row_mxfp4_cuda;
case GGML_TYPE_NVFP4:
return dequantize_row_nvfp4_cuda;
case GGML_TYPE_F8_E4M3:
return convert_unary_cont_cuda<ggml_fp8_e4m3_t>;
case GGML_TYPE_F32:
return convert_unary_cont_cuda<float>;
case GGML_TYPE_BF16:
@@ -662,8 +658,6 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
return dequantize_row_mxfp4_cuda;
case GGML_TYPE_NVFP4:
return dequantize_row_nvfp4_cuda;
case GGML_TYPE_F8_E4M3:
return convert_unary_cont_cuda<ggml_fp8_e4m3_t>;
case GGML_TYPE_F16:
return convert_unary_cont_cuda<half>;
case GGML_TYPE_BF16:
@@ -691,8 +685,6 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
case GGML_TYPE_Q8_0:
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
case GGML_TYPE_F8_E4M3:
return convert_unary_cuda<ggml_fp8_e4m3_t>;
case GGML_TYPE_BF16:
return convert_unary_cuda<nv_bfloat16>;
default:
@@ -718,8 +710,6 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
case GGML_TYPE_Q8_0:
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
case GGML_TYPE_F8_E4M3:
return convert_unary_cuda<ggml_fp8_e4m3_t>;
case GGML_TYPE_F16:
return convert_unary_cuda<half, nv_bfloat16>;
default:
@@ -745,8 +735,6 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
case GGML_TYPE_Q8_0:
return dequantize_block_cuda<QK8_0, QR8_0, dequantize_q8_0>;
case GGML_TYPE_F8_E4M3:
return convert_unary_cuda<ggml_fp8_e4m3_t>;
case GGML_TYPE_BF16:
return convert_unary_cuda<nv_bfloat16, float>;
default:
-30
View File
@@ -35,36 +35,6 @@ template<typename dst_t, typename src_t>
__host__ __device__ inline dst_t ggml_cuda_cast(src_t x) {
if constexpr (std::is_same_v<dst_t, src_t>) {
return x;
} else if constexpr (std::is_same_v<dst_t, ggml_fp8_e4m3_t>) {
return { ggml_cuda_fp32_to_f8_e4m3(ggml_cuda_cast<float>(x)) };
} else if constexpr (std::is_same_v<src_t, ggml_fp8_e4m3_t>) {
#if defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
#if CUDART_VERSION >= 13040
if constexpr (std::is_same_v<dst_t, nv_bfloat16>) {
return static_cast<nv_bfloat16>(__nv_cvt_fp8_to_bf16raw(x.bits, __NV_E4M3));
}
#endif
if constexpr (std::is_same_v<dst_t, half>) {
return static_cast<half>(__nv_cvt_fp8_to_halfraw(x.bits, __NV_E4M3));
}
#endif
return ggml_cuda_cast<dst_t>(ggml_cuda_f8_e4m3_to_fp32(x.bits));
#if defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
} else if constexpr (std::is_same_v<src_t, __nv_fp8x2_e4m3> && std::is_same_v<dst_t, nv_bfloat162>) {
#if CUDART_VERSION >= 13040
return static_cast<nv_bfloat162>(__nv_cvt_fp8x2_to_bf162raw(x.__x, __NV_E4M3));
#else
return make_bfloat162(
ggml_cuda_cast<nv_bfloat16>(ggml_cuda_f8_e4m3_to_fp32(x.__x & 0xFF)),
ggml_cuda_cast<nv_bfloat16>(ggml_cuda_f8_e4m3_to_fp32(x.__x >> 8)));
#endif
} else if constexpr (std::is_same_v<src_t, __nv_fp8x2_e4m3> && std::is_same_v<dst_t, float2>) {
#if CUDART_VERSION >= 13040
return ggml_cuda_cast<float2>(ggml_cuda_cast<nv_bfloat162>(x));
#else
return __half22float2(static_cast<half2>(__nv_cvt_fp8x2_to_halfraw2(x.__x, __NV_E4M3)));
#endif
#endif
} else if constexpr(std::is_same_v<dst_t, nv_bfloat16>) {
return __float2bfloat16(float(x));
} else if constexpr(std::is_same_v<src_t, nv_bfloat16>) {
-35
View File
@@ -477,41 +477,6 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg
} else if (ggml_cuda_cpy_as_memcpy_2d(src0, src1, mc_width, mc_height, mc_spitch, mc_dpitch)) {
CUDA_CHECK(cudaMemcpy2DAsync(src1_ddc, mc_dpitch, src0_ddc, mc_spitch,
mc_width, mc_height, cudaMemcpyDeviceToDevice, main_stream));
} else if (src0->type == GGML_TYPE_F8_E4M3 && src1->type == GGML_TYPE_F8_E4M3) {
ggml_cpy_scalar_cuda<uint8_t, uint8_t>
(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F8_E4M3) {
if (contiguous_srcs) {
ggml_cpy_scalar_contiguous_cuda<float, ggml_fp8_e4m3_t>
(src0_ddc, src1_ddc, ne, main_stream);
} else {
ggml_cpy_scalar_cuda<float, ggml_fp8_e4m3_t>
(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
}
} else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F8_E4M3) {
if (contiguous_srcs) {
ggml_cpy_scalar_contiguous_cuda<half, ggml_fp8_e4m3_t>
(src0_ddc, src1_ddc, ne, main_stream);
} else {
ggml_cpy_scalar_cuda<half, ggml_fp8_e4m3_t>
(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
}
} else if (src0->type == GGML_TYPE_BF16 && src1->type == GGML_TYPE_F8_E4M3) {
if (contiguous_srcs) {
ggml_cpy_scalar_contiguous_cuda<nv_bfloat16, ggml_fp8_e4m3_t>
(src0_ddc, src1_ddc, ne, main_stream);
} else {
ggml_cpy_scalar_cuda<nv_bfloat16, ggml_fp8_e4m3_t>
(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
}
} else if (src0->type == GGML_TYPE_F8_E4M3 && src1->type == GGML_TYPE_F32) {
if (contiguous_srcs) {
ggml_cpy_scalar_contiguous_cuda<ggml_fp8_e4m3_t, float>
(src0_ddc, src1_ddc, ne, main_stream);
} else {
ggml_cpy_scalar_cuda<ggml_fp8_e4m3_t, float>
(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
}
} else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) {
if (can_be_transposed) {
ggml_cpy_scalar_cuda<float, float, true>
-45
View File
@@ -146,34 +146,6 @@ static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_bf16(
return sum;
}
template <int D, int nthreads>
static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_f8_e4m3(
const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
const uint8_t * K_f8 = (const uint8_t *) K_c;
GGML_UNUSED(Q_q8);
GGML_UNUSED(Q_ds_v);
float sum = 0.0f;
#pragma unroll
for (int k_KQ_0 = 0; k_KQ_0 < D; k_KQ_0 += nthreads*4) {
uint8_t tmp[4];
ggml_cuda_memcpy_1<sizeof(tmp)>(tmp, K_f8 + k_KQ_0 + (threadIdx.x % nthreads)*4);
#pragma unroll
for (int k_KQ_1 = 0; k_KQ_1 < 4; ++k_KQ_1) {
const float k = ggml_cuda_f8_e4m3_to_fp32(tmp[k_KQ_1]);
#ifdef V_DOT2_F32_F16_AVAILABLE
sum += k * __half2float(((const half *) Q_v)[k_KQ_0/nthreads + k_KQ_1]);
#else
sum += k * ((const float *) Q_v)[k_KQ_0/nthreads + k_KQ_1];
#endif // V_DOT2_F32_F16_AVAILABLE
}
}
return sum;
}
template<int D, int nthreads>
static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_q4_0(
const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
@@ -434,19 +406,6 @@ static __device__ __forceinline__ void dequantize_V_bf16(const void * __restrict
}
}
template <typename T, int ne>
static __device__ __forceinline__ void dequantize_V_f8_e4m3(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
static_assert(ne == 2 || ne == 4, "bad ne");
static_assert(std::is_same_v<T, half> || std::is_same_v<T, float>, "unsupported type");
ggml_fp8_e4m3_t tmp[ne];
ggml_cuda_memcpy_1<sizeof(tmp)>(tmp, (const ggml_fp8_e4m3_t *) vx + i0);
#pragma unroll
for (int l = 0; l < ne; ++l) {
((T *) dst)[l] = ggml_cuda_cast<T>(tmp[l]);
}
}
template <typename T, int ne>
static __device__ __forceinline__ void dequantize_V_q4_0(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
const block_q4_0 * x = (const block_q4_0 *) vx;
@@ -685,8 +644,6 @@ constexpr __device__ vec_dot_KQ_t get_vec_dot_KQ() {
return vec_dot_fattn_vec_KQ_q8_0<D, nthreads>;
} else if constexpr (type_K == GGML_TYPE_BF16) {
return vec_dot_fattn_vec_KQ_bf16<D, nthreads>;
} else if constexpr (type_K == GGML_TYPE_F8_E4M3) {
return vec_dot_fattn_vec_KQ_f8_e4m3<D, nthreads>;
} else {
static_assert(type_K == -1, "bad type");
return nullptr;
@@ -709,8 +666,6 @@ constexpr __device__ dequantize_V_t get_dequantize_V() {
return dequantize_V_q8_0<T, ne>;
} else if constexpr (type_V == GGML_TYPE_BF16) {
return dequantize_V_bf16<float, ne>;
} else if constexpr (type_V == GGML_TYPE_F8_E4M3) {
return dequantize_V_f8_e4m3<T, ne>;
} else {
static_assert(type_V == -1, "bad type");
return nullptr;
+36 -152
View File
@@ -36,21 +36,11 @@ struct fattn_mma_config {
} \
static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_config_ampere(const int DKQ, const int DV, const int ncols) {
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 8, 128, 2, 128, 24, 24, 24, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 16, 128, 2, 64, 24, 24, 24, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 32, 128, 2, 64, 24, 24, 24, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 64, 128, 2, 64, 24, 24, 24, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 8, 128, 2, 128, 32, 32, 32, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 16, 128, 2, 64, 32, 32, 32, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 32, 128, 2, 64, 32, 32, 32, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 64, 128, 2, 64, 32, 32, 32, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 8, 128, 2, 128, 40, 40, 40, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 16, 128, 2, 64, 40, 40, 40, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 32, 128, 2, 64, 40, 40, 40, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 64, 128, 2, 64, 40, 40, 40, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 8, 128, 2, 128, 40, 40, 40, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 16, 128, 2, 64, 40, 40, 40, 2, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 32, 128, 2, 64, 40, 40, 40, 2, true);
@@ -133,29 +123,16 @@ static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_co
GGML_CUDA_FATTN_MMA_CONFIG_CASE(576, 512, 32, 128, 2, 32, 160, 128, 64, 1, false);
GGML_CUDA_FATTN_MMA_CONFIG_CASE(576, 512, 64, 256, 1, 32, 160, 128, 64, 1, false);
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 32, 128, 2, 32, 128, 128, 64, 1, false);
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 64, 128, 2, 32, 128, 64, 128, 2, false);
// TODO tune specifically for Volta
return ggml_cuda_fattn_mma_get_config_ampere(DKQ, DV, ncols);
}
static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_config_rdna(const int DKQ, const int DV, const int ncols) {
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 8, 128, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 16, 128, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 32, 128, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 64, 128, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 8, 128, 2, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 16, 128, 2, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 32, 128, 2, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 64, 128, 2, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 8, 64, 2, 32, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 16, 64, 2, 32, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 32, 128, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 64, 128, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 8, 64, 2, 32, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 16, 64, 2, 32, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 32, 128, 2, 64, 40, 40, 40, 1, true);
@@ -203,21 +180,11 @@ static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_co
}
static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_config_cdna(const int DKQ, const int DV, const int ncols) {
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 8, 256, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 16, 256, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 32, 256, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 40, 40, 64, 256, 2, 64, 24, 24, 24, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 8, 128, 1, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 16, 256, 2, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 32, 256, 2, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 64, 64, 64, 256, 3, 64, 32, 32, 32, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 8, 256, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 16, 256, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 32, 256, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 72, 72, 64, 256, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 8, 256, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 16, 256, 2, 64, 40, 40, 40, 1, true);
GGML_CUDA_FATTN_MMA_CONFIG_CASE( 80, 80, 32, 256, 2, 64, 40, 40, 40, 1, true);
@@ -652,9 +619,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
constexpr bool Q_in_reg = ggml_cuda_fattn_mma_get_Q_in_reg (DKQ, DV, ncols);
constexpr int nstages = ggml_cuda_fattn_mma_get_nstages (DKQ, DV, ncols1, ncols2, use_sparse);
constexpr int DKQp = GGML_PAD(DKQ, 16);
constexpr int DVp = GGML_PAD(DV, 16);
constexpr bool swz = ggml_cuda_fattn_mma_get_swizzled(DKQ, DV, ncols);
constexpr int stride_tile_K = swz ? nbatch_K2 : nbatch_K2 + 4;
constexpr int stride_tile_V = V_is_K_view ? stride_tile_K : (swz ? nbatch_V2 : nbatch_V2 + 4);
@@ -671,12 +635,12 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
if constexpr (nstages > 1) {
static_assert(!oob_check, "OOB check incompatible with multi-stage pipeline");
static_assert(!V_is_K_view, "K data reuse not implemented multi-stage loading");
static_assert(nbatch_K2 == DKQp/2, "batching not implemented for multi stage loading");
static_assert(nbatch_K2 == DKQ/2, "batching not implemented for multi stage loading");
constexpr bool use_cp_async = true;
cp_async_wait_all();
__syncthreads();
flash_attn_ext_f16_load_tile<stride_tile_V, nwarps, nbatch_fa, use_cp_async, oob_check, use_sparse>
(V_h2, tile_V, min(nbatch_V2, DV/2), stride_V, k_VKQ_0, k_VKQ_sup, indices);
(V_h2, tile_V, nbatch_V2, stride_V, k_VKQ_0, k_VKQ_sup, indices);
} else {
// the sparse mask values are gathered per element, always load them synchronously
constexpr bool use_cp_async = nstages == 1 && !use_sparse;
@@ -689,11 +653,11 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
// For MLA K and V have the same data.
// Therefore, iterate over K in reverse and later re-use the data if possible.
#pragma unroll
for (int k0_start = (DKQp/2-1) - (DKQp/2-1) % nbatch_K2; k0_start >= 0; k0_start -= nbatch_K2) {
const int k0_stop = k0_start + nbatch_K2 < DKQp/2 ? k0_start + nbatch_K2 : DKQp/2;
for (int k0_start = (DKQ/2-1) - (DKQ/2-1) % nbatch_K2; k0_start >= 0; k0_start -= nbatch_K2) {
const int k0_stop = k0_start + nbatch_K2 < DKQ/2 ? k0_start + nbatch_K2 : DKQ/2;
if constexpr (nstages <= 1) {
const int k0_diff = min(k0_stop, DKQ/2) - k0_start; // Only load the real head size from global memory.
const int k0_diff = k0_stop - k0_start;
constexpr bool use_cp_async = nstages == 1;
flash_attn_ext_f16_load_tile<stride_tile_K, nwarps, nbatch_fa, use_cp_async, oob_check, use_sparse>
(K_h2 + k0_start, tile_K, k0_diff, stride_K, k_VKQ_0, k_VKQ_sup, indices);
@@ -727,7 +691,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
}
}
} else {
constexpr int stride_tile_Q = DKQp/2 + 4;
constexpr int stride_tile_Q = DKQ/2 + 4;
#pragma unroll
for (int k_KQ_0 = k0_start; k_KQ_0 < k0_stop; k_KQ_0 += T_A_KQ::J) {
load_ldmatrix(Q_B[0], tile_Q + (threadIdx.y / np)*(T_B_KQ::I*stride_tile_Q) + k_KQ_0, stride_tile_Q);
@@ -954,7 +918,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
if constexpr (cols_per_warp == 8) {
const half2 KQ_max_scale_h2 = make_half2(KQ_max_scale[0], KQ_max_scale[cols_per_thread - 1]);
#pragma unroll
for (int i = 0; i < DVp/T_C_VKQ::I; ++i) {
for (int i = 0; i < DV/T_C_VKQ::I; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale_h2;
@@ -965,7 +929,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
for (int col = 0; col < cols_per_thread; ++col) {
const half2 KQ_max_scale_h2 = make_half2(KQ_max_scale[col], KQ_max_scale[col]);
#pragma unroll
for (int i = 0; i < (DVp/2)/T_C_VKQ::J; ++i) {
for (int i = 0; i < (DV/2)/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l0 = 0; l0 < T_C_VKQ::ne; l0 += 2) {
VKQ_C[i].x[l0 + col] *= KQ_max_scale_h2;
@@ -977,7 +941,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
if constexpr (std::is_same_v<decltype(T_C_VKQ::x), half2[T_C_VKQ::ne]>) {
const half2 KQ_max_scale_h2 = make_half2(KQ_max_scale[0], KQ_max_scale[0]);
#pragma unroll
for (int i = 0; i < (DVp/2)/T_C_VKQ::J; ++i) {
for (int i = 0; i < (DV/2)/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale_h2;
@@ -986,7 +950,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
} else {
static_assert(std::is_same_v<decltype(T_C_VKQ::x), float[T_C_VKQ::ne]>, "bad VKQ type");
#pragma unroll
for (int i = 0; i < DVp/T_C_VKQ::J; ++i) {
for (int i = 0; i < DV/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale[0];
@@ -997,7 +961,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
const half2 KQ_max_scale_h2 = make_half2(
KQ_max_scale[(threadIdx.x / 2) % 2], KQ_max_scale[(threadIdx.x / 2) % 2]);
#pragma unroll
for (int i = 0; i < (DVp/2)/T_C_VKQ::J; ++i) {
for (int i = 0; i < (DV/2)/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale_h2;
@@ -1033,19 +997,19 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
(mask_h, tile_mask, stride_mask, k_VKQ_0 + nbatch_fa, k_VKQ_sup, jt*ncols1, ne01, nullptr);
}
flash_attn_ext_f16_load_tile<stride_tile_K, nwarps, nbatch_fa, use_cp_async, oob_check, use_sparse>
(K_h2, tile_K, min(nbatch_K2, DKQ/2), stride_K, k_VKQ_0 + nbatch_fa, k_VKQ_sup, indices);
(K_h2, tile_K, nbatch_K2, stride_K, k_VKQ_0 + nbatch_fa, k_VKQ_sup, indices);
}
}
// Calculate VKQ tile, need to use logical rather than physical elements for i0 due to transposition of V:
#pragma unroll
for (int i0_start = 0; i0_start < DVp; i0_start += 2*nbatch_V2) {
static_assert(DVp % (2*nbatch_V2) == 0, "bad loop size");
for (int i0_start = 0; i0_start < DV; i0_start += 2*nbatch_V2) {
static_assert(DV % (2*nbatch_V2) == 0, "bad loop size");
const int i0_stop = i0_start + 2*nbatch_V2;
if constexpr (nstages <= 1) {
const int i0_diff = min(i0_stop, DV) - i0_start; // Only load the real head size from global memory.
const int i0_diff = i0_stop - i0_start;
if (!V_is_K_view || i0_stop > 2*nbatch_K2) {
constexpr bool use_cp_async = nstages == 1;
flash_attn_ext_f16_load_tile<stride_tile_V, nwarps, nbatch_fa, use_cp_async, oob_check, use_sparse>
@@ -1182,22 +1146,6 @@ template<int ncols> struct mma_tile_sizes<112, ncols> {
using T_B_VKQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // column-major
using T_C_VKQ = tile<16, 16, float, DATA_LAYOUT_I_MAJOR>; // column-major
};
template<int ncols> struct mma_tile_sizes<40, ncols> {
using T_A_KQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // row-major
using T_B_KQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // column-major
using T_C_KQ = tile<16, 16, float, DATA_LAYOUT_I_MAJOR>; // column-major
using T_A_VKQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // row-major
using T_B_VKQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // column-major
using T_C_VKQ = tile<16, 16, float, DATA_LAYOUT_I_MAJOR>; // column-major
};
template<int ncols> struct mma_tile_sizes<72, ncols> {
using T_A_KQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // row-major
using T_B_KQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // column-major
using T_C_KQ = tile<16, 16, float, DATA_LAYOUT_I_MAJOR>; // column-major
using T_A_VKQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // row-major
using T_B_VKQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED>; // column-major
using T_C_VKQ = tile<16, 16, float, DATA_LAYOUT_I_MAJOR>; // column-major
};
#else
template<int DKQ, int ncols> struct mma_tile_sizes {
using T_A_KQ = tile<16, 8, half2, DATA_LAYOUT_I_MAJOR>; // row-major
@@ -1223,22 +1171,6 @@ template<int ncols> struct mma_tile_sizes<112, ncols> {
using T_B_VKQ = tile<16, 8, half2>; // column-major
using T_C_VKQ = tile<16, 8, half2>; // column-major
};
template<int ncols> struct mma_tile_sizes<40, ncols> {
using T_A_KQ = tile<16, 8, half2>; // row-major
using T_B_KQ = tile<16, 8, half2>; // column-major
using T_C_KQ = tile<16, 16, float>; // column-major
using T_A_VKQ = tile<16, 8, half2>; // row-major
using T_B_VKQ = tile<16, 8, half2>; // column-major
using T_C_VKQ = tile<16, 8, half2>; // column-major
};
template<int ncols> struct mma_tile_sizes<72, ncols> {
using T_A_KQ = tile<16, 8, half2>; // row-major
using T_B_KQ = tile<16, 8, half2>; // column-major
using T_C_KQ = tile<16, 16, float>; // column-major
using T_A_VKQ = tile<16, 8, half2>; // row-major
using T_B_VKQ = tile<16, 8, half2>; // column-major
using T_C_VKQ = tile<16, 8, half2>; // column-major
};
#endif // RDNA3
#elif defined(AMD_MFMA_AVAILABLE)
template<int DKQ, int ncols> struct mma_tile_sizes {
@@ -1308,9 +1240,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
constexpr bool Q_in_reg = ggml_cuda_fattn_mma_get_Q_in_reg (DKQ, DV, ncols);
constexpr int nstages = ggml_cuda_fattn_mma_get_nstages (DKQ, DV, ncols1, ncols2, use_sparse);
constexpr int DKQp = GGML_PAD(DKQ, 16);
constexpr int DVp = GGML_PAD(DV, 16);
if (cols_per_warp > ncols) {
NO_DEVICE_CODE;
return;
@@ -1319,7 +1248,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
static_assert(nwarps * (cols_per_warp/ncols2) % ncols1 == 0, "bad nwarps");
constexpr bool swz = ggml_cuda_fattn_mma_get_swizzled(DKQ, DV, ncols);
constexpr int stride_tile_Q = DKQp/2 + 4;
constexpr int stride_tile_Q = DKQ/2 + 4;
constexpr int stride_tile_K = swz ? nbatch_K2 : nbatch_K2 + 4;
constexpr int stride_tile_V = V_is_K_view ? stride_tile_K : (swz ? nbatch_V2 : nbatch_V2 + 4);
constexpr int stride_tile_KV_max = stride_tile_K > stride_tile_V ? stride_tile_K : stride_tile_V;
@@ -1329,17 +1258,17 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
half2 * tile_V = nstages > 1 ? tile_K + nbatch_fa * stride_tile_K : tile_K;
half * tile_mask = (half *) (nstages > 1 ? tile_V + nbatch_fa * stride_tile_V : tile_V + nbatch_fa * stride_tile_KV_max);
T_B_KQ Q_B[(Q_in_reg ? DKQp/(2*T_B_KQ::J) : 1)];
T_B_KQ Q_B[(Q_in_reg ? DKQ/(2*T_B_KQ::J) : 1)];
#if defined(TURING_MMA_AVAILABLE)
T_C_VKQ VKQ_C[cols_per_warp == 8 ? DVp/T_C_VKQ::I : DVp/(2*T_C_VKQ::J)];
T_C_VKQ VKQ_C[cols_per_warp == 8 ? DV/T_C_VKQ::I : DV/(2*T_C_VKQ::J)];
#elif defined(AMD_WMMA_256B_AVAILABLE)
T_C_VKQ VKQ_C[DVp % 32 != 0 ? DVp/T_C_VKQ::J : DVp/(2*T_C_VKQ::J)];
T_C_VKQ VKQ_C[DV % 32 != 0 ? DV/T_C_VKQ::J : DV/(2*T_C_VKQ::J)];
#elif defined(AMD_MFMA_AVAILABLE)
T_C_VKQ VKQ_C[ DVp/T_C_VKQ::J];
T_C_VKQ VKQ_C[ DV/T_C_VKQ::J];
#elif defined(AMD_WMMA_AVAILABLE)
T_C_VKQ VKQ_C[ DVp/(2*T_C_VKQ::J)];
T_C_VKQ VKQ_C[ DV/(2*T_C_VKQ::J)];
#else // Volta
T_C_VKQ VKQ_C[ DVp/(2*T_C_VKQ::J)];
T_C_VKQ VKQ_C[ DV/(2*T_C_VKQ::J)];
#endif // defined(TURING_MMA_AVAILABLE)
float KQ_rowsum[cols_per_thread] = {0.0f};
@@ -1355,8 +1284,8 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
const half2 scale_h2 = make_half2(scale, scale);
#pragma unroll
for (int stride_k : {warp_size, warp_size/2, warp_size/4, warp_size/8}) {
const int k0_start = stride_k == warp_size ? 0 : DKQp/2 - (DKQp/2) % (2*stride_k);
const int k0_stop = DKQp/2 - (DKQp/2) % (1*stride_k);
const int k0_start = stride_k == warp_size ? 0 : DKQ/2 - (DKQ/2) % (2*stride_k);
const int k0_stop = DKQ/2 - (DKQ/2) % (1*stride_k);
const int stride_jc = warp_size / stride_k;
if (k0_start == k0_stop) {
@@ -1379,8 +1308,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
for (int k0 = k0_start; k0 < k0_stop; k0 += stride_k) {
const int k = k0 + (stride_k == warp_size ? threadIdx.x : threadIdx.x % stride_k);
const float2 tmp = DKQp == DKQ || k < DKQ/2 ?
Q_f2[(jt*ncols1 + j)*stride_Q1 + c*stride_Q2 + k] : make_float2(0.0f, 0.0f);
const float2 tmp = Q_f2[(jt*ncols1 + j)*stride_Q1 + c*stride_Q2 + k];
tile_Q[jc*stride_tile_Q + k] = scale_h2 * make_half2(tmp.x, tmp.y);
}
} else {
@@ -1400,50 +1328,19 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
const int j0 = (threadIdx.y / np) * cols_per_warp;
#pragma unroll
for (int k0 = 0; k0 < DKQp/2; k0 += T_B_KQ::J) {
for (int k0 = 0; k0 < DKQ/2; k0 += T_B_KQ::J) {
load_ldmatrix(Q_B[k0/T_B_KQ::J], tile_Q + j0*stride_tile_Q + k0, stride_tile_Q);
}
}
__syncthreads();
// Zero the pad columns of the K/V tiles once, the tile loads only write the first DKQ/2 (DV/2) half2 of a row.
// Done after Q_B is loaded because tile_K may alias tile_Q.
if constexpr (DKQp != DKQ || DVp != DV) {
static_assert(!swz, "pad zero-fill assumes unswizzled tiles");
static_assert(nbatch_K2 == DKQp/2 && nbatch_V2 == DVp/2, "K/V tile batching not implemented for padded head sizes");
static_assert(DKQ == DV, "pad zero-fill for DKQ != DV requires separate handling of the aliased V tile");
const half2 zero[4] = {{0.0f, 0.0f}, {0.0f, 0.0f}, {0.0f, 0.0f}, {0.0f, 0.0f}};
const int tid = threadIdx.y*warp_size + threadIdx.x;
#pragma unroll
for (int i0 = 0; i0 < nbatch_fa; i0 += nwarps*warp_size) {
const int i = i0 + tid;
if (i0 + nwarps*warp_size > nbatch_fa && i >= nbatch_fa) {
break;
}
#pragma unroll
for (int k = DKQ/2; k < DKQp/2; k += 4) {
ggml_cuda_memcpy_1<16>(tile_K + i*stride_tile_K + k, zero);
}
if constexpr (nstages > 1) { // tile_V is only a separate buffer with multi-stage loading.
#pragma unroll
for (int k = DV/2; k < DVp/2; k += 4) {
ggml_cuda_memcpy_1<16>(tile_V + i*stride_tile_V + k, zero);
}
}
}
__syncthreads();
}
int kb0 = kb0_start;
// Preload mask and K data for first iteration when using cp_async with multiple stages:
if constexpr (nstages > 1) {
static_assert(!use_sparse, "sparse gather not implemented for multi-stage loading");
static_assert(nbatch_K2 == DKQp/2, "batching not implemented for multi-stage pipeline");
static_assert(nbatch_K2 == DKQ/2, "batching not implemented for multi-stage pipeline");
constexpr bool use_cp_async = true;
constexpr bool oob_check = false;
constexpr int k_VKQ_sup = nbatch_fa;
@@ -1452,7 +1349,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
(mask_h, tile_mask, stride_mask, kb0*nbatch_fa, k_VKQ_sup, jt*ncols1, ne01, nullptr);
}
flash_attn_ext_f16_load_tile<stride_tile_K, nwarps, nbatch_fa, use_cp_async, oob_check, use_sparse>
(K_h2, tile_K, min(nbatch_K2, DKQ/2), stride_K, kb0*nbatch_fa, k_VKQ_sup, indices);
(K_h2, tile_K, nbatch_K2, stride_K, kb0*nbatch_fa, k_VKQ_sup, indices);
}
// kb0_start is always < kb0_stop so the last iter can be executed unconditionally.
@@ -1557,7 +1454,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
if constexpr (cols_per_warp == 8) {
const half2 KQ_max_scale_h2 = make_half2(KQ_max_scale[0], KQ_max_scale[cols_per_thread - 1]);
#pragma unroll
for (int i = 0; i < DVp/T_C_VKQ::I; ++i) {
for (int i = 0; i < DV/T_C_VKQ::I; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale_h2;
@@ -1568,7 +1465,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
for (int col = 0; col < cols_per_thread; ++col) {
const half2 KQ_max_scale_h2 = make_half2(KQ_max_scale[col], KQ_max_scale[col]);
#pragma unroll
for (int i = 0; i < (DVp/2)/T_C_VKQ::J; ++i) {
for (int i = 0; i < (DV/2)/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l0 = 0; l0 < T_C_VKQ::ne; l0 += 2) {
VKQ_C[i].x[l0 + col] *= KQ_max_scale_h2;
@@ -1580,7 +1477,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
if constexpr (std::is_same_v<decltype(T_C_VKQ::x), half2[T_C_VKQ::ne]>) {
const half2 KQ_max_scale_h2 = make_half2(KQ_max_scale[0], KQ_max_scale[0]);
#pragma unroll
for (int i = 0; i < (DVp/2)/T_C_VKQ::J; ++i) {
for (int i = 0; i < (DV/2)/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale_h2;
@@ -1589,7 +1486,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
} else {
static_assert(std::is_same_v<decltype(T_C_VKQ::x), float[T_C_VKQ::ne]>, "bad VKQ type");
#pragma unroll
for (int i = 0; i < DVp/T_C_VKQ::J; ++i) {
for (int i = 0; i < DV/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale[0];
@@ -1600,7 +1497,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
const int col = (threadIdx.x / 2) % 2;
const half2 KQ_max_scale_h2 = make_half2(KQ_max_scale[col], KQ_max_scale[col]);
#pragma unroll
for (int i = 0; i < (DVp/2)/T_C_VKQ::J; ++i) {
for (int i = 0; i < (DV/2)/T_C_VKQ::J; ++i) {
#pragma unroll
for (int l = 0; l < T_C_VKQ::ne; ++l) {
VKQ_C[i].x[l] *= KQ_max_scale_h2;
@@ -1614,7 +1511,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
// So also write VKQ accumulators to shared memory in column-major format if np == 1.
constexpr int tile_stride = nbatch_combine + 4;
static_assert((DVp/2) % nbatch_combine == 0, "bad nbatch_combine");
static_assert((DV/2) % nbatch_combine == 0, "bad nbatch_combine");
if constexpr (cols_per_warp == 8) {
const int jc_cwmo = (threadIdx.x % (2*T_C_VKQ::J)) / T_C_VKQ::J; // jc combine write meta offset
@@ -1758,7 +1655,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
}
#pragma unroll
for (int k00 = 0; k00 < DVp/2; k00 += nbatch_combine) {
for (int k00 = 0; k00 < DV/2; k00 += nbatch_combine) {
if constexpr (cols_per_warp == 8) {
static_assert(std::is_same_v<decltype(T_C_VKQ::x), half2[T_C_VKQ::ne]>, "bad VKQ type");
const int jc_cwd = threadIdx.y*T_B_KQ::I + T_B_KQ::get_i(-1); // jc combine write data
@@ -1857,10 +1754,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
for (int k0 = k0_start; k0 < k0_stop; k0 += stride_k) {
const int k = k0 + (stride_k == warp_size ? threadIdx.x : threadIdx.x % stride_k);
if (DVp != DV && k00 + k >= DV/2) {
continue;
}
float2 dstk_val = make_float2(0.0f, 0.0f);
#pragma unroll
for (int ip = 0; ip < np; ++ip) {
@@ -2141,7 +2034,6 @@ void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml
const int nwarps = nthreads / warp_size_host;
constexpr bool V_is_K_view = DKQ == 576; // Guaranteed by the kernel selection logic in fattn.cu
constexpr int DKQp = GGML_PAD(DKQ, 16);
// KV tile strides must match flash_attn_ext_f16_iter / _process_tile.
const bool swz = ggml_cuda_fattn_mma_get_swizzled(DKQ, DV, ncols, cc);
@@ -2149,7 +2041,7 @@ void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml
const int stride_tile_V = V_is_K_view ? stride_tile_K : (swz ? nbatch_V2 : nbatch_V2 + 4);
const size_t nbytes_shared_KV_1stage = nbatch_fa * std::max(stride_tile_K, stride_tile_V) * sizeof(half2);
const size_t nbytes_shared_KV_2stage = nbatch_fa * (stride_tile_K + stride_tile_V) * sizeof(half2);
const size_t nbytes_shared_Q = ncols * (DKQp/2 + 4) * sizeof(half2);
const size_t nbytes_shared_Q = ncols * (DKQ/2 + 4) * sizeof(half2);
const size_t nbytes_shared_mask = ncols1 * (nbatch_fa/2 + 4) * sizeof(half2);
const size_t nbytes_shared_combine = nwarps*cols_per_warp * (nbatch_combine + 4) * sizeof(half2);
@@ -2234,36 +2126,28 @@ void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml
extern DECL_FATTN_MMA_F16_CASE(DKQ, DV, (ncols)/ 8, 8); \
extern DECL_FATTN_MMA_F16_CASE(DKQ, DV, (ncols)/16, 16); \
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 40, 40, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 64, 64, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 72, 72, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 80, 80, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 96, 96, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(112, 112, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(128, 128, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(256, 256, 8)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 40, 40, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 64, 64, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 72, 72, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 80, 80, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 96, 96, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(112, 112, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(128, 128, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(256, 256, 16)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 40, 40, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 64, 64, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 72, 72, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 80, 80, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 96, 96, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(112, 112, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(128, 128, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(256, 256, 32)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 40, 40, 64)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 64, 64, 64)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 72, 72, 64)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 80, 80, 64)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2( 96, 96, 64)
DECL_FATTN_MMA_F16_CASE_ALL_NCOLS2(112, 112, 64)
+1 -26
View File
@@ -96,7 +96,7 @@ static __global__ void flash_attn_ext_vec(
constexpr int V_cols_per_iter = WARP_SIZE / nthreads_V;
constexpr vec_dot_KQ_t vec_dot_KQ = get_vec_dot_KQ<type_K, D, nthreads_KQ>();
constexpr bool Q_q8_1 = type_K != GGML_TYPE_F16 && type_K != GGML_TYPE_BF16 && type_K != GGML_TYPE_F8_E4M3;
constexpr bool Q_q8_1 = type_K != GGML_TYPE_F16 && type_K != GGML_TYPE_BF16;
#ifdef V_DOT2_F32_F16_AVAILABLE
constexpr dequantize_V_t dequantize_V = get_dequantize_V<type_V, half, V_rows_per_thread>();
#else
@@ -203,27 +203,6 @@ static __global__ void flash_attn_ext_vec(
}
__syncthreads();
} else if constexpr (type_K == GGML_TYPE_F8_E4M3) {
#pragma unroll
for (int j = 0; j < ncols; ++j) {
const float * Q_j = (const float *) (Q + j*nb01);
#pragma unroll
for (int i0 = 0; i0 < D; i0 += nthreads_KQ*4) {
const int i = i0 + (threadIdx.x % nthreads_KQ)*4;
float4 tmp = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
if (ncols == 1 || ic0 + j < int(ne01.z)) {
ggml_cuda_memcpy_1<sizeof(tmp)>(&tmp, Q_j + i);
}
#ifdef V_DOT2_F32_F16_AVAILABLE
const half2 scale_h2 = make_half2(scale, scale);
Q_reg[j][i0/(2*nthreads_KQ) + 0] = __hmul2(make_half2(tmp.x, tmp.y), scale_h2);
Q_reg[j][i0/(2*nthreads_KQ) + 1] = __hmul2(make_half2(tmp.z, tmp.w), scale_h2);
#else
Q_reg[j][i0/(2*nthreads_KQ) + 0] = make_float2(tmp.x*scale, tmp.y*scale);
Q_reg[j][i0/(2*nthreads_KQ) + 1] = make_float2(tmp.z*scale, tmp.w*scale);
#endif
}
}
} else {
#ifdef V_DOT2_F32_F16_AVAILABLE
const half2 scale_h2 = make_half2(scale, scale);
@@ -608,7 +587,6 @@ void ggml_cuda_flash_attn_ext_vec_case(ggml_backend_cuda_context & ctx, ggml_ten
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q5_1); \
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q8_0); \
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_BF16); \
extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_F8_E4M3); \
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_F16)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q4_0)
@@ -617,7 +595,6 @@ EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q5_0)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q5_1)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q8_0)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_BF16)
EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_F8_E4M3)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_F16)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q4_0)
@@ -626,7 +603,6 @@ EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q5_0)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q5_1)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q8_0)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_BF16)
EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_F8_E4M3)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_F16)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q4_0)
@@ -635,4 +611,3 @@ EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q5_0)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q5_1)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q8_0)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_BF16)
EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_F8_E4M3)
+1 -26
View File
@@ -314,18 +314,10 @@ static void ggml_cuda_flash_attn_ext_mma_f16(ggml_backend_cuda_context & ctx, gg
const ggml_tensor * mask = dst->src[3];
switch (Q->ne[0]) {
case 40:
GGML_ASSERT(V->ne[0] == 40);
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols2< 40, 40>(ctx, dst);
break;
case 64:
GGML_ASSERT(V->ne[0] == 64);
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols2< 64, 64>(ctx, dst);
break;
case 72:
GGML_ASSERT(V->ne[0] == 72);
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols2< 72, 72>(ctx, dst);
break;
case 80:
GGML_ASSERT(V->ne[0] == 80);
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols2< 80, 80>(ctx, dst);
@@ -470,7 +462,6 @@ static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, co
FATTN_VEC_CASES_ALL_D(Q5_1, F16)
FATTN_VEC_CASES_ALL_D(Q8_0, F16)
FATTN_VEC_CASES_ALL_D(BF16, F16)
FATTN_VEC_CASES_ALL_D(F8_E4M3, F16)
FATTN_VEC_CASES_ALL_D(F16, Q4_0)
FATTN_VEC_CASES_ALL_D(Q4_0, Q4_0)
@@ -479,7 +470,6 @@ static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, co
FATTN_VEC_CASES_ALL_D(Q5_1, Q4_0)
FATTN_VEC_CASES_ALL_D(Q8_0, Q4_0)
FATTN_VEC_CASES_ALL_D(BF16, Q4_0)
FATTN_VEC_CASES_ALL_D(F8_E4M3, Q4_0)
FATTN_VEC_CASES_ALL_D(F16, Q4_1)
FATTN_VEC_CASES_ALL_D(Q4_0, Q4_1)
@@ -488,7 +478,6 @@ static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, co
FATTN_VEC_CASES_ALL_D(Q5_1, Q4_1)
FATTN_VEC_CASES_ALL_D(Q8_0, Q4_1)
FATTN_VEC_CASES_ALL_D(BF16, Q4_1)
FATTN_VEC_CASES_ALL_D(F8_E4M3, Q4_1)
FATTN_VEC_CASES_ALL_D(F16, Q5_0)
FATTN_VEC_CASES_ALL_D(Q4_0, Q5_0)
@@ -497,7 +486,6 @@ static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, co
FATTN_VEC_CASES_ALL_D(Q5_1, Q5_0)
FATTN_VEC_CASES_ALL_D(Q8_0, Q5_0)
FATTN_VEC_CASES_ALL_D(BF16, Q5_0)
FATTN_VEC_CASES_ALL_D(F8_E4M3, Q5_0)
FATTN_VEC_CASES_ALL_D(F16, Q5_1)
FATTN_VEC_CASES_ALL_D(Q4_0, Q5_1)
@@ -506,7 +494,6 @@ static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, co
FATTN_VEC_CASES_ALL_D(Q5_1, Q5_1)
FATTN_VEC_CASES_ALL_D(Q8_0, Q5_1)
FATTN_VEC_CASES_ALL_D(BF16, Q5_1)
FATTN_VEC_CASES_ALL_D(F8_E4M3, Q5_1)
FATTN_VEC_CASES_ALL_D(F16, Q8_0)
FATTN_VEC_CASES_ALL_D(Q4_0, Q8_0)
@@ -515,7 +502,6 @@ static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, co
FATTN_VEC_CASES_ALL_D(Q5_1, Q8_0)
FATTN_VEC_CASES_ALL_D(Q8_0, Q8_0)
FATTN_VEC_CASES_ALL_D(BF16, Q8_0)
FATTN_VEC_CASES_ALL_D(F8_E4M3, Q8_0)
FATTN_VEC_CASES_ALL_D(F16, BF16)
FATTN_VEC_CASES_ALL_D(Q4_0, BF16)
@@ -524,16 +510,6 @@ static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, co
FATTN_VEC_CASES_ALL_D(Q5_1, BF16)
FATTN_VEC_CASES_ALL_D(Q8_0, BF16)
FATTN_VEC_CASES_ALL_D(BF16, BF16)
FATTN_VEC_CASES_ALL_D(F8_E4M3, BF16)
FATTN_VEC_CASES_ALL_D(F16, F8_E4M3)
FATTN_VEC_CASES_ALL_D(Q4_0, F8_E4M3)
FATTN_VEC_CASES_ALL_D(Q4_1, F8_E4M3)
FATTN_VEC_CASES_ALL_D(Q5_0, F8_E4M3)
FATTN_VEC_CASES_ALL_D(Q5_1, F8_E4M3)
FATTN_VEC_CASES_ALL_D(Q8_0, F8_E4M3)
FATTN_VEC_CASES_ALL_D(BF16, F8_E4M3)
FATTN_VEC_CASES_ALL_D(F8_E4M3, F8_E4M3)
return nullptr;
}
@@ -577,7 +553,6 @@ static bool ggml_cuda_fattn_kv_type_supported(const ggml_type type) {
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_F8_E4M3:
return true;
default:
return false;
@@ -691,7 +666,7 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
const bool can_use_vector_kernel = Q->ne[0] <= 256 && Q->ne[0] % 64 == 0 && Q->ne[0] != 192 && K->ne[1] % FATTN_KQ_STRIDE == 0;
// If Turing tensor cores are available, use them:
if (turing_mma_available(cc)) {
if (turing_mma_available(cc) && Q->ne[0] != 40 && Q->ne[0] != 72) {
if (can_use_vector_kernel) {
if (!ggml_is_quantized(K->type) && !ggml_is_quantized(V->type)) {
// the sparse gather exists only in the MMA kernel: (DKQ, DV, 1, 8) with GQA > 4
-206
View File
@@ -1,206 +0,0 @@
#include "fp8.cuh"
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
struct fp8_abs_src {
const float * x;
int64_t ne0;
int64_t ne1;
int64_t ne2;
int64_t s1;
int64_t s2;
int64_t s3;
__device__ float operator()(int64_t i) const {
const int64_t i0 = i % ne0;
i /= ne0;
const int64_t i1 = i % ne1;
i /= ne1;
const int64_t i2 = i % ne2;
const int64_t i3 = i / ne2;
const float value = fabsf(x[i0 + i1*s1 + i2*s2 + i3*s3]);
return isfinite(value) ? value : 448.0f;
}
};
static __global__ void fp8_amax_partials(fp8_abs_src src, int64_t ne, float * partials) {
float amax = 0.0f;
for (int64_t i = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; i < ne; i += (int64_t) blockDim.x*gridDim.x) {
amax = fmaxf(amax, src(i));
}
__shared__ float shared[WARP_SIZE];
amax = block_reduce<block_reduce_method::MAX, 256>(amax, shared);
if (threadIdx.x == 0) {
partials[blockIdx.x] = amax;
}
}
static __global__ void fp8_amax_final(const float * partials, int n, float * amax) {
float value = 0.0f;
for (int i = threadIdx.x; i < n; i += blockDim.x) {
value = fmaxf(value, partials[i]);
}
__shared__ float shared[WARP_SIZE];
value = block_reduce<block_reduce_method::MAX, 256>(value, shared);
if (threadIdx.x == 0) {
*amax = value;
}
}
static __global__ void quantize_fp8_e4m3(
const float * __restrict__ x, uint8_t * __restrict__ y, const float * __restrict__ amax,
float * __restrict__ scale, int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne, int64_t s1, int64_t s2, int64_t s3) {
const int64_t i = (int64_t) blockIdx.x*blockDim.x + threadIdx.x;
if (i >= ne) {
return;
}
int64_t tmp = i / ne0;
const int64_t i0 = i - tmp*ne0;
const int64_t i1 = tmp % ne1;
tmp /= ne1;
const int64_t i2 = tmp % ne2;
const int64_t i3 = tmp / ne2;
const float d = *amax > 0.0f ? *amax / 448.0f : 1.0f;
y[i] = ggml_cuda_fp32_to_f8_e4m3(x[i0 + i1*s1 + i2*s2 + i3*s3] / d);
if (i == 0) {
*scale = d;
}
}
static void fp8_set_batch(cublasLtMatrixLayout_t desc, int32_t count, int64_t stride) {
CUBLAS_CHECK(cublasLtMatrixLayoutSetAttribute(desc, CUBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &count, sizeof(count)));
CUBLAS_CHECK(cublasLtMatrixLayoutSetAttribute(desc, CUBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &stride, sizeof(stride)));
}
static void fp8_destroy_matmul(
cublasLtMatmulDesc_t op_desc, cublasLtMatrixLayout_t a_desc, cublasLtMatrixLayout_t b_desc,
cublasLtMatrixLayout_t d_desc) {
CUBLAS_CHECK(cublasLtMatrixLayoutDestroy(d_desc));
CUBLAS_CHECK(cublasLtMatrixLayoutDestroy(b_desc));
CUBLAS_CHECK(cublasLtMatrixLayoutDestroy(a_desc));
CUBLAS_CHECK(cublasLtMatmulDescDestroy(op_desc));
}
bool ggml_cuda_mul_mat_fp8(
ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
const ggml_tensor * weight_scale) {
const int cc = ggml_cuda_info().devices[ctx.device].cc;
if (!fp8_mma_hardware_available(cc) || src0->type != GGML_TYPE_F8_E4M3 || src1->type != GGML_TYPE_F32 ||
dst->type != GGML_TYPE_F32 || !ggml_is_contiguous(dst) || src0->ne[0] % 16 != 0 || src0->ne[1] % 16 != 0 ||
src0->nb[0] != sizeof(uint8_t) || src0->nb[1] != (size_t) src0->ne[0] || src1->nb[0] != sizeof(float) ||
!ggml_cuda_is_aligned(src0, 16) || !ggml_cuda_is_aligned(dst, 16)) {
return false;
}
GGML_ASSERT(!weight_scale || (weight_scale->type == GGML_TYPE_F32 && ggml_is_contiguous(weight_scale) && ggml_nelements(weight_scale) == 1));
GGML_TENSOR_BINARY_OP_LOCALS
GGML_ASSERT(ne10 == ne00);
GGML_ASSERT(ne0 == ne01);
GGML_ASSERT(ne12 % ne02 == 0);
GGML_ASSERT(ne13 % ne03 == 0);
cudaStream_t stream = ctx.stream();
const int64_t ne_src1 = ggml_nelements(src1);
ggml_cuda_pool_alloc<uint8_t> src1_fp8(ctx.pool(), ne_src1);
ggml_cuda_pool_alloc<float> src1_scale(ctx.pool(), 1);
ggml_cuda_pool_alloc<float> src1_amax(ctx.pool(), 1);
const fp8_abs_src abs_src = {
(const float *) src1->data, ne10, ne11, ne12,
(int64_t) (nb11 / sizeof(float)), (int64_t) (nb12 / sizeof(float)), (int64_t) (nb13 / sizeof(float))
};
const int reduce_blocks = std::min<int64_t>((ne_src1 + 255)/256, 1024);
ggml_cuda_pool_alloc<float> reduce_tmp(ctx.pool(), reduce_blocks);
fp8_amax_partials<<<reduce_blocks, 256, 0, stream>>>(abs_src, ne_src1, reduce_tmp.ptr);
fp8_amax_final<<<1, 256, 0, stream>>>(reduce_tmp.ptr, reduce_blocks, src1_amax.ptr);
quantize_fp8_e4m3<<<(ne_src1 + 255)/256, 256, 0, stream>>>(
(const float *) src1->data, src1_fp8.ptr, src1_amax.ptr, src1_scale.ptr, ne10, ne11, ne12, ne_src1,
nb11 / sizeof(float), nb12 / sizeof(float), nb13 / sizeof(float));
CUDA_CHECK(cudaGetLastError());
cublasLtMatmulDesc_t op_desc;
cublasLtMatrixLayout_t a_desc;
cublasLtMatrixLayout_t b_desc;
cublasLtMatrixLayout_t d_desc;
CUBLAS_CHECK(cublasLtMatmulDescCreate(&op_desc, CUBLAS_COMPUTE_32F, CUDA_R_32F));
const cublasOperation_t trans_a = CUBLAS_OP_T;
CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(op_desc, CUBLASLT_MATMUL_DESC_TRANSA, &trans_a, sizeof(trans_a)));
CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(
op_desc, CUBLASLT_MATMUL_DESC_B_SCALE_POINTER, &src1_scale.ptr, sizeof(src1_scale.ptr)));
if (weight_scale) {
const float * scale_ptr = (const float *) weight_scale->data;
CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(
op_desc, CUBLASLT_MATMUL_DESC_A_SCALE_POINTER, &scale_ptr, sizeof(scale_ptr)));
}
const int64_t r2 = ne12 / ne02;
const int64_t r3 = ne13 / ne03;
// src1 batches that share a src0 matrix are adjacent in the compact src1 copy and in dst, fold them into N
const int64_t n = ne11*(ne02 == 1 ? r2*r3 : r2);
// batch the remaining matmuls with a fixed src0 stride, or issue one batched call per i3 if there is none
int64_t batch_count = ne02;
int64_t stride_a = nb02;
int64_t n_calls = ne13;
if (ne02 == 1) {
batch_count = ne03;
stride_a = nb03;
n_calls = 1;
} else if (r3 == 1 && (ne03 == 1 || nb03 == ne02*nb02)) {
batch_count = ne02*ne03;
n_calls = 1;
}
CUBLAS_CHECK(cublasLtMatrixLayoutCreate(&a_desc, CUDA_R_8F_E4M3, ne00, ne01, ne00));
CUBLAS_CHECK(cublasLtMatrixLayoutCreate(&b_desc, CUDA_R_8F_E4M3, ne10, n, ne10));
CUBLAS_CHECK(cublasLtMatrixLayoutCreate(&d_desc, CUDA_R_32F, ne0, n, ne0));
if (batch_count > 1) {
fp8_set_batch(a_desc, batch_count, stride_a);
fp8_set_batch(b_desc, batch_count, ne10*n);
fp8_set_batch(d_desc, batch_count, ne0*n);
}
// TODO: Allow workspace in the heuristic search for potential performance gains.
// cuBLASLt takes caller-owned workspace per matmul call.
cublasLtMatmulPreference_t preference;
CUBLAS_CHECK(cublasLtMatmulPreferenceCreate(&preference));
cublasLtMatmulHeuristicResult_t heuristic;
int returned = 0;
const cublasStatus_t status = cublasLtMatmulAlgoGetHeuristic(
ctx.cublaslt_handle(), op_desc, a_desc, b_desc, d_desc, d_desc, preference, 1, &heuristic, &returned);
CUBLAS_CHECK(cublasLtMatmulPreferenceDestroy(preference));
if (status != CUBLAS_STATUS_SUCCESS || returned == 0) {
fp8_destroy_matmul(op_desc, a_desc, b_desc, d_desc);
return false;
}
const float alpha = 1.0f;
const float beta = 0.0f;
for (int64_t i3 = 0; i3 < n_calls; ++i3) {
const char * a = (const char *) src0->data + (i3/r3)*nb03;
const uint8_t * b = src1_fp8.ptr + i3*ne12*ne11*ne10;
float * d = (float *) ((char *) dst->data + i3*dst->nb[3]);
CUBLAS_CHECK(cublasLtMatmul(ctx.cublaslt_handle(), op_desc, &alpha, a, a_desc, b, b_desc,
&beta, d, d_desc, d, d_desc, &heuristic.algo, nullptr, 0, stream));
}
fp8_destroy_matmul(op_desc, a_desc, b_desc, d_desc);
return true;
}
#else
bool ggml_cuda_mul_mat_fp8(
ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
const ggml_tensor * weight_scale) {
GGML_UNUSED_VARS(ctx, src0, src1, dst, weight_scale);
return false;
}
#endif
-7
View File
@@ -1,7 +0,0 @@
#pragma once
#include "common.cuh"
bool ggml_cuda_mul_mat_fp8(
ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
const ggml_tensor * weight_scale = nullptr);
-4
View File
@@ -316,10 +316,6 @@ static void ggml_cuda_get_rows_switch_src0_type(
get_rows_cuda_float((const nv_bfloat16 *) src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
break;
case GGML_TYPE_F8_E4M3:
get_rows_cuda_float((const ggml_fp8_e4m3_t *) src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
break;
case GGML_TYPE_Q1_0:
get_rows_cuda_q<QK1_0, QR1_0, dequantize_q1_0>(src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
+34 -242
View File
@@ -4,7 +4,6 @@
#include "ggml-cuda/allreduce.cuh"
#include "ggml-cuda/common.cuh"
#include "ggml-cuda/graph.cuh"
#include "ggml-cuda/acc.cuh"
#include "ggml-cuda/add-id.cuh"
#include "ggml-cuda/arange.cuh"
@@ -28,7 +27,6 @@
#include "ggml-cuda/diagmask.cuh"
#include "ggml-cuda/diag.cuh"
#include "ggml-cuda/fattn.cuh"
#include "ggml-cuda/fp8.cuh"
#include "ggml-cuda/fwht.cuh"
#include "ggml-cuda/getrows.cuh"
#include "ggml-cuda/im2col.cuh"
@@ -585,7 +583,6 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool {
CUdeviceptr pool_addr = 0;
size_t pool_used = 0;
size_t pool_size = 0;
size_t max_size = CUDA_POOL_VMM_MAX_SIZE;
size_t granularity;
#if defined(GGML_USE_HIP)
std::vector<std::pair<CUdeviceptr, size_t>> mappings;
@@ -607,7 +604,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool {
#else
CU_CHECK(cuMemUnmap(pool_addr, pool_size));
#endif
CU_CHECK(cuMemAddressFree(pool_addr, max_size));
CU_CHECK(cuMemAddressFree(pool_addr, CUDA_POOL_VMM_MAX_SIZE));
}
}
@@ -623,21 +620,7 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool {
size_t reserve_size = size - avail;
reserve_size = granularity * ((reserve_size + granularity - 1) / granularity);
// reserve virtual address space (if not already reserved)
if (pool_addr == 0) {
CUresult err = cuMemAddressReserve(&pool_addr, max_size, 0, 0, 0);
while (err == CUDA_ERROR_OUT_OF_MEMORY && (max_size >> 1) >= reserve_size && (max_size >> 1) >= granularity) {
max_size >>= 1;
err = cuMemAddressReserve(&pool_addr, max_size, 0, 0, 0);
}
CU_CHECK(err);
if (max_size < CUDA_POOL_VMM_MAX_SIZE) {
GGML_LOG_WARN(GGML_CUDA_NAME " pool[%d]: reserved %.2f GiB of virtual memory instead of %.2f GiB\n",
device, max_size / (1024.0 * 1024.0 * 1024.0), CUDA_POOL_VMM_MAX_SIZE / (1024.0 * 1024.0 * 1024.0));
}
}
GGML_ASSERT(pool_size + reserve_size <= max_size);
GGML_ASSERT(pool_size + reserve_size <= CUDA_POOL_VMM_MAX_SIZE);
// allocate more physical memory
CUmemAllocationProp prop = {};
@@ -647,6 +630,11 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool {
CUmemGenericAllocationHandle handle;
CU_CHECK(cuMemCreate(&handle, reserve_size, &prop, 0));
// reserve virtual address space (if not already reserved)
if (pool_addr == 0) {
CU_CHECK(cuMemAddressReserve(&pool_addr, CUDA_POOL_VMM_MAX_SIZE, 0, 0, 0));
}
// map at the end of the pool
CUdeviceptr start_ptr = (CUdeviceptr)((char *)(pool_addr) + pool_size);
CU_CHECK(cuMemMap(start_ptr, reserve_size, 0, handle, 0));
@@ -747,124 +735,6 @@ std::unique_ptr<ggml_cuda_pool> ggml_backend_cuda_context::new_pool_for_device(g
return std::unique_ptr<ggml_cuda_pool>(new ggml_cuda_pool_leg(ctx, device));
}
ggml_backend_cuda_context::ggml_backend_cuda_context(int device) :
device(device),
name(GGML_CUDA_NAME + std::to_string(device)),
concurrent_stream_context(std::make_unique<ggml_cuda_stream_context>()) {
}
#ifdef USE_CUDA_GRAPH
ggml_cuda_graph * ggml_backend_cuda_context::cuda_graph(const void * first_node_ptr, int64_t time_now) {
// Give all cached graphs 10s to be reused after an idle period, without updating each graph.
if (!cuda_graphs.empty() && time_now - last_graph_use_time >= 10'000'000) {
next_graph_eviction_time = time_now + 10'000'000;
}
last_graph_use_time = time_now;
// Sweep every 5s during active use, evicting graphs unused for at least 10s.
if (time_now >= next_graph_eviction_time) {
next_graph_eviction_time = time_now + 5'000'000;
for (auto it = cuda_graphs.begin(); it != cuda_graphs.end(); ) {
if (it->first != first_node_ptr && time_now - it->second->last_used_time >= 10'000'000) {
it = cuda_graphs.erase(it);
} else {
++it;
}
}
}
auto it = cuda_graphs.find(first_node_ptr);
if (it == cuda_graphs.end()) {
it = cuda_graphs.emplace(first_node_ptr, std::make_unique<ggml_cuda_graph>()).first;
}
it->second->last_used_time = time_now;
return it->second.get();
}
bool ggml_backend_cuda_context::any_cuda_graph_enabled() const {
for (const auto & [key, graph] : cuda_graphs) {
if (graph && graph->is_enabled()) {
return true;
}
}
return false;
}
bool ggml_backend_cuda_context::any_cuda_graph_has_instance() const {
for (const auto & [key, graph] : cuda_graphs) {
if (graph && !graph->instances.empty()) {
return true;
}
}
return false;
}
#endif // USE_CUDA_GRAPH
#if defined(GGML_CUDA_USE_PDL)
bool ggml_cuda_kernel_can_use_pdl(const void * kernel) {
static const bool env_pdl_enabled = []() {
const char * env = getenv("GGML_CUDA_PDL");
return env == nullptr || std::atoi(env) != 0;
}();
if (!env_pdl_enabled) {
return false;
}
const int device = ggml_cuda_get_device();
struct cache_key {
int device;
const void * kernel;
bool operator==(const cache_key & other) const { return device == other.device && kernel == other.kernel; }
};
struct cache_key_hash {
// MurmurHash3 mixing function for better hash distribution (vs. just std::hash which in some implementations simply returns the identity)
static size_t hash_mix(size_t x) {
std::uint64_t y = x;
const std::uint64_t m = 0xe9846af9b1a615d;
y ^= y >> 32;
y *= m;
y ^= y >> 32;
y *= m;
y ^= y >> 28;
return static_cast<size_t>(y);
}
size_t operator()(const cache_key & key) const {
// Use a nonzero seed to avoid mapping all-zero keys to zero
size_t h = 42;
h = hash_mix(h + key.device);
h = hash_mix(h + reinterpret_cast<size_t>(key.kernel));
return h;
}
};
static std::mutex cache_mutex;
static std::unordered_map<cache_key, bool, cache_key_hash> cache;
const cache_key key = { device, kernel };
std::lock_guard<std::mutex> lock(cache_mutex);
const auto it = cache.find(key);
if (it != cache.end()) {
return it->second;
}
cudaFuncAttributes attr = {};
CUDA_CHECK(cudaFuncGetAttributes(&attr, kernel));
// PDL device-side primitives are emitted only for PTX versions >= 90.
// We have to guard on a loaded kernel's PTX version so a kernel forward-JIT'ed
// from pre-Hopper PTX to a Hopper-or-newer GPU does not opt into PDL.
const bool can_use_pdl = attr.ptxVersion >= 90;
cache.emplace(key, can_use_pdl);
return can_use_pdl;
}
#endif // defined(GGML_CUDA_USE_PDL)
// destroying a cuBLAS handle while a graph is being captured in a different thread can result in a CUDA error
// this lock is used to ensure that no cuBLAS handle is destroyed while a graph is being captured
@@ -891,11 +761,6 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() {
CUDA_CHECK(cudaFree(cublas_workspaces[i][j]));
}
}
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
if (cublaslt_handles[i] != nullptr) {
CUBLAS_CHECK(cublasLtDestroy(cublaslt_handles[i]));
}
#endif
}
}
@@ -2004,9 +1869,9 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
const int cc = ggml_cuda_info().devices[ctx.device].cc;
const ggml_prec prec = ggml_prec_get_acc(dst);
const ggml_prec prec = (ggml_prec) ggml_get_op_params_i32(dst, 0);
ggml_type compute_type = src0->type;
if (ggml_is_quantized(compute_type) || compute_type == GGML_TYPE_F8_E4M3) {
if (ggml_is_quantized(compute_type)) {
compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc)) {
compute_type = GGML_TYPE_F32;
@@ -2177,7 +2042,7 @@ static bool ggml_cuda_should_fuse_mul_mat(const ggml_tensor * ffn_up,
return true;
}
static bool ggml_cuda_should_fuse_mul_mat_vec_f(const ggml_tensor * tensor, const ggml_tensor * gate = nullptr) {
static bool ggml_cuda_should_fuse_mul_mat_vec_f(const ggml_tensor * tensor) {
ggml_tensor * src0 = tensor->src[0];
ggml_tensor * src1 = tensor->src[1];
const ggml_tensor * dst = tensor;
@@ -2185,17 +2050,12 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_f(const ggml_tensor * tensor, cons
const bool is_mul_mat_id = tensor->op == GGML_OP_MUL_MAT_ID;
bool use_mul_mat_vec_f =
(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_BF16 ||
src0->type == GGML_TYPE_F8_E4M3) &&
(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_BF16) &&
src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32;
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
use_mul_mat_vec_f = use_mul_mat_vec_f && ggml_cuda_should_use_mmvf(src0, cc, warp_size, is_mul_mat_id ? src1->ne[2] : src1->ne[1]);
if (gate && !ggml_cuda_should_use_mmvf(gate->src[0], cc, warp_size, is_mul_mat_id ? src1->ne[2] : src1->ne[1])) {
return false;
}
use_mul_mat_vec_f = use_mul_mat_vec_f && ggml_cuda_should_use_mmvf(src0->type, cc, warp_size, src0->ne, src0->nb, is_mul_mat_id ? src1->ne[2] : src1->ne[1]);
//we only support fusion for ncols_dst = 1
if (tensor->op == GGML_OP_MUL_MAT && dst->ne[1] != 1) {
@@ -2219,11 +2079,11 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) {
ggml_nbytes(src0) != ggml_backend_buffer_get_alloc_size(src0->buffer, src0) &&
src0->view_src;
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
bool use_mul_mat_vec_q = ggml_cuda_should_use_mmvq(src0->type, cc, src1->ne[1]) && !bad_padding_clear &&
src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32;
bool use_mul_mat_vec_q = ggml_is_quantized(src0->type) && !bad_padding_clear && src1->type == GGML_TYPE_F32 &&
dst->type == GGML_TYPE_F32 && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE;
// fusion is not universally faster on Pascal
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
if (cc <= GGML_CUDA_CC_PASCAL) {
return false;
}
@@ -2291,7 +2151,7 @@ static bool ggml_cuda_match_shared_expert(const ggml_cgraph * graph, int routed_
return true;
}
// The weight scale that scale_node (MUL) applies to the output of mm_node (MUL_MAT), or nullptr.
// The NVFP4 weight scale that scale_node (MUL) applies to the output of mm_node (MUL_MAT), or nullptr.
// Shared by ggml_cuda_try_fuse and graph_optimize so that both accept the same patterns.
static const ggml_tensor * ggml_cuda_get_mul_mat_scale(const ggml_tensor * scale_node, const ggml_tensor * mm_node) {
const bool scale_lhs_mm = scale_node->src[0] == mm_node;
@@ -2301,7 +2161,7 @@ static const ggml_tensor * ggml_cuda_get_mul_mat_scale(const ggml_tensor * scale
}
const ggml_tensor * scale = scale_lhs_mm ? scale_node->src[1] : scale_node->src[0];
if ((mm_node->src[0]->type != GGML_TYPE_NVFP4 && mm_node->src[0]->type != GGML_TYPE_F8_E4M3) || scale_node->type != GGML_TYPE_F32 ||
if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 ||
scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != 1 ||
!ggml_are_same_shape(scale_node, mm_node)) {
return nullptr;
@@ -2322,7 +2182,7 @@ static const ggml_tensor * ggml_cuda_get_mul_mat_id_scale(const ggml_tensor * re
}
const ggml_tensor * scale = reshape->src[0];
if ((mm_node->src[0]->type != GGML_TYPE_NVFP4 && mm_node->src[0]->type != GGML_TYPE_F8_E4M3) || scale_node->type != GGML_TYPE_F32 ||
if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 ||
scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm_node->src[0]->ne[2] ||
!ggml_are_same_shape(scale_node, mm_node)) {
return nullptr;
@@ -2420,17 +2280,17 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
const int cc = ggml_cuda_info().devices[ctx.device].cc;
const int warp_size = ggml_cuda_info().devices[ctx.device].warp_size;
if (ggml_cuda_should_use_mmvf(src0, cc, warp_size, ne11)) {
// The custom vector kernel can be used over batched cuBLAS GEMM.
if (ggml_cuda_should_use_mmvf(src0->type, cc, warp_size, src0->ne, src0->nb, ne11)) {
// The custom F16 vector kernel can be used over batched cuBLAS GEMM.
// But this is only faster for GPUs without tensor cores or with a thin src0 matrix (particularly KQV in attention)
ggml_cuda_mul_mat_vec_f(ctx, src0, src1, nullptr, dst);
return;
}
// A transposed vector can still use MMVF (i.e. ne01 == 1)
// A transposed vector can still use MMVQ (i.e. ne01 == 1)
if (ne01 == 1 && ne11 > MMVF_MAX_BATCH_SIZE && ne2 == 1 && ne3 == 1
&& src0->type == GGML_TYPE_F32
&& ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst)
&& ggml_cuda_should_use_mmvf(src1, cc, warp_size, /*ne11 =*/ 1)) {
&& ggml_cuda_should_use_mmvf(src1->type, cc, warp_size, src1->ne, src1->nb, /*ne11 =*/ 1)) {
ggml_tensor dst_vec = *dst;
dst_vec.ne[0] = ne11;
dst_vec.ne[1] = 1;
@@ -2452,9 +2312,6 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
ggml_cuda_mul_mat_q(ctx, src0, src1, nullptr, dst);
return;
}
if (src0->type == GGML_TYPE_F8_E4M3 && ggml_cuda_mul_mat_fp8(ctx, src0, src1, dst)) {
return;
}
ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
}
@@ -2463,6 +2320,7 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
static bool ggml_cuda_mul_mat_id_needs_sync(const ggml_tensor * dst, const int cc) {
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
if (src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
@@ -2474,10 +2332,7 @@ static bool ggml_cuda_mul_mat_id_needs_sync(const ggml_tensor * dst, const int c
if (dst->ne[2] <= get_mmvq_mmid_max_batch(src0->type, cc)) {
return false;
}
} else if (src0->type == GGML_TYPE_F8_E4M3 &&
ggml_cuda_should_use_mmvf(src0, cc, warp_size, dst->ne[2])) {
return false;
} else if (src0->type != GGML_TYPE_F8_E4M3 && GGML_CUDA_CC_IS_AMD(cc)) {
} else if (GGML_CUDA_CC_IS_AMD(cc)) {
return false;
}
}
@@ -2503,7 +2358,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
GGML_TENSOR_BINARY_OP_LOCALS
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
// [TAG_MUL_MAT_ID_CUDA_GRAPHS]
@@ -2516,12 +2371,8 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
ggml_cuda_mul_mat_vec_q(ctx, src0, src1, ids, dst);
return;
}
} else if (src0->type == GGML_TYPE_F8_E4M3 &&
ggml_cuda_should_use_mmvf(src0, cc, warp_size, ne2)) {
ggml_cuda_mul_mat_vec_f(ctx, src0, src1, ids, dst);
return;
} else {
if (src0->type != GGML_TYPE_F8_E4M3 && GGML_CUDA_CC_IS_AMD(cc)) {
if (GGML_CUDA_CC_IS_AMD(cc)) {
ggml_cuda_mul_mat_vec_f(ctx, src0, src1, ids, dst);
return;
}
@@ -2547,7 +2398,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
GGML_ASSERT(nb2 % nb1 == 0);
const ggml_type type_src1_sorted = (src0->type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc))
|| ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3 ? GGML_TYPE_F32 : src0->type;
|| ggml_is_quantized(src0->type) ? GGML_TYPE_F32 : src0->type;
const ggml_type type_dst_sorted = GGML_TYPE_F32;
const size_t ts_src1_sorted = ggml_type_size(type_src1_sorted);
const size_t ts_dst_sorted = ggml_type_size(type_dst_sorted);
@@ -4573,13 +4424,6 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
fusion_data.glu_op = ggml_get_glu_op(glu);
fusion_data.glu_limit = ggml_get_op_params_f32(glu, 3);
if (ggml_cuda_should_fuse_mul_mat_vec_f(up_n, gate_n)) {
ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data);
fused_mul_mat_vec = true;
fused_node_count = n_ops;
break;
}
if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) {
ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data);
fused_mul_mat_vec = true;
@@ -4682,13 +4526,6 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
fusion_data.glu_op = ggml_get_glu_op(glu);
fusion_data.glu_limit = ggml_get_op_params_f32(glu, 3);
if (ggml_cuda_should_fuse_mul_mat_vec_f(up_n, gate_n)) {
ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data);
fused_mul_mat_vec = true;
fused_node_count = n_ops;
break;
}
if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) {
ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data);
fused_mul_mat_vec = true;
@@ -4746,7 +4583,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
const ggml_tensor * src1 = up_n->src[1];
const ggml_tensor * ids = up_n->src[2];
if (ggml_cuda_should_fuse_mul_mat_vec_f(up_n, gate_n)) {
if (ggml_cuda_should_fuse_mul_mat_vec_f(up_n)) {
ggml_cuda_mm_fusion_args_host fusion_data{};
fusion_data.gate = gate_n->src[0];
fusion_data.x_bias = up_bias_tensor;
@@ -4789,7 +4626,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
const ggml_tensor * src1 = up->src[1];
const ggml_tensor * ids = up->src[2];
if (ggml_cuda_should_fuse_mul_mat_vec_f(up, gate)) {
if (ggml_cuda_should_fuse_mul_mat_vec_f(up)) {
ggml_cuda_mm_fusion_args_host fusion_data{};
fusion_data.gate = gate->src[0];
fusion_data.glu_op = ggml_get_glu_op(glu);
@@ -4907,19 +4744,6 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
fusion_data.x_bias = bias;
fusion_data.x_scale = scale;
if (ggml_cuda_should_fuse_mul_mat_vec_f(mm_node)) {
ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, out_node, &fusion_data);
fused_mul_mat_vec = true;
fused_node_count = n_ops;
break;
}
if (op == GGML_OP_MUL_MAT && !with_bias && src0->type == GGML_TYPE_F8_E4M3 &&
ggml_get_op_params_i32(mm_node, 3) != GGML_PREC_F32 &&
ggml_cuda_mul_mat_fp8(*cuda_ctx, src0, src1, scale_node, scale)) {
return 1;
}
if (ggml_cuda_should_fuse_mul_mat_vec_q(mm_node)) {
ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, out_node, &fusion_data);
fused_mul_mat_vec = true;
@@ -5319,9 +5143,7 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend,
ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key);
if (graph->is_enabled()) {
// graph->uid is set only for a compatible graph (see ggml_cuda_graph_update_required), so a reused graph needs no new check
const bool graph_reused = cgraph->uid != 0 && cgraph->uid == graph->uid;
const bool graph_compatible = graph_reused || ggml_cuda_graph_check_compability(cgraph);
const bool graph_compatible = ggml_cuda_graph_check_compability(cgraph);
if (graph_compatible) {
const bool properties_changed = ggml_cuda_graph_update_required(cuda_ctx, cgraph);
@@ -6156,16 +5978,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
if (b->type == GGML_TYPE_F16 && a->type != GGML_TYPE_F16 && !ggml_cuda_op_mul_mat_use_fwht(op)) {
return false;
}
const int cc_dev = ggml_cuda_info().devices[dev_ctx->device].cc;
const bool f32_src1_ok = GGML_CUDA_CC_IS_NVIDIA(cc_dev) && cc_dev < GGML_CUDA_CC_AMPERE;
const enum ggml_prec prec_src1 = ggml_prec_get_src(op, 1);
const bool fp8_id_mmvf = op->op == GGML_OP_MUL_MAT_ID &&
a->type == GGML_TYPE_F8_E4M3 && b->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&
ggml_cuda_should_use_mmvf(a, cc_dev, ggml_cuda_info().devices[dev_ctx->device].warp_size, b->ne[2]);
if (prec_src1 == GGML_PREC_F32 && !fp8_id_mmvf && (op->op == GGML_OP_MUL_MAT_ID || !f32_src1_ok || a->type != GGML_TYPE_F32)) {
return false;
}
if (prec_src1 == GGML_PREC_BF16 && (op->op == GGML_OP_MUL_MAT_ID || a->type != GGML_TYPE_F32)) {
if (op->op == GGML_OP_MUL_MAT_ID && (ggml_get_op_params_i32(op, 3) == GGML_PREC_BF16 || ggml_get_op_params_i32(op, 3) == GGML_PREC_F32)) {
return false;
}
#ifdef GGML_USE_MUSA
@@ -6193,7 +6006,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_TYPE_Q8_0:
case GGML_TYPE_MXFP4:
case GGML_TYPE_NVFP4:
case GGML_TYPE_F8_E4M3:
case GGML_TYPE_Q2_K:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
@@ -6231,7 +6043,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_F8_E4M3:
case GGML_TYPE_Q2_K:
case GGML_TYPE_Q3_K:
case GGML_TYPE_Q4_K:
@@ -6261,7 +6072,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
{
return (
(
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 || op->type == GGML_TYPE_F8_E4M3 ||
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 ||
op->type == GGML_TYPE_Q4_0 || op->type == GGML_TYPE_Q4_1 || op->type == GGML_TYPE_Q5_0 ||
op->type == GGML_TYPE_Q5_1 || op->type == GGML_TYPE_Q8_0 || op->type == GGML_TYPE_IQ4_NL) &&
op->src[0]->type == GGML_TYPE_F32
@@ -6282,16 +6093,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
{
ggml_type src0_type = op->src[0]->type;
ggml_type src1_type = op->src[1]->type;
if (src0_type == GGML_TYPE_F8_E4M3 && src1_type == GGML_TYPE_F8_E4M3) {
return true;
}
if ((src0_type == GGML_TYPE_F32 || src0_type == GGML_TYPE_F16 || src0_type == GGML_TYPE_BF16) &&
src1_type == GGML_TYPE_F8_E4M3) {
return true;
}
if (src0_type == GGML_TYPE_F8_E4M3 && src1_type == GGML_TYPE_F32) {
return true;
}
if ((src0_type == GGML_TYPE_F32 || src0_type == GGML_TYPE_BF16 || src0_type == GGML_TYPE_F16) &&
(src1_type == GGML_TYPE_F32 || src1_type == GGML_TYPE_BF16 || src1_type == GGML_TYPE_F16)
) {
@@ -6709,21 +6510,12 @@ static ggml_backend_feature * ggml_backend_cuda_get_features(ggml_backend_reg_t
{
const auto & info = ggml_cuda_info();
bool has_fp8 = false;
bool has_fp4 = false;
for (int id = 0; id < info.device_count; ++id) {
has_fp8 |= fp8_mma_hardware_available(info.devices[id].cc);
has_fp4 |= blackwell_mma_available(info.devices[id].cc);
if (has_fp8 && has_fp4) {
if (blackwell_mma_available(info.devices[id].cc)) {
features.push_back({ "BLACKWELL_NATIVE_FP4", "1"});
break;
}
}
if (has_fp8) {
features.push_back({ "NATIVE_FP8", "1"});
}
if (has_fp4) {
features.push_back({ "BLACKWELL_NATIVE_FP4", "1"});
}
}
#undef _STRINGIFY
-191
View File
@@ -1,191 +0,0 @@
#pragma once
// CUDA graph capture and concurrent stream state, only needed by ggml-cuda.cu
#include "common.cuh"
#include <algorithm>
#include <unordered_map>
#include <utility>
#include <vector>
struct ggml_cuda_graph {
#ifdef USE_CUDA_GRAPH
~ggml_cuda_graph() {
for (cudaGraphExec_t instance : instances) {
CUDA_CHECK(cudaGraphExecDestroy(instance));
}
}
// the ggml graph is captured as a series of CUDA graphs, so the GPU can run one chunk while the next is enqueued
std::vector<cudaGraphExec_t> instances;
size_t num_nodes = 0;
std::vector<cudaGraphNode_t> nodes;
bool disable_due_to_gpu_arch = false;
bool warmup_complete = false;
uint64_t uid = 0;
int64_t last_used_time = 0;
struct node_properties {
ggml_tensor node;
void * node_src_data_ptrs[GGML_MAX_SRC];
int64_t node_src_ne[GGML_MAX_SRC][GGML_MAX_DIMS];
size_t node_src_nb[GGML_MAX_SRC][GGML_MAX_DIMS];
};
std::vector<node_properties> node_props;
bool is_enabled() const {
static const bool disable_cuda_graphs_due_to_env = (getenv("GGML_CUDA_DISABLE_GRAPHS") != nullptr);
return !(disable_due_to_gpu_arch || disable_cuda_graphs_due_to_env);
}
#endif
};
struct ggml_cuda_concurrent_event {
std::vector<cudaEvent_t> join_events;
cudaEvent_t fork_event = nullptr;
int n_streams = 0;
std::unordered_map<const ggml_tensor *, int> stream_mapping;
// Original order of nodes in this concurrent region (before interleaving)
// Used to restore grouping for fusion within streams
std::vector<const ggml_tensor *> original_order;
const ggml_tensor * join_node;
ggml_cuda_concurrent_event() = default;
ggml_cuda_concurrent_event(const ggml_cuda_concurrent_event &) = delete;
ggml_cuda_concurrent_event & operator=(const ggml_cuda_concurrent_event &) = delete;
explicit ggml_cuda_concurrent_event(int n_streams) : n_streams(n_streams) {
join_events.resize(n_streams);
for (size_t i = 0; i < join_events.size(); ++i) {
CUDA_CHECK(cudaEventCreateWithFlags(&join_events[i], cudaEventDisableTiming));
}
CUDA_CHECK(cudaEventCreateWithFlags(&fork_event, cudaEventDisableTiming));
}
ggml_cuda_concurrent_event(ggml_cuda_concurrent_event && other) noexcept
: join_events(std::move(other.join_events))
, fork_event(other.fork_event)
, n_streams(other.n_streams)
, stream_mapping(std::move(other.stream_mapping))
, original_order(std::move(other.original_order))
, join_node(other.join_node) {
other.fork_event = nullptr;
}
// 1. check if any branches write to overlapping memory ranges (except the join node)
// 2. check whether all srcs are either within the branch or outside the nodes covered by ggml_cuda_concurrent_event
// we assume all nodes have the same buffer
bool is_valid() const {
std::vector<std::vector<std::pair<int64_t, int64_t>>> write_ranges;
write_ranges.resize(n_streams);
// get join_node's memory range to exclude from overlap checking.
// multiple nodes can use join_node's buffer; we synchronize on the join node.
const ggml_tensor * join_t = join_node->view_src ? join_node->view_src : join_node;
const int64_t join_start = (int64_t) join_t->data;
const int64_t join_end = join_start + ggml_nbytes(join_t);
for (const auto & [tensor, stream] : stream_mapping) {
const ggml_tensor * t = tensor->view_src ? tensor->view_src : tensor;
const int64_t t_start = (int64_t) t->data;
const int64_t t_end = t_start + ggml_nbytes(t);
// skip tensors that overlap with join_node's buffer.
if ((t_start <= join_start && join_start < t_end) || (join_start <= t_start && t_start < join_end)) {
continue;
}
// concurrent streams begin from 1
write_ranges[stream - 1].emplace_back(t_start, t_end);
}
for (int i = 0; i < n_streams; ++i) {
// sorts first by start then by end of write range
std::sort(write_ranges[i].begin(), write_ranges[i].end());
}
bool writes_overlap = false;
bool dependent_srcs = false;
for (const auto & [tensor, stream] : stream_mapping) {
const ggml_tensor * t = tensor->view_src ? tensor->view_src : tensor;
const int64_t t_start = (int64_t) t->data;
const int64_t t_end = t_start + ggml_nbytes(t);
// skip tensors that overlap with join_node's buffer
if ((t_start <= join_start && join_start < t_end) || (join_start <= t_start && t_start < join_end)) {
continue;
}
// check if this buffer's write data overlaps with another stream's
std::pair<int64_t, int64_t> data_range = std::make_pair(t_start, t_end);
for (int i = 0; i < n_streams; ++i) {
if (i == stream - 1) {
continue;
}
auto it = std::lower_bound(write_ranges[i].begin(), write_ranges[i].end(), data_range);
if (it != write_ranges[i].end()) {
const std::pair<int64_t, int64_t> & other = *it;
// std::lower_bound returns the first element where other >= data_range (lexicographically).
// This guarantees other.first >= data_range.first.
// Therefore, overlap occurs iff other.first < data_range.second
// (i.e., the other range starts before this range ends).
if (other.first < data_range.second) {
GGML_LOG_DEBUG("Writes overlap for %s", tensor->name);
writes_overlap = true;
break;
}
}
}
//check if all srcs are either in branch or don't have a branch
for (int i = 0; i < GGML_MAX_SRC; ++i) {
if (!tensor->src[i]) {
continue;
}
auto it = stream_mapping.find(tensor->src[i]);
if (it == stream_mapping.end()) {
continue;
}
if (it->second != stream) {
dependent_srcs = true;
break;
}
}
if (dependent_srcs || writes_overlap) {
break;
}
}
return !writes_overlap && !dependent_srcs;
}
~ggml_cuda_concurrent_event() {
if (fork_event != nullptr) {
CUDA_CHECK(cudaEventDestroy(fork_event));
}
for (cudaEvent_t e : join_events) {
if (e != nullptr) {
CUDA_CHECK(cudaEventDestroy(e));
}
}
}
};
struct ggml_cuda_stream_context {
std::unordered_map<const ggml_tensor *, ggml_cuda_concurrent_event> concurrent_events;
void reset() {
concurrent_events.clear();
}
};
+13 -106
View File
@@ -5,14 +5,6 @@
#include "mmf.cuh"
#include "convert.cuh"
static __device__ __forceinline__ float2 mmvf_load_fp8x2(const ggml_fp8_e4m3_t * x, const int col2) {
#if defined(FP8_AVAILABLE) && !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
return ggml_cuda_cast<float2>(reinterpret_cast<const __nv_fp8x2_e4m3 *>(x)[col2]);
#else
return make_float2(ggml_cuda_cast<float>(x[2*col2]), ggml_cuda_cast<float>(x[2*col2 + 1]));
#endif
}
template <typename T, typename type_acc, int ncols_dst, int block_size, bool has_fusion = false, bool is_multi_token_id = false>
static __global__ void mul_mat_vec_f(
const T * x_ptr, const float * y_ptr, const int32_t * ids_ptr, const ggml_cuda_mm_fusion_args_device fusion, float * dst_ptr,
@@ -64,24 +56,16 @@ static __global__ void mul_mat_vec_f(
bool use_gate = false;
bool use_bias = false;
bool use_gate_bias = false;
bool use_scale = false;
bool use_gate_scale = false;
ggml_glu_op glu_op = ggml_glu_op::GGML_GLU_OP_SWIGLU;
float glu_limit = 0.0f;
const T * gate_x = nullptr;
const float * x_bias = nullptr;
const float * gate_bias = nullptr;
const float * x_scale = nullptr;
const float * gate_scale = nullptr;
if constexpr (has_fusion) {
use_gate = fusion.gate != nullptr;
use_bias = fusion.x_bias != nullptr;
use_gate_bias = fusion.gate_bias != nullptr;
if constexpr (std::is_same_v<T, ggml_fp8_e4m3_t>) {
use_scale = fusion.x_scale != nullptr;
use_gate_scale = fusion.gate_scale != nullptr && use_gate;
}
glu_op = fusion.glu_op;
glu_limit = fusion.glu_limit;
@@ -97,12 +81,6 @@ static __global__ void mul_mat_vec_f(
} else {
use_gate_bias = false;
}
if (use_scale) {
x_scale = static_cast<const float *>(fusion.x_scale);
}
if (use_gate_scale) {
gate_scale = static_cast<const float *>(fusion.gate_scale);
}
}
if (use_gate) {
@@ -119,6 +97,8 @@ static __global__ void mul_mat_vec_f(
}
}
const float2 * y2 = (const float2 *) y;
extern __shared__ char data_mmv[];
float * buf_iw = (float *) data_mmv;
[[maybe_unused]] float * buf_iw_gate = nullptr;
@@ -148,7 +128,6 @@ static __global__ void mul_mat_vec_f(
}
if constexpr (std::is_same_v<T, float>) {
const float2 * y2 = (const float2 *) y;
const float2 * x2 = (const float2 *) x;
[[maybe_unused]] const float2 * gate_x2 = nullptr;
if constexpr (has_fusion) {
@@ -181,7 +160,6 @@ static __global__ void mul_mat_vec_f(
}
}
} else if constexpr (std::is_same_v<T, half>) {
const float2 * y2 = (const float2 *) y;
const half2 * x2 = (const half2 *) x;
[[maybe_unused]] const half2 * gate_x2 = nullptr;
if constexpr (has_fusion) {
@@ -257,7 +235,6 @@ static __global__ void mul_mat_vec_f(
#endif // FP16_AVAILABLE
}
} else if constexpr (std::is_same_v<T, nv_bfloat16>) {
const float2 * y2 = (const float2 *) y;
//TODO: add support for ggml_cuda_mad for hip_bfloat162
#if defined(GGML_USE_HIP)
const int * x2 = (const int *) x;
@@ -324,28 +301,6 @@ static __global__ void mul_mat_vec_f(
}
}
#endif
} else if constexpr (std::is_same_v<T, ggml_fp8_e4m3_t>) {
const float2 * y2 = (const float2 *) y;
for (int col2 = tid; col2 < ncols2; col2 += block_size) {
const float2 tmpx = mmvf_load_fp8x2(x, col2);
float2 tmpx_gate = {};
if constexpr (has_fusion) {
if (use_gate) {
tmpx_gate = mmvf_load_fp8x2(gate_x, col2);
}
}
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
const float2 tmpy = y2[j*stride_col_y2 + col2];
ggml_cuda_mad(sumf[j], tmpx, tmpy);
if constexpr (has_fusion) {
if (use_gate) {
ggml_cuda_mad(sumf_gate[j], tmpx_gate, tmpy);
}
}
}
}
} else {
static_assert(std::is_same_v<T, void>, "unsupported type");
}
@@ -393,18 +348,12 @@ static __global__ void mul_mat_vec_f(
float value = sumf[tid];
if constexpr (has_fusion) {
if (use_scale) {
value *= x_scale[ids ? channel_x : 0];
}
if (use_bias) {
value += x_bias[tid*stride_col_dst + row];
}
if (use_gate) {
float gate_value = sumf_gate[tid];
if (use_gate_scale) {
gate_value *= gate_scale[ids ? channel_x : 0];
}
if (use_gate_bias) {
gate_value += gate_bias[tid*stride_col_dst + row];
}
@@ -431,11 +380,7 @@ static __global__ void mul_mat_vec_f(
dst[tid*stride_col_dst + row] = value;
if constexpr (!has_fusion) {
GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, use_scale, use_gate_scale, glu_op, glu_limit, gate_x, x_bias, gate_bias,
x_scale, gate_scale, sumf_gate);
}
if constexpr (!std::is_same_v<T, ggml_fp8_e4m3_t>) {
GGML_UNUSED_VARS(use_scale, use_gate_scale, x_scale, gate_scale);
GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, glu_op, glu_limit, gate_x, x_bias, gate_bias, sumf_gate);
}
}
@@ -450,8 +395,7 @@ static void mul_mat_vec_f_switch_fusion(
const ggml_cuda_kernel_launch_params launch_params = {block_nums, block_dims, nbytes_shared, stream};
const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr ||
fusion.x_scale != nullptr || fusion.gate_scale != nullptr;
const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr;
if constexpr (ncols_dst == 1) {
if (has_fusion) {
ggml_cuda_kernel_launch(mul_mat_vec_f<T, type_acc, ncols_dst, block_size, true, is_multi_token_id>, launch_params,
@@ -672,7 +616,7 @@ static void mul_mat_vec_f_cuda(
const int64_t ids_stride, enum ggml_prec prec, cudaStream_t stream) {
if constexpr(std::is_same_v<T, half>) {
if (prec == GGML_PREC_UNDEFINED) {
if (prec == GGML_PREC_DEFAULT) {
mul_mat_vec_f_cuda_switch_ncols_dst<T, half>
(x, y, ids, fusion, dst, ncols, nrows, ncols_dst, stride_row, stride_col_y, stride_col_dst,
nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y,
@@ -705,11 +649,9 @@ void ggml_cuda_mul_mat_vec_f(ggml_backend_cuda_context & ctx, const ggml_tensor
GGML_ASSERT( nb10 == ts_src1);
GGML_ASSERT(!ids || ids->nb[0] == ggml_type_size(ids->type));
GGML_ASSERT( nb0 == ts_dst);
GGML_ASSERT(src0->type != GGML_TYPE_F8_E4M3 || (uintptr_t) src0->data % alignof(uint16_t) == 0);
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
// the fused GLU path gets the GLU node as dst, which carries no matmul precision, so it accumulates in F32
const enum ggml_prec prec = fast_fp16_available(cc) && !(fusion && fusion->gate) ? ggml_prec_get_acc(dst) : GGML_PREC_F32;
const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec(dst->op_params[0]) : GGML_PREC_F32;
const float * src1_d = (const float *) src1->data;
const int32_t * ids_d = ids ? (const int32_t *) ids->data : nullptr;
@@ -720,7 +662,6 @@ void ggml_cuda_mul_mat_vec_f(ggml_backend_cuda_context & ctx, const ggml_tensor
if (fusion) {
GGML_ASSERT( !ids || dst->ne[2] == 1);
GGML_ASSERT( ids || dst->ne[1] == 1);
GGML_ASSERT((fusion->x_scale == nullptr && fusion->gate_scale == nullptr) || src0->type == GGML_TYPE_F8_E4M3);
if (fusion->x_bias) {
GGML_ASSERT(fusion->x_bias->type == GGML_TYPE_F32);
GGML_ASSERT(fusion->x_bias->ne[0] == dst->ne[0]);
@@ -729,7 +670,6 @@ void ggml_cuda_mul_mat_vec_f(ggml_backend_cuda_context & ctx, const ggml_tensor
}
if (fusion->gate) {
GGML_ASSERT(fusion->gate->type == src0->type && ggml_are_same_stride(fusion->gate, src0));
GGML_ASSERT(src0->type != GGML_TYPE_F8_E4M3 || (uintptr_t) fusion->gate->data % alignof(uint16_t) == 0);
fusion_local.gate = fusion->gate->data;
}
if (fusion->gate_bias) {
@@ -738,18 +678,6 @@ void ggml_cuda_mul_mat_vec_f(ggml_backend_cuda_context & ctx, const ggml_tensor
GGML_ASSERT(!ids || fusion->gate_bias->ne[1] == src0->ne[2]);
fusion_local.gate_bias = fusion->gate_bias->data;
}
if (fusion->x_scale) {
GGML_ASSERT(fusion->x_scale->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_contiguous(fusion->x_scale));
GGML_ASSERT(ggml_nelements(fusion->x_scale) == (ids ? src0->ne[2] : 1));
fusion_local.x_scale = fusion->x_scale->data;
}
if (fusion->gate_scale) {
GGML_ASSERT(fusion->gate_scale->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_contiguous(fusion->gate_scale));
GGML_ASSERT(ggml_nelements(fusion->gate_scale) == (ids ? src0->ne[2] : 1));
fusion_local.gate_scale = fusion->gate_scale->data;
}
fusion_local.glu_op = fusion->glu_op;
fusion_local.glu_limit = fusion->glu_limit;
}
@@ -794,12 +722,6 @@ void ggml_cuda_mul_mat_vec_f(ggml_backend_cuda_context & ctx, const ggml_tensor
ne02, nchannels_y, nchannels_dst, s02, stride_channel_y, stride_channel_dst,
ne03, ne3, s03, s13, s3, ids_stride, prec, ctx.stream());
} break;
case GGML_TYPE_F8_E4M3: {
const ggml_fp8_e4m3_t * src0_d = (const ggml_fp8_e4m3_t *) src0->data;
mul_mat_vec_f_cuda(src0_d, src1_d, ids_d, fusion_local, dst_d, ne00, ne01, ncols_dst, s01, stride_col_y, stride_col_dst,
ne02, nchannels_y, nchannels_dst, s02, stride_channel_y, stride_channel_dst,
ne03, ne3, s03, s13, s3, ids_stride, prec, ctx.stream());
} break;
default:
GGML_ABORT("unsupported type: %s", ggml_type_name(src0->type));
}
@@ -821,7 +743,7 @@ void ggml_cuda_op_mul_mat_vec_f(
const int id = ggml_cuda_get_device();
const int cc = ggml_cuda_info().devices[id].cc;
const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec_get_acc(dst) : GGML_PREC_F32;
const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec(dst->op_params[0]) : GGML_PREC_F32;
// ggml_cuda_op provides single, contiguous matrices
const int64_t stride_row = ne00;
@@ -866,11 +788,7 @@ void ggml_cuda_op_mul_mat_vec_f(
GGML_UNUSED_VARS(ctx, src1, dst, src1_ddq_i, src1_ncols, src1_padded_row_size);
}
bool ggml_cuda_should_use_mmvf(const ggml_tensor * src0, int cc, int warp_size, int64_t ne11) {
const ggml_type type = src0->type;
const int64_t * src0_ne = src0->ne;
const size_t * src0_nb = src0->nb;
bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, int warp_size, const int64_t * src0_ne, const size_t * src0_nb, int64_t ne11) {
if (src0_ne[0] % 2 != 0) {
return false;
}
@@ -880,8 +798,11 @@ bool ggml_cuda_should_use_mmvf(const ggml_tensor * src0, int cc, int warp_size,
return false;
}
if (!ggml_cuda_is_aligned(src0, 2*ts)) {
return false;
// Pointers not aligned to the size of half2/nv_bfloat162/float2 would result in a crash:
for (size_t i = 1; i < GGML_MAX_DIMS; ++i) {
if (src0_nb[i] % (2*ts) != 0) {
return false;
}
}
switch (type) {
@@ -941,18 +862,6 @@ bool ggml_cuda_should_use_mmvf(const ggml_tensor * src0, int cc, int warp_size,
if (src0_small && !ggml_cuda_should_use_mmf(type, cc, warp_size, src0_ne, src0_nb, ne11, /*mul_mat_id =*/ false)) {
return ne11 <= MMVF_MAX_BATCH_SIZE;
}
// MMF tiles 32 rows per block, so a matrix of few rows leaves most SMs idle while MMVF runs a block
// per row: with a long enough K, MMVF wins for the small batches of speculative decoding as long as
// its work (rows x batch) stays small relative to the SM count, with a budget that grows with K
if (src0_small && src0_ne[0] >= 1024) {
const bool blackwell = cc >= GGML_CUDA_CC_BLACKWELL;
const int nsm = ggml_cuda_info().devices[ggml_cuda_get_device()].nsm;
const int64_t max_ncols = blackwell ? MMVF_MAX_BATCH_SIZE : 5;
const int64_t max_k = blackwell ? 6144 : 4096;
if (ne11 <= max_ncols && src0_ne[1]*ne11*256 <= nsm*std::min(src0_ne[0], max_k)) {
return true;
}
}
if (ampere_mma_available(cc)) {
return src0_small && ne11 == 1;
}
@@ -973,8 +882,6 @@ bool ggml_cuda_should_use_mmvf(const ggml_tensor * src0, int cc, int warp_size,
return ne11 <= MMVF_MAX_BATCH_SIZE;
}
return ne11 <= MMVF_MAX_BATCH_SIZE;
case GGML_TYPE_F8_E4M3:
return ne11 <= MMVF_MAX_BATCH_SIZE;
default:
return false;
}
+1 -1
View File
@@ -11,4 +11,4 @@ void ggml_cuda_op_mul_mat_vec_f(
const char * src1_ddq_i, float * dst_dd_i, const int64_t row_low, const int64_t row_high, const int64_t src1_ncols,
const int64_t src1_padded_row_size, cudaStream_t stream);
bool ggml_cuda_should_use_mmvf(const ggml_tensor * src0, int cc, int warp_size, int64_t ne11);
bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, int warp_size, const int64_t * src0_ne, const size_t * src0_nb, int64_t ne11);
+1 -1
View File
@@ -1555,7 +1555,7 @@ void ggml_cuda_mul_mat_vec_q(
const int64_t s11 = src1->nb[1] / ts_src1;
const int64_t s12 = src1->nb[2] / ts_src1;
const int64_t s13 = src1->nb[3] / ts_src1;
quantize_row_q8_1_cuda(src1_d, src1_q8_1.get(), ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream);
quantize_row_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream);
}
const int64_t s01 = src0->nb[1] / ts_src0;
+3 -1
View File
@@ -557,9 +557,10 @@ static __global__ void quantize_mmq_q8_1(
}
void quantize_row_q8_1_cuda(
const float * x, void * vy,
const float * x, const int32_t * ids, void * vy, const ggml_type type_src0,
const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) {
GGML_ASSERT(!ids);
GGML_ASSERT(ne0 % QK8_1 == 0);
const uint3 ne2_fastdiv = init_fastdiv_values(ne2);
@@ -569,6 +570,7 @@ void quantize_row_q8_1_cuda(
const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE, 1, 1);
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(num_blocks, block_size, 0, stream);
ggml_cuda_kernel_launch(quantize_q8_1, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv);
GGML_UNUSED(type_src0);
}
void quantize_mmq_q8_1_cuda(
+7 -2
View File
@@ -11,9 +11,14 @@
static_assert(MATRIX_ROW_PADDING % CUDA_QUANTIZE_BLOCK_SIZE == 0, "Risk of out-of-bounds access.");
static_assert(MATRIX_ROW_PADDING % (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ) == 0, "Risk of out-of-bounds access.");
typedef void (*quantize_cuda_t)(
const float * x, const int32_t * ids, void * vy,
ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03,
int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream);
void quantize_row_q8_1_cuda(
const float * x, void * vy,
int64_t ne00, int64_t s01, int64_t s02, int64_t s03,
const float * x, const int32_t * ids, void * vy,
ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03,
int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream);
void quantize_mmq_q8_1_cuda(
-10
View File
@@ -257,16 +257,6 @@ static void set_rows_cuda(ggml_backend_cuda_context & ctx, const ggml_tensor * s
nb1, nb2, nb3,
stream
);
} else if (dst->type == GGML_TYPE_F8_E4M3) {
set_rows_cuda(
src0_d, src1_d, (ggml_fp8_e4m3_t *) dst->data,
ne00, ne01, ne02, ne03,
ne10, ne11, ne12, ne13,
nb01, nb02, nb03,
nb10, nb11, nb12,
nb1, nb2, nb3,
stream
);
} else if (dst->type == GGML_TYPE_Q4_0) {
set_rows_cuda_quant<idx_t, block_q4_0, QK4_0, quantize_f32_q4_0_block>(
src0_d, src1_d, (block_q4_0*)dst->data,
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 1, 8);
DECL_FATTN_MMA_F16_CASE(64, 64, 1, 8);
DECL_FATTN_MMA_F16_CASE(72, 72, 1, 8);
DECL_FATTN_MMA_F16_CASE(80, 80, 1, 8);
DECL_FATTN_MMA_F16_CASE(96, 96, 1, 8);
DECL_FATTN_MMA_F16_CASE(112, 112, 1, 8);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 16, 1);
DECL_FATTN_MMA_F16_CASE(64, 64, 16, 1);
DECL_FATTN_MMA_F16_CASE(72, 72, 16, 1);
DECL_FATTN_MMA_F16_CASE(80, 80, 16, 1);
DECL_FATTN_MMA_F16_CASE(96, 96, 16, 1);
DECL_FATTN_MMA_F16_CASE(112, 112, 16, 1);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 16, 2);
DECL_FATTN_MMA_F16_CASE(64, 64, 16, 2);
DECL_FATTN_MMA_F16_CASE(72, 72, 16, 2);
DECL_FATTN_MMA_F16_CASE(80, 80, 16, 2);
DECL_FATTN_MMA_F16_CASE(96, 96, 16, 2);
DECL_FATTN_MMA_F16_CASE(112, 112, 16, 2);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 16, 4);
DECL_FATTN_MMA_F16_CASE(64, 64, 16, 4);
DECL_FATTN_MMA_F16_CASE(72, 72, 16, 4);
DECL_FATTN_MMA_F16_CASE(80, 80, 16, 4);
DECL_FATTN_MMA_F16_CASE(96, 96, 16, 4);
DECL_FATTN_MMA_F16_CASE(112, 112, 16, 4);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 2, 4);
DECL_FATTN_MMA_F16_CASE(64, 64, 2, 4);
DECL_FATTN_MMA_F16_CASE(72, 72, 2, 4);
DECL_FATTN_MMA_F16_CASE(80, 80, 2, 4);
DECL_FATTN_MMA_F16_CASE(96, 96, 2, 4);
DECL_FATTN_MMA_F16_CASE(112, 112, 2, 4);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 2, 8);
DECL_FATTN_MMA_F16_CASE(64, 64, 2, 8);
DECL_FATTN_MMA_F16_CASE(72, 72, 2, 8);
DECL_FATTN_MMA_F16_CASE(80, 80, 2, 8);
DECL_FATTN_MMA_F16_CASE(96, 96, 2, 8);
DECL_FATTN_MMA_F16_CASE(112, 112, 2, 8);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 32, 1);
DECL_FATTN_MMA_F16_CASE(64, 64, 32, 1);
DECL_FATTN_MMA_F16_CASE(72, 72, 32, 1);
DECL_FATTN_MMA_F16_CASE(80, 80, 32, 1);
DECL_FATTN_MMA_F16_CASE(96, 96, 32, 1);
DECL_FATTN_MMA_F16_CASE(112, 112, 32, 1);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 32, 2);
DECL_FATTN_MMA_F16_CASE(64, 64, 32, 2);
DECL_FATTN_MMA_F16_CASE(72, 72, 32, 2);
DECL_FATTN_MMA_F16_CASE(80, 80, 32, 2);
DECL_FATTN_MMA_F16_CASE(96, 96, 32, 2);
DECL_FATTN_MMA_F16_CASE(112, 112, 32, 2);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 4, 2);
DECL_FATTN_MMA_F16_CASE(64, 64, 4, 2);
DECL_FATTN_MMA_F16_CASE(72, 72, 4, 2);
DECL_FATTN_MMA_F16_CASE(80, 80, 4, 2);
DECL_FATTN_MMA_F16_CASE(96, 96, 4, 2);
DECL_FATTN_MMA_F16_CASE(112, 112, 4, 2);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 4, 4);
DECL_FATTN_MMA_F16_CASE(64, 64, 4, 4);
DECL_FATTN_MMA_F16_CASE(72, 72, 4, 4);
DECL_FATTN_MMA_F16_CASE(80, 80, 4, 4);
DECL_FATTN_MMA_F16_CASE(96, 96, 4, 4);
DECL_FATTN_MMA_F16_CASE(112, 112, 4, 4);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 4, 8);
DECL_FATTN_MMA_F16_CASE(64, 64, 4, 8);
DECL_FATTN_MMA_F16_CASE(72, 72, 4, 8);
DECL_FATTN_MMA_F16_CASE(80, 80, 4, 8);
DECL_FATTN_MMA_F16_CASE(96, 96, 4, 8);
DECL_FATTN_MMA_F16_CASE(112, 112, 4, 8);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 64, 1);
DECL_FATTN_MMA_F16_CASE(64, 64, 64, 1);
DECL_FATTN_MMA_F16_CASE(72, 72, 64, 1);
DECL_FATTN_MMA_F16_CASE(80, 80, 64, 1);
DECL_FATTN_MMA_F16_CASE(96, 96, 64, 1);
DECL_FATTN_MMA_F16_CASE(112, 112, 64, 1);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 8, 1);
DECL_FATTN_MMA_F16_CASE(64, 64, 8, 1);
DECL_FATTN_MMA_F16_CASE(72, 72, 8, 1);
DECL_FATTN_MMA_F16_CASE(80, 80, 8, 1);
DECL_FATTN_MMA_F16_CASE(96, 96, 8, 1);
DECL_FATTN_MMA_F16_CASE(112, 112, 8, 1);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 8, 2);
DECL_FATTN_MMA_F16_CASE(64, 64, 8, 2);
DECL_FATTN_MMA_F16_CASE(72, 72, 8, 2);
DECL_FATTN_MMA_F16_CASE(80, 80, 8, 2);
DECL_FATTN_MMA_F16_CASE(96, 96, 8, 2);
DECL_FATTN_MMA_F16_CASE(112, 112, 8, 2);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 8, 4);
DECL_FATTN_MMA_F16_CASE(64, 64, 8, 4);
DECL_FATTN_MMA_F16_CASE(72, 72, 8, 4);
DECL_FATTN_MMA_F16_CASE(80, 80, 8, 4);
DECL_FATTN_MMA_F16_CASE(96, 96, 8, 4);
DECL_FATTN_MMA_F16_CASE(112, 112, 8, 4);
@@ -2,9 +2,7 @@
#include "../fattn-mma-f16.cuh"
DECL_FATTN_MMA_F16_CASE(40, 40, 8, 8);
DECL_FATTN_MMA_F16_CASE(64, 64, 8, 8);
DECL_FATTN_MMA_F16_CASE(72, 72, 8, 8);
DECL_FATTN_MMA_F16_CASE(80, 80, 8, 8);
DECL_FATTN_MMA_F16_CASE(96, 96, 8, 8);
DECL_FATTN_MMA_F16_CASE(112, 112, 8, 8);

Some files were not shown because too many files have changed in this diff Show More