diff --git a/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw_option_defs.inc b/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw_option_defs.inc index 081be721abc4b8..b125e39b3ee3c5 100644 --- a/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw_option_defs.inc +++ b/src/plugins/intel_npu/src/al/include/intel_npu/config/npuw_option_defs.inc @@ -90,3 +90,6 @@ INTEL_NPU_NPUW_SIMPLE_OPT(NPUW_TEXT_EMBED, bool, false, ov::intel_npu::npuw::tex INTEL_NPU_NPUW_SIMPLE_OPT(NPUW_KOKORO, bool, false, ov::intel_npu::npuw::kokoro, enabled, "NPUW_KOKORO", KOKORO, EXPOSED, UNCACHED, ALL) INTEL_NPU_NPUW_SIMPLE_OPT(NPUW_KOKORO_BLOCK_SIZE, uint64_t, 200, ov::intel_npu::npuw::kokoro, block_size, "NPUW_KOKORO_BLOCK_SIZE", KOKORO, EXPOSED, UNCACHED, ALL) INTEL_NPU_NPUW_SIMPLE_OPT(NPUW_KOKORO_OVERLAP_SIZE, uint64_t, 20, ov::intel_npu::npuw::kokoro, overlap_size, "NPUW_KOKORO_OVERLAP_SIZE", KOKORO, EXPOSED, UNCACHED, ALL) +INTEL_NPU_NPUW_SIMPLE_OPT(NPUW_LLM_ASYM_VOCAB_AS_INPUT, bool, true, ov::intel_npu::npuw::llm, asym_vocab_as_input, "NPUW_LLM_ASYM_VOCAB_AS_INPUT", LLM, EXPOSED, CACHED, ALL) +INTEL_NPU_NPUW_SIMPLE_OPT(NPUW_LLM_ASYM_I8_VOCAB_AS_INPUT, bool, true, ov::intel_npu::npuw::llm, asym_i8_vocab_as_input, "NPUW_LLM_ASYM_I8_VOCAB_AS_INPUT", LLM, EXPOSED, CACHED, ALL) +INTEL_NPU_NPUW_SIMPLE_OPT(NPUW_LLM_MATMUL_FIRST_VOCAB, bool, true, ov::intel_npu::npuw::llm, matmul_first_vocab, "NPUW_LLM_MATMUL_FIRST_VOCAB", LLM, EXPOSED, CACHED, ALL) diff --git a/src/plugins/intel_npu/src/plugin/CMakeLists.txt b/src/plugins/intel_npu/src/plugin/CMakeLists.txt index 1e9e219aebe13a..900df9244884d6 100644 --- a/src/plugins/intel_npu/src/plugin/CMakeLists.txt +++ b/src/plugins/intel_npu/src/plugin/CMakeLists.txt @@ -50,7 +50,7 @@ cross_compiled_file(${TARGET_NAME} ARCH AVX2 ANY npuw/util_xarch.cpp API npuw/util_xarch.hpp - NAME unpack_i4i8 unpack_u4i8 unpack_i4f16 unpack_i4f16_scale unpack_i4f16_z unpack_u4f16 unpack_u4f16_scale_zp unpack_u4f16_asymm_zp unpack_u4f16_z unpack_u4f32 unpack_i8f16 unpack_i8f16_scale unpack_u8f16 to_f16 copy_row_as_column transpose_i4 transpose_f16 transpose_f32 unpack_f8f16_scale + NAME unpack_i4i8 unpack_u4i8 unpack_i4f16 unpack_i4f16_scale unpack_i4f16_z unpack_u4f16 unpack_u4f16_scale_zp unpack_u4f16_asymm_zp unpack_u4f16_z unpack_u4f32 unpack_i8f16 unpack_i8f16_scale unpack_i8f16_scale_zp unpack_u8f16 to_f16 copy_row_as_column transpose_i4 transpose_f16 transpose_f32 unpack_f8f16_scale NAMESPACE ov::npuw::util::XARCH ) diff --git a/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.cpp b/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.cpp index 3d3c137e327fd7..dd07a0087557bd 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.cpp @@ -37,8 +37,11 @@ Const::Const(const std::shared_ptr& n) : m_node(n) { m_offset = weightless_cache_attr->second.as().bin_offset; } else { // See the comment in serialize() for more details - LOG_WARN("Some pattern introduced a new Constant node not present in the original weights file. We need to " - "keep it in case export occurs. This will increase memory consumption."); + LOG_WARN("Some pattern introduced a new Constant node, " + << m_node + << ", not present in the " + "original weights file. We need to keep it in case export occurs. This will increase " + "memory consumption."); m_copied_if_not_in_model = ov::npuw::util::copy_tensor_from_const(m_node); } } @@ -57,9 +60,9 @@ bool Const::operator==(const Const& other) const { m_cached_ptr == other.m_cached_ptr); } -ov::Tensor Const::eval() const { +ov::Tensor Const::eval_view() const { if (m_node) { - return ov::npuw::util::copy_tensor_from_const(m_node); + return ov::npuw::util::tensor_from_const(m_node); } // Weightless import case. Mmmap CPU weight on demand to avoid allocating all weights at once. @@ -83,6 +86,15 @@ ov::Tensor Const::eval() const { return m_read_from_bin; } +ov::Tensor Const::eval() const { + if (m_node) { + return ov::npuw::util::copy_tensor_from_const(m_node); + } + // The import branches of eval_view() don't copy - the bank takes + // ownership on its side when required + return eval_view(); +} + LazyTensor::Meta Const::eval_meta() const { if (m_node) { return {m_node->get_shape(), m_node->get_element_type()}; @@ -356,6 +368,54 @@ void Gather::detach() { w.detach(); } +std::size_t Sub128::hash() const { + std::size_t seed = std::hash()(7u) + 0x9e3779b9; + seed ^= tensor.get_hash() + 0x9e3779b9; + return seed; +} + +bool Sub128::operator==(const Sub128& other) const { + return tensor == other.tensor; +} + +ov::Tensor Sub128::eval() const { + const auto trs = tensor.get_transformations(); + + ov::Tensor src; + if (trs.size() == 1 && std::holds_alternative(trs.front())) { + // Fused path: read straight through a zero-copy view of the source, + // skipping the intermediate copy Const::eval() would make. The view is + // only ever READ here, so this is correct for all Const flavors, + // including the deserialized ones (read-only mmap / cached bin tensor) + src = std::get(trs.front()).eval_view(); + } else { + src = tensor.eval(); + } + + const auto src_type = src.get_element_type(); + NPUW_ASSERT(src_type == ov::element::u8 || src_type == ov::element::i8); + + ov::Tensor dst(ov::element::i8, src.get_shape()); + const auto* s = static_cast(src.data()); + auto* d = dst.data(); + for (std::size_t i = 0, n = src.get_size(); i < n; ++i) { + d[i] = static_cast(static_cast(s[i]) - 128); + } + return dst; +} + +LazyTensor::Meta Sub128::eval_meta() const { + return {tensor.eval_meta().shape, ov::element::i8}; +} + +void Sub128::read_weight(const ov::npuw::s11n::WeightsContext& ctx) { + tensor.read_weight(ctx); +} + +void Sub128::detach() { + tensor.detach(); +} + } // namespace op // Stable, permanently assigned op-type IDs. @@ -368,6 +428,7 @@ enum class TransformType : std::uint16_t { PERMUTE = 4, CONVERT = 5, GATHER = 6, + SUB128 = 7, }; struct LazyTensorImpl { @@ -430,6 +491,10 @@ ov::npuw::weights::TransformType get_transform_type(const ov::npuw::weights::op: return ov::npuw::weights::TransformType::GATHER; } +ov::npuw::weights::TransformType get_transform_type(const ov::npuw::weights::op::Sub128&) { + return ov::npuw::weights::TransformType::SUB128; +} + } // namespace namespace ov { @@ -500,6 +565,10 @@ void Gather::serialize(ov::npuw::orc::Stream& stream) { } } +void Sub128::serialize(ov::npuw::orc::Stream& stream) { + stream & tensor; +} + } // namespace op void LazyTensorImpl::serialize(ov::npuw::orc::Stream& stream) { @@ -537,6 +606,9 @@ void LazyTensorImpl::serialize(ov::npuw::orc::Stream& stream) { case TransformType::GATHER: m_transform.emplace(ov::npuw::orc::load_versioned_payload(section)); break; + case TransformType::SUB128: + m_transform.emplace(ov::npuw::orc::load_versioned_payload(section)); + break; default: OPENVINO_THROW("ORC LazyTensor: unknown op_type ", section.type, " — please upgrade NPUW"); break; @@ -649,6 +721,10 @@ void LazyTensorImpl::get_transformations(std::vector& vec auto next_tr = op.w.get_transformations(); vec.insert(vec.end(), next_tr.begin(), next_tr.end()); }, + [&vec](const op::Sub128& op) { + auto next_tr = op.tensor.get_transformations(); + vec.insert(vec.end(), next_tr.begin(), next_tr.end()); + }, }, m_transform); } @@ -688,6 +764,12 @@ LazyTensor LazyTensor::convert(const ov::element::Type& type) { return new_lt; } +LazyTensor LazyTensor::sub128() { + LazyTensor new_lt; + new_lt.m_impl = std::make_shared(op::Sub128(*this)); + return new_lt; +} + bool LazyTensor::operator==(const LazyTensor& other) const { if (!m_impl && !other.m_impl) { return true; diff --git a/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.hpp b/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.hpp index 42c5aec7251835..e1f2257ff0fb39 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.hpp +++ b/src/plugins/intel_npu/src/plugin/npuw/lazy_tensor.hpp @@ -27,6 +27,7 @@ class Unpack; class Permute; class Convert; class Gather; +class Sub128; } // namespace op class LazyTensor { @@ -41,7 +42,8 @@ class LazyTensor { ov::npuw::weights::op::Unpack, ov::npuw::weights::op::Permute, ov::npuw::weights::op::Convert, - ov::npuw::weights::op::Gather>; + ov::npuw::weights::op::Gather, + ov::npuw::weights::op::Sub128>; LazyTensor() = default; LazyTensor(const std::shared_ptr& const_ptr); @@ -58,6 +60,7 @@ class LazyTensor { LazyTensor permute(const std::vector& axes); LazyTensor convert(const ov::element::Type& type); + LazyTensor sub128(); // byte-wise -128 zero-point shift (u8/i8 -> symmetric i8) bool operator==(const LazyTensor& other) const; bool operator!=(const LazyTensor& other) const; @@ -97,6 +100,7 @@ class Const { std::size_t hash() const; bool operator==(const Const& other) const; ov::Tensor eval() const; + ov::Tensor eval_view() const; LazyTensor::Meta eval_meta() const; void read_weight(const ov::npuw::s11n::WeightsContext& ctx); void detach(); @@ -240,6 +244,31 @@ class Gather { ov::element::Type dst_type; ov::Shape dst_shape; }; + +class Sub128 { + friend struct ov::npuw::weights::LazyTensorImpl; + +public: + static constexpr std::uint16_t kVersion = 0u; + + // rt_info marker: a graph-level pass puts this key on a Constant to request + // the -128 shift to be applied when the weight is lifted into a closure + static constexpr const char* rt_key = "npuw::sub128"; + + Sub128() = default; + explicit Sub128(const LazyTensor& _tensor) : tensor(_tensor) {} + + std::size_t hash() const; + bool operator==(const Sub128& other) const; + ov::Tensor eval() const; + LazyTensor::Meta eval_meta() const; + void read_weight(const ov::npuw::s11n::WeightsContext& ctx); + void detach(); + void serialize(ov::npuw::orc::Stream& stream); + +private: + LazyTensor tensor; +}; } // namespace op } // namespace weights diff --git a/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp b/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp index 7a3f938b18a596..45b483bb1fdcab 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/llm_compiled_model.cpp @@ -63,6 +63,157 @@ bool is_aligned_to(T value, T alignment) { } // namespace +class ConvertVocabAsymU8ToI8 : public ov::pass::MatcherPass { +public: + OPENVINO_MATCHER_PASS_RTTI("ov::npuw::ConvertVocabU8ToI8"); + explicit ConvertVocabAsymU8ToI8() { + auto qweight = opp::wrap_type(); + auto qcoeff = opp::wrap_type(); + auto qzerop = opp::wrap_type(); + auto qcvtw = opp::wrap_type({qweight}); + auto qcvtz = opp::wrap_type({qzerop}); + auto qsub = opp::wrap_type({qcvtw, qcvtz}); + auto qmuls = opp::wrap_type({qsub, qcoeff}); + auto qcvtm = opp::wrap_type({qmuls}); + auto qmmi = opp::any_input(); + auto qmm = opp::wrap_type({qmmi, qcvtm}); + auto qres = opp::wrap_type({qmm}); + + auto callback = [=](opp::Matcher& m) { + auto& node_to_output = m.get_pattern_value_map(); + + auto matched_qweight = node_to_output.at(qweight).get_node_shared_ptr(); + if (matched_qweight->get_element_type() != ov::element::u8) { + return false; + } + if (matched_qweight->get_shape().size() != 2) { + return false; + } + auto matched_qcoeff = node_to_output.at(qcoeff).get_node_shared_ptr(); + auto matched_qzerop = node_to_output.at(qzerop).get_node_shared_ptr(); + auto qcoeff_shape = std::static_pointer_cast(matched_qcoeff)->get_shape(); + auto matched_matmul = + std::static_pointer_cast(node_to_output.at(qmm).get_node_shared_ptr()); + + if (qcoeff_shape.size() == 2 && qcoeff_shape[1] == 1 && !matched_matmul->get_transpose_a() && + matched_matmul->get_transpose_b()) { + auto matched_qweight_const = std::static_pointer_cast(matched_qweight); + auto matched_qzerop_const = std::static_pointer_cast(matched_qzerop); + + auto reinterpret_u8_as_i8 = [](const std::shared_ptr& src) { + OPENVINO_ASSERT(src->get_element_type() == ov::element::u8); + auto dst = std::make_shared( + ov::element::i8, + src->get_shape(), + src->get_data_ptr(), + src); // <-- 'so': source node kept alive => no dangling, no copy + dst->set_friendly_name(src->get_friendly_name()); + // FIXME: This copies weightless attribute to not preserve a constant as a new for weightless + // import. NOTE: ov::copy_runtime_info(src, dst) doesn't lead to the desired behavior, failing the + // LazyTensor to find Weightless attribute! + dst->get_rt_info() = src->get_rt_info(); + return dst; + }; + + // To not mmap and allocate vocab memory here, its shifting will be deferred to the LazyTensor unpacking + // stage. + auto i8_qweight_constant = reinterpret_u8_as_i8(matched_qweight_const); + // Inform partitioning, that this Const is special and needs to be unpacked with -128 shift applied. + i8_qweight_constant->get_rt_info()[ov::npuw::weights::op::Sub128::rt_key] = true; + ov::replace_node(matched_qweight_const, i8_qweight_constant); + auto i8_qzerop_constant = reinterpret_u8_as_i8(matched_qzerop_const); + // Inform partitioning, that this Const is special and needs to be unpacked with -128 shift applied. + i8_qzerop_constant->get_rt_info()[ov::npuw::weights::op::Sub128::rt_key] = true; + ov::replace_node(matched_qzerop_const, i8_qzerop_constant); + return true; + } + return false; + }; + register_matcher(std::make_shared(qres, "ConvertVocabAsymU8ToI8"), std::move(callback)); + } +}; + +// Rewrites the LM head DQ chain so the heavy MatMul consumes raw u8/i8 weight and the +// per-row (weight - zerop) * scale dequant is applied after it: +// logits = (x @ W^T) * s - sum(x, axis=-1) * (z * s) +class MatMulFirstAsymVocab : public ov::pass::MatcherPass { +public: + OPENVINO_MATCHER_PASS_RTTI("ov::npuw::MatMulFirstAsymVocab"); + explicit MatMulFirstAsymVocab() { + auto qweight = opp::wrap_type(); + auto qcoeff = opp::wrap_type(); + auto qzerop = opp::wrap_type(); + auto qcvtw = opp::wrap_type({qweight}); + auto qcvtz = opp::wrap_type({qzerop}); + auto qsub = opp::wrap_type({qcvtw, qcvtz}); + auto qmuls = opp::wrap_type({qsub, qcoeff}); + auto qcvtm = opp::wrap_type({qmuls}); + auto qmmi = opp::any_input(); + auto qmm = opp::wrap_type({qmmi, qcvtm}); + auto qres = opp::wrap_type({qmm}); + + auto callback = [=](opp::Matcher& m) { + auto& node_to_output = m.get_pattern_value_map(); + auto matched_qweight = node_to_output.at(qweight).get_node_shared_ptr(); + if (matched_qweight->get_element_type() != ov::element::u8 && + matched_qweight->get_element_type() != ov::element::i8) { + return false; + } + if (matched_qweight->get_shape().size() != 2) { + return false; + } + auto matched_qcoeff = node_to_output.at(qcoeff).get_node_shared_ptr(); + auto matched_qzerop = node_to_output.at(qzerop).get_node_shared_ptr(); + auto qcoeff_shape = matched_qcoeff->get_shape(); + auto matched_matmul = + std::static_pointer_cast(node_to_output.at(qmm).get_node_shared_ptr()); + auto matched_result = + std::static_pointer_cast(node_to_output.at(qres).get_node_shared_ptr()); + + if (qcoeff_shape.size() != 2 || qcoeff_shape[1] != 1 || matched_matmul->get_transpose_a() || + !matched_matmul->get_transpose_b()) { + return false; + } + + auto hidden = node_to_output.at(qmmi); + const auto vocab_size = static_cast(matched_qweight->get_shape()[0]); + + // // (x @ W^T) with raw u8 weight cast to f32; compiler may fuse the Convert into MatMul. + auto w_f32 = std::make_shared(matched_qweight, ov::element::f32); + auto new_matmul = std::make_shared(hidden, w_f32, false, true); + + auto s_f32 = std::make_shared(matched_qcoeff, ov::element::f32); + auto s_reshape = + ov::op::v0::Constant::create(ov::element::i64, ov::Shape{3}, std::vector{1, 1, vocab_size}); + auto scale_reshaped = std::make_shared(s_f32, s_reshape, false); + auto scaled = std::make_shared(new_matmul, scale_reshaped); + scaled->set_friendly_name("scale_after_matmul"); + + auto z_f32 = std::make_shared(matched_qzerop, ov::element::f32); + auto zp_scale = std::make_shared(z_f32, s_f32); + zp_scale->set_friendly_name("zero_point_scale"); + auto zp_reshape = + ov::op::v0::Constant::create(ov::element::i64, ov::Shape{3}, std::vector{1, 1, vocab_size}); + auto zp_reshaped = std::make_shared(zp_scale, zp_reshape, false); + + auto reduce_axis = ov::op::v0::Constant::create(ov::element::i64, ov::Shape{1}, std::vector{-1}); + auto sum_h = std::make_shared(hidden, reduce_axis, true); + sum_h->set_friendly_name("reduce_sum_h"); + auto zp_correction = std::make_shared(sum_h, zp_reshaped); + zp_correction->set_friendly_name("zp_scale_after_matmul"); + + auto logits = std::make_shared(scaled, zp_correction); + logits->set_friendly_name("subtract_after_matmul"); + // Preserve tensor names (e.g. "logits") from the old MatMul output so consumers keep working. + logits->output(0).set_names(matched_matmul->output(0).get_names()); + + matched_result->input(0).replace_source_output(logits->output(0)); + return true; + }; + register_matcher(std::make_shared(qres, "MatMulFirstAsymVocab"), std::move(callback)); + } +}; + class CutLMHead : public ov::pass::MatcherPass { public: OPENVINO_MATCHER_PASS_RTTI("ov::npuw::CutLMHead"); @@ -148,6 +299,22 @@ class CutLMHead : public ov::pass::MatcherPass { }; namespace { +bool convert_vocab_to_i8(const std::shared_ptr& model) { + ov::pass::GraphRewrite rewr; + rewr.add_matcher(); + auto id_converted = rewr.run_on_model(model); + model->validate_nodes_and_infer_types(); + return id_converted; +} + +bool apply_matmul_first_vocab(const std::shared_ptr& model) { + ov::pass::GraphRewrite rewr; + rewr.add_matcher(); + auto ran = rewr.run_on_model(model); + model->validate_nodes_and_infer_types(); + return ran; +} + std::shared_ptr cut_lm_head(const std::shared_ptr& model) { ov::pass::GraphRewrite rewr; std::shared_ptr lm_head_model = nullptr; @@ -260,6 +427,14 @@ std::optional extract_npu_descriptor(const std::shared_ptr +std::optional get_option(const ov::AnyMap& config, const std::string& option_name) { + if (auto it = config.find(option_name); it != config.end()) { + return std::make_optional(it->second.as()); + } + return std::nullopt; +} + std::optional pop_option(ov::AnyMap& config, const std::string& option_name) { if (auto it = config.find(option_name); it != config.end()) { std::optional found = std::make_optional(it->second); @@ -833,7 +1008,15 @@ ov::npuw::LLMCompiledModel::LLMCompiledModel(const std::shared_ptr& m ov::npuw::ReplaceDeepstackScatterWithAdd().run_on_model(kvcache_model); + if (m_cfg.get<::intel_npu::NPUW_LLM_ASYM_I8_VOCAB_AS_INPUT>()) { + NPUW_ASSERT(convert_vocab_to_i8(kvcache_model)); + } auto lm_head_model = check_and_cut_lm_head(kvcache_model, m_cfg); + if (lm_head_model) { + if (m_cfg.get<::intel_npu::NPUW_LLM_MATMUL_FIRST_VOCAB>()) { + NPUW_ASSERT(apply_matmul_first_vocab(lm_head_model)); + } + } if (!m_is_whisper) { LOG_DEBUG("Try patch sliding window attention mask (Phi-3, Gemma-2, Gemma-3, Gemma-4), if it exists."); @@ -1213,6 +1396,11 @@ ov::npuw::LLMCompiledModel::LLMCompiledModel(const std::shared_ptr& m apply_weights_bank_name(lm_head_config, weights_bank_name); + auto vocab_as_input = get_option(other_props, std::string("NPUW_ASYM_VOCAB_AS_INPUT")); + if (m_cfg.get<::intel_npu::NPUW_LLM_ASYM_I8_VOCAB_AS_INPUT>() || + m_cfg.get<::intel_npu::NPUW_LLM_ASYM_VOCAB_AS_INPUT>()) { + lm_head_config["NPUW_HOST_GATHER"] = "NO"; + } m_lm_head_compiled = m_compiled_model_factory(lm_head_model, plugin, lm_head_config); NPUW_ASSERT(m_lm_head_compiled); } diff --git a/src/plugins/intel_npu/src/plugin/npuw/partitioning/partitioning.cpp b/src/plugins/intel_npu/src/plugin/npuw/partitioning/partitioning.cpp index 6457727cc150b5..7f44d47c81f8af 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/partitioning/partitioning.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/partitioning/partitioning.cpp @@ -1752,6 +1752,18 @@ void Partitioner::matchResults(const std::string& func_name) { LOG_VERB("Done"); } +namespace { +ov::npuw::weights::LazyTensor put_to_closure(const std::shared_ptr& input_node) { + auto const_node = std::static_pointer_cast(input_node); + ov::npuw::weights::LazyTensor lt(const_node); + if (const_node->get_rt_info().count(ov::npuw::weights::op::Sub128::rt_key) > 0) { + LOG_DEBUG("Sub128 marker found on " << const_node->get_friendly_name() << " - applying to the LazyTensor"); + lt = lt.sub128(); + } + return lt; +} +} // anonymous namespace + void Partitioner::createFunction(FunctionPipeline& func_ggg) { using namespace ov::npuw::weights; @@ -1814,8 +1826,15 @@ void Partitioner::createFunction(FunctionPipeline& func_ggg) { new_param_idx++; LOG_DEBUG("Register " << prod_output << " in the function closure"); - funcall._lazy_closure.push_back( - LazyTensor(std::static_pointer_cast(input_node))); // (n)/1/i/c + funcall._lazy_closure.push_back(put_to_closure(input_node)); // (n)/1/i/c + } else if (ov::op::util::is_constant(input_node) && + input_node->get_rt_info().count(ov::npuw::weights::op::Sub128::rt_key) > 0) { + // A Sub128-marked Constant landed in consts_to_keep and stays inline + // in the function body - the shift would be silently skipped there, + // producing wrong numerics. Fail loudly instead. + OPENVINO_THROW("NPUW: Sub128-marked Constant ", + input_node->get_friendly_name(), + " is kept in the function body and won't be transformed"); } else if (ov::op::util::is_parameter(input_node)) { LOG_DEBUG("Handling a Parameter input " << prod_output); LOG_BLOCK(); @@ -1985,8 +2004,7 @@ void Partitioner::matchRepeatedSubgraphs(const std::string& func_name) { std::make_pair(proto_layer_name, input_desc.get_index())); // (t)/1/b LOG_DEBUG("Register " << prod_output << " in the function closure[" << param_idx << "] (via prototype " << proto_layer_name << ")"); - funcall._lazy_closure[param_idx - function._param_offset] = - LazyTensor(std::static_pointer_cast(input_node)); // (t)/1/c + funcall._lazy_closure[param_idx - function._param_offset] = put_to_closure(input_node); // (t)/1/c } } // for (inputs) } // for(nodes) diff --git a/src/plugins/intel_npu/src/plugin/npuw/partitioning/patterns/opt.cpp b/src/plugins/intel_npu/src/plugin/npuw/partitioning/patterns/opt.cpp index d39851ddd580e8..615e1f9bf50417 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/partitioning/patterns/opt.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/partitioning/patterns/opt.cpp @@ -1745,8 +1745,9 @@ DQUnpackDictMatMulCWu::DQUnpackDictMatMulCWu(Context::Ref ctx) { auto qcoeff_shape = matched_qcoeff->output(0).get_shape(); - if (ov::element::u8 == matched_qweight->get_element_type() && qcoeff_shape[1] == 1 && - !matched_matmul->get_transpose_a() && matched_matmul->get_transpose_b()) { + if ((ov::element::u8 == matched_qweight->get_element_type() || + ov::element::i8 == matched_qweight->get_element_type()) && + qcoeff_shape[1] == 1 && !matched_matmul->get_transpose_a() && matched_matmul->get_transpose_b()) { auto new_cvt_a = std::make_shared(matched_mmi, ov::element::f16); auto new_wi = ctx.get().unpack(matched_qweight, matched_qzerop, matched_qcoeff, ov::element::f16); @@ -1977,7 +1978,9 @@ PreserveConstDictMatMulAsymm::PreserveConstDictMatMulAsymm(Context::Ref ctx, // Pre-transposed layout: weight [IC, OC], scale [1, OC], transpose_b=false const bool pretransposed_layout = qcoeff_shape.size() == 2 && qcoeff_shape[0] == 1 && !matched_matmul->get_transpose_a() && !matched_matmul->get_transpose_b(); - if (ov::element::u8 == matched_qweight->get_element_type() && (standard_layout || pretransposed_layout)) { + if ((ov::element::u8 == matched_qweight->get_element_type() || + ov::element::i8 == matched_qweight->get_element_type()) && + (standard_layout || pretransposed_layout)) { to_keep.get().push_back(matched_qweight); to_keep.get().push_back(matched_qzerop); to_keep.get().push_back(matched_qcoeff); diff --git a/src/plugins/intel_npu/src/plugin/npuw/util.cpp b/src/plugins/intel_npu/src/plugin/npuw/util.cpp index 07fd8249de09cd..a41aa54d007b00 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/util.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/util.cpp @@ -295,6 +295,10 @@ void ov::npuw::util::unpack(const ov::SoPtr& from, NPUW_ASSERT(type_zerop == ov::element::u8); NPUW_ASSERT(type_scale == ov::element::f16); NPUW_ASSERT(type_to == ov::element::f16); + } else if (type_from == ov::element::i8) { + NPUW_ASSERT(type_zerop == ov::element::i8); + NPUW_ASSERT(type_scale == ov::element::f16); + NPUW_ASSERT(type_to == ov::element::f16); } else { NPUW_ASSERT(false && "Unsupported combination"); } @@ -384,6 +388,52 @@ void ov::npuw::util::unpack(const ov::SoPtr& from, } else { NPUW_ASSERT(false); } + } else if (type_from == ov::element::i8) { + if (scale_shape.size() == 3 && scale_shape[1] == 1 && scale_shape[2] == 1) { + // Special case for broadcasting vocab by 2 dimensions + // FIXME: all this logic probably should be in some specific unpack or another util function + const auto& from_strides = from->get_strides(); + const auto& zerop_strides = zerop->get_strides(); + const auto& scale_strides = scale->get_strides(); + ov::Tensor wraped_from(from->get_element_type(), + ov::Shape{from_shape[0], from_shape[1] * from_shape[2]}, + from->data(), + ov::Strides{from_strides[0], from_strides[2]}); + ov::Tensor wraped_zerop(zerop->get_element_type(), + ov::Shape{zerop_shape[0], zerop_shape[1] * zerop_shape[2]}, + zerop->data(), + ov::Strides{zerop_strides[0], zerop_strides[2]}); + ov::Tensor wraped_scale(scale->get_element_type(), + ov::Shape{scale_shape[0], scale_shape[1] * scale_shape[2]}, + scale->data(), + ov::Strides{scale_strides[0], scale_strides[2]}); + + ov::npuw::util::XARCH::unpack_i8f16_scale_zp(ov::get_tensor_impl(wraped_from), + ov::get_tensor_impl(wraped_zerop), + ov::get_tensor_impl(wraped_scale), + to, + unpack_options); + } else if (scale_shape.size() == 3 && scale_shape[0] == 1 && scale_shape[2] == 1) { + // Special case for broadcasting vocab by 2 dimensions + // FIXME: all this logic probably should be in some specific unpack or another util function + ov::Tensor wraped_from(from->get_element_type(), ov::Shape{from_shape[1], from_shape[2]}, from->data()); + ov::Tensor wraped_zerop(zerop->get_element_type(), + ov::Shape{zerop_shape[1], zerop_shape[2]}, + zerop->data()); + ov::Tensor wraped_scale(scale->get_element_type(), + ov::Shape{scale_shape[1], scale_shape[2]}, + scale->data()); + + ov::npuw::util::XARCH::unpack_i8f16_scale_zp(ov::get_tensor_impl(wraped_from), + ov::get_tensor_impl(wraped_zerop), + ov::get_tensor_impl(wraped_scale), + to, + unpack_options); + } else if (scale_shape.size() == 2 && scale_shape[0] == from_shape[0] && scale_shape[1] == 1) { + ov::npuw::util::XARCH::unpack_i8f16_scale_zp(from, zerop, scale, to, unpack_options); + } else { + NPUW_ASSERT(false); + } } } diff --git a/src/plugins/intel_npu/src/plugin/npuw/util_xarch.cpp b/src/plugins/intel_npu/src/plugin/npuw/util_xarch.cpp index f1c85ad7db8aea..42458f0b91ce5a 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/util_xarch.cpp +++ b/src/plugins/intel_npu/src/plugin/npuw/util_xarch.cpp @@ -204,6 +204,14 @@ inline __m128i avx2_i8tof16(__m128i vi8, __m256 s) { return _mm256_cvtps_ph(f32scl, _MM_FROUND_TO_NEAREST_INT); // convert: 8 x f32 -> 8 x f16 [128b] } +inline __m128i avx2_i8tof16(__m128i vi8, __m256 z, __m256 s) { + __m256i i32vec = _mm256_cvtepi8_epi32(vi8); // extend: 8 x i8 -> 8 x i32 [256b of 256b] + __m256 f32vec = _mm256_cvtepi32_ps(i32vec); // convert: 8 x i32 -> 8 x f32 [256b of 256b] + __m256 f32sub = _mm256_sub_ps(f32vec, z); // subtract: 8 x f32 -> 8 x f32 [256b of 256b] + __m256 f32scl = _mm256_mul_ps(f32sub, s); // scale: 8 x f32 -> 8 x f32 [256b of 256b] + return _mm256_cvtps_ph(f32scl, _MM_FROUND_TO_NEAREST_INT); // convert: 8 x f32 -> 8 x f16 [128b] +} + inline __m128i avx2_u8tof16_hi(__m128i vu8, __m256 z, __m256 s) { __m256i u32vec = _mm256_cvtepu8_epi32(vu8); // extend: 8 x u8 -> 8 x i32 [256b of 256b] __m256 f32vec = _mm256_cvtepi32_ps(u32vec); // convert: 8 x i32 -> 8 x f32 [256b of 256b] @@ -1485,6 +1493,60 @@ void ov::npuw::util::XARCH::unpack_i8f16_scale(const ov::SoPtr& fro #endif } +void ov::npuw::util::XARCH::unpack_i8f16_scale_zp(const ov::SoPtr& from, + const ov::SoPtr& zerop, + const ov::SoPtr& scale, + const ov::SoPtr& to, + const ov::npuw::util::UnpackOptions& _options) { + NPUW_ASSERT(from->is_continuous()); + NPUW_ASSERT(zerop->is_continuous()); + NPUW_ASSERT(scale->is_continuous()); + NPUW_ASSERT(to->is_continuous()); + NPUW_ASSERT(from->get_size() == to->get_size()); + NPUW_ASSERT(from->get_size() % 8 == 0); + NPUW_ASSERT(scale->get_shape()[0] == from->get_shape()[0]); + NPUW_ASSERT(scale->get_shape()[1] == 1); + NPUW_ASSERT(zerop->get_shape()[0] == from->get_shape()[0]); + NPUW_ASSERT(zerop->get_shape()[1] == 1); + + const auto scale_elem_type = scale->get_element_type(); + NPUW_ASSERT(scale_elem_type == ov::element::f32 || scale_elem_type == ov::element::f16); + + const auto zerop_elem_type = zerop->get_element_type(); + NPUW_ASSERT(zerop_elem_type == ov::element::i8); + +#if defined(HAVE_AVX2) + constexpr std::size_t VECSIZE = 8; + + const std::size_t total = from->get_size(); + const std::size_t stotal = scale->get_size(); + const int8_t* pSrc = from->data(); + const int8_t* pZrp = zerop->data(); + const int8_t* pScl = static_cast(scale->data()); + int16_t* pDst = static_cast(to->data()); + + for (std::size_t sindex = 0u; sindex < stotal; sindex++) { + __m256 svec = avx2_load_scale(pScl, scale_elem_type); + __m128i u8zp = _mm_set1_epi8(*pZrp); // bcast: 8 x i8 + __m256i u32zp = _mm256_cvtepi8_epi32(u8zp); // i32 zero point + __m256 f32zp = _mm256_cvtepi32_ps(u32zp); // f32 zero point + for (std::size_t index = 0u; index < (total / stotal); index += VECSIZE) { + const __m128i* pSrcV = reinterpret_cast(pSrc); + __m128i* pDstV = reinterpret_cast<__m128i*>(pDst); + __m128i i8vec = _mm_loadl_epi64(pSrcV); // load: 8 x i8 [ 64b of 128b] + __m128i f16vec = avx2_i8tof16(i8vec, f32zp, svec); // convert & scale + _mm_store_si128(pDstV, f16vec); // store: 8 x f16 [128b] + pSrc += VECSIZE; + pDst += VECSIZE; + } // index + pScl += scale_elem_type.size(); + pZrp++; + } // sindex +#else + OPENVINO_THROW("AVX2 support is necessary but it's not enabled!"); +#endif +} + void ov::npuw::util::XARCH::unpack_u8f16(const ov::SoPtr& from, const ov::SoPtr& zerop, const ov::SoPtr& scale, diff --git a/src/plugins/intel_npu/src/plugin/npuw/util_xarch.hpp b/src/plugins/intel_npu/src/plugin/npuw/util_xarch.hpp index ab65a1c9bc54aa..1df8b98708c7aa 100644 --- a/src/plugins/intel_npu/src/plugin/npuw/util_xarch.hpp +++ b/src/plugins/intel_npu/src/plugin/npuw/util_xarch.hpp @@ -74,6 +74,12 @@ void unpack_i8f16_scale(const ov::SoPtr& from, const ov::SoPtr& to, const ov::npuw::util::UnpackOptions& unpack_options); +void unpack_i8f16_scale_zp(const ov::SoPtr& from, + const ov::SoPtr& zerop, + const ov::SoPtr& scale, + const ov::SoPtr& to, + const ov::npuw::util::UnpackOptions& _options); + void unpack_u8f16(const ov::SoPtr& from, const ov::SoPtr& zerop, const ov::SoPtr& scale,