diff --git a/src/common/low_precision_transformations/tests/mark_dequantization_subgraph_transformation.cpp b/src/common/low_precision_transformations/tests/mark_dequantization_subgraph_transformation.cpp index 92eb885c249..1bb38b64647 100644 --- a/src/common/low_precision_transformations/tests/mark_dequantization_subgraph_transformation.cpp +++ b/src/common/low_precision_transformations/tests/mark_dequantization_subgraph_transformation.cpp @@ -8,6 +8,7 @@ #include "transformations/fp16_compression/mark_decompression_convert_constant_folding.hpp" #include "transformations/rt_info/decompression.hpp" #include "transformations/rt_info/dequantization_node.hpp" +#include "transformations/rt_info/keep_const_precision.hpp" #include "common_test_utils/ov_test_utils.hpp" @@ -38,7 +39,8 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformation) { // // After MarkDequantizationSubgraph all Subtract and Multiply nodes from above graph // are marked with 'DequantizationNode' attribute. - // Also all 'Convert(DCF)' nodes from above graph are marked with 'DisableConstantFolding' attribute + // All 'Convert(DCF)' nodes from above graph are marked with 'DisableConstantFolding' attribute + // Weights and zero points are marked with 'KeepConstPrecision' attribute { auto parameter = std::make_shared(element::f32, Shape{1, 16, 14, 14}); @@ -107,10 +109,12 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformation) { } std::shared_ptr weights = opset10::Constant::create(element::i8, Shape{4, 16, 1, 1}, {-2}); + enable_keep_const_precision(weights); { auto convert = std::make_shared(weights, element::f32); pass::disable_constant_folding(convert); auto zero_point = opset10::Constant::create(element::i8, Shape{}, {127}); + enable_keep_const_precision(zero_point); auto convert_on_zero_point = std::make_shared(zero_point, element::f32); pass::disable_constant_folding(convert_on_zero_point); auto subtract = std::make_shared(convert, convert_on_zero_point); @@ -157,6 +161,7 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationNoZeroPoint // After MarkDequantizationSubgraph all Multiply nodes from above graph // are marked with 'DequantizationNode' attribute. // Also 'Convert(DCF)' node from above graph is marked with 'DisableConstantFolding' attribute + // Weights node is marked with 'KeepConstPrecision' attribute { auto parameter = std::make_shared(element::f32, Shape{1, 16, 14, 14}); @@ -214,6 +219,7 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationNoZeroPoint } std::shared_ptr weights = opset10::Constant::create(element::i8, Shape{4, 16, 1, 1}, {-2}); + enable_keep_const_precision(weights); { auto convert = std::make_shared(weights, element::f32); pass::disable_constant_folding(convert); @@ -259,6 +265,7 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationNoZeroPoint // After MarkDequantizationSubgraph all Multiply nodes from above graph // are marked with 'DequantizationNode' attribute. // Also 'Convert(DCF)' node from above graph is marked with 'DisableConstantFolding' attribute + // Weights node is marked with 'KeepConstPrecision' attribute { auto parameter = std::make_shared(element::f32, Shape{1, 16, 14, 14}); @@ -323,6 +330,7 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationNoZeroPoint } std::shared_ptr weights = opset10::Constant::create(element::i8, Shape{4, 16, 1, 1}, {-2}); + enable_keep_const_precision(weights); { auto convert = std::make_shared(weights, element::f32); pass::disable_constant_folding(convert); @@ -373,6 +381,7 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationNotConstant // After MarkDequantizationSubgraph all Subtract and Multiply nodes from above graph // are marked with 'DequantizationNode' attribute. // Also all 'Convert(DCF)' nodes from above graph are marked with 'DisableConstantFolding' attribute + // Weights and zero point nodes are marked with 'KeepConstPrecision' attribute { auto parameter = std::make_shared(element::f32, Shape{1, 16, 14, 14}); @@ -395,10 +404,12 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationNotConstant } std::shared_ptr weights = opset10::Constant::create(element::i8, Shape{4, 16, 1, 1}, {-3}); + enable_keep_const_precision(weights); { auto clamp = std::make_shared(weights, -2, 2); auto convert = std::make_shared(clamp, element::f32); auto zero_point = opset10::Constant::create(element::i8, Shape{}, {127}); + enable_keep_const_precision(zero_point); auto convert_on_zero_point = std::make_shared(zero_point, element::f32); auto subtract = std::make_shared(convert, convert_on_zero_point); auto scale = opset10::Constant::create(element::f32, Shape{}, {0.2}); @@ -488,6 +499,7 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationFoldSubCons // are marked with 'DequantizationNode' attribute. // Also all 'Convert(DCF)' node before weights is marked with 'DisableConstantFolding' attribute // but Convert before Dequantization Sub const isn't because fold_subtract_const is set to true + // Weights node is marked with 'KeepConstPrecision' attribute { auto weights = opset10::Constant::create(element::u8, Shape{4, 16, 1, 1}, {3}); @@ -505,6 +517,7 @@ TEST_F(TransformationTestsF, MarkDequantizationSubgraphTransformationFoldSubCons { auto weights = opset10::Constant::create(element::u8, Shape{4, 16, 1, 1}, {3}); + enable_keep_const_precision(weights); auto convert = std::make_shared(weights, element::f32); pass::disable_constant_folding(convert); auto zero_point = opset10::Constant::create(element::f32, Shape{}, {127}); diff --git a/src/common/transformations/src/transformations/low_precision/mark_dequantization_subgraph.cpp b/src/common/transformations/src/transformations/low_precision/mark_dequantization_subgraph.cpp index d79ac3bb873..c2662dc77e2 100644 --- a/src/common/transformations/src/transformations/low_precision/mark_dequantization_subgraph.cpp +++ b/src/common/transformations/src/transformations/low_precision/mark_dequantization_subgraph.cpp @@ -9,6 +9,7 @@ #include "openvino/pass/pattern/op/wrap_type.hpp" #include "transformations/rt_info/dequantization_node.hpp" #include "transformations/rt_info/disable_constant_folding.hpp" +#include "transformations/rt_info/keep_const_precision.hpp" #include "transformations/utils/utils.hpp" ov::pass::MarkDequantizationSubgraph::MarkDequantizationSubgraph(const element::TypeVector& precisions, @@ -62,6 +63,16 @@ ov::pass::MarkDequantizationSubgraph::MarkDequantizationSubgraph(const element:: if (ov::op::util::is_on_constant_path(input)) { // disable ConstantFolding if dequantization subgraph is on constant data ov::disable_constant_folding(convert); + // It is also necessary to avoid precision conversion for constant nodes with input_precision + auto keep_const_precision = [&](Node* node) { + if (auto constant = ov::as_type(node)) { + const auto& const_et = constant->get_element_type(); + if (std::find(precisions.begin(), precisions.end(), const_et) != precisions.end()) + ov::enable_keep_const_precision(convert->get_input_node_shared_ptr(0)); + } + }; + std::unordered_set visited; + ov::op::util::visit_shape_path(input.get_node(), visited, keep_const_precision); } if (subtract_it != pattern_map.end()) { @@ -75,6 +86,7 @@ ov::pass::MarkDequantizationSubgraph::MarkDequantizationSubgraph(const element:: // so we don't have to constantfold it and then convert it back to // low precision in LP transformations ov::disable_constant_folding(zero_point); + ov::enable_keep_const_precision(zero_point->get_input_node_shared_ptr(0)); } } diff --git a/src/plugins/intel_gpu/include/intel_gpu/primitives/fully_connected.hpp b/src/plugins/intel_gpu/include/intel_gpu/primitives/fully_connected.hpp index dc438b691aa..21e72b9e0e5 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/primitives/fully_connected.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/primitives/fully_connected.hpp @@ -3,6 +3,7 @@ // #pragma once +#include "intel_gpu/runtime/optionals.hpp" #include "primitive.hpp" #include @@ -110,6 +111,7 @@ struct fully_connected : public primitive_base { bool compressed_weights = false; primitive_id decompression_scale = ""; primitive_id decompression_zero_point = ""; + optional_value decompression_zero_point_scalar = optional_value(); /// @brief Primitive dimension size. size_t input_size = 2; @@ -124,6 +126,8 @@ struct fully_connected : public primitive_base { seed = hash_combine(seed, compressed_weights); seed = hash_combine(seed, !decompression_scale.empty()); seed = hash_combine(seed, !decompression_zero_point.empty()); + seed = hash_combine(seed, decompression_zero_point_scalar.has_value()); + seed = hash_combine(seed, decompression_zero_point_scalar.value_or(0.0f)); return seed; } @@ -135,7 +139,11 @@ struct fully_connected : public primitive_base { return input_size == rhs_casted.input_size && weights_rank == rhs_casted.weights_rank && - bias.empty() == rhs_casted.bias.empty(); + bias.empty() == rhs_casted.bias.empty() && + compressed_weights == rhs_casted.compressed_weights && + decompression_scale.empty() == rhs_casted.decompression_scale.empty() && + decompression_zero_point.empty() == rhs_casted.decompression_zero_point.empty() && + decompression_zero_point_scalar.value_or(0.0f) == rhs_casted.decompression_zero_point_scalar.value_or(0.0f); } void save(BinaryOutputBuffer& ob) const override { @@ -147,6 +155,13 @@ struct fully_connected : public primitive_base { ob << decompression_zero_point; ob << input_size; ob << weights_rank; + + if (decompression_zero_point_scalar.has_value()) { + ob << true; + ob << make_data(&decompression_zero_point_scalar.value(), sizeof(float)); + } else { + ob << false; + } } void load(BinaryInputBuffer& ib) override { @@ -158,6 +173,16 @@ struct fully_connected : public primitive_base { ib >> decompression_zero_point; ib >> input_size; ib >> weights_rank; + + bool has_value; + ib >> has_value; + if (has_value) { + float decompression_zero_point_value = 0.f; + ib >> make_data(&decompression_zero_point_value, sizeof(float)); + decompression_zero_point_scalar = decompression_zero_point_value; + } else { + decompression_zero_point_scalar = optional_value(); + } } protected: diff --git a/src/plugins/intel_gpu/src/graph/fully_connected.cpp b/src/plugins/intel_gpu/src/graph/fully_connected.cpp index ed396907d7e..4d65e5223e2 100644 --- a/src/plugins/intel_gpu/src/graph/fully_connected.cpp +++ b/src/plugins/intel_gpu/src/graph/fully_connected.cpp @@ -222,6 +222,9 @@ std::string fully_connected_inst::to_string(fully_connected_node const& node) { if (desc->compressed_weights) { fc_info.add("decompression scale id", desc->decompression_scale); fc_info.add("decompression zp id", desc->decompression_zero_point); + if (desc->decompression_zero_point_scalar.has_value()) { + fc_info.add("decompression zp value", desc->decompression_zero_point_scalar.value()); + } } node_info->add("fully connected info", fc_info); diff --git a/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp b/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp index cc9d8602316..167b77c74af 100644 --- a/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp +++ b/src/plugins/intel_gpu/src/graph/graph_optimizer/prepare_primitive_fusing.cpp @@ -1200,8 +1200,25 @@ void prepare_primitive_fusing::fuse_constant_transposes(program& p) { next_node->is_type() || next_node->is_type()) { size_t weights_offset = next_node->get_primitive()->input_size(); - return &next_node->get_dependency(weights_offset) == node ? next_node - : nullptr; + std::vector valid_weights_indices = {next_node->get_primitive()->input_size()}; + if (next_node->is_type()) { + auto& fc = next_node->as(); + auto desc = fc.get_primitive(); + if (desc->compressed_weights) { + size_t scale_idx = weights_offset + (fc.bias_term() ? 2 : 1); + valid_weights_indices.push_back(scale_idx); + if (!desc->decompression_zero_point.empty()) { + valid_weights_indices.push_back(scale_idx + 1); + } + } + } + + for (auto& widx : valid_weights_indices) { + if (&next_node->get_dependency(widx) == node) { + return next_node; + } + } + return nullptr; } if (node->is_constant() && node->get_users().size() == 1) diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp index 19007a48157..e3ac31cccc7 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp @@ -161,6 +161,10 @@ public: if (with_zp) { params.has_decompression_zp = true; params.decompression_zero_point = convert_data_tensor(input_layouts[3]); + } else if (primitive->decompression_zero_point_scalar.has_value()) { + params.has_decompression_zp = true; + params.scalar_zp = true; + params.zp_value = primitive->decompression_zero_point_scalar.value(); } } diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp index f0872d37029..e9b98a57e00 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/kernel_selector_helper.cpp @@ -123,6 +123,10 @@ kernel_selector::data_type to_data_type(data_types dt) { switch (dt) { case cldnn::data_types::u1: return kernel_selector::data_type::BINARY; + case cldnn::data_types::i4: + return kernel_selector::data_type::INT4; + case cldnn::data_types::u4: + return kernel_selector::data_type::UINT4; case cldnn::data_types::i8: return kernel_selector::data_type::INT8; case cldnn::data_types::u8: @@ -144,6 +148,10 @@ data_types from_data_type(kernel_selector::data_type dt) { switch (dt) { case kernel_selector::data_type::BINARY: return cldnn::data_types::u1; + case kernel_selector::data_type::INT4: + return cldnn::data_types::i4; + case kernel_selector::data_type::UINT4: + return cldnn::data_types::u4; case kernel_selector::data_type::INT8: return cldnn::data_types::i8; case kernel_selector::data_type::UINT8: @@ -165,6 +173,10 @@ kernel_selector::weights_type to_weights_type(data_types dt) { switch (dt) { case cldnn::data_types::u1: return kernel_selector::weights_type::BINARY; + case cldnn::data_types::u4: + return kernel_selector::weights_type::UINT4; + case cldnn::data_types::i4: + return kernel_selector::weights_type::INT4; case cldnn::data_types::i8: return kernel_selector::weights_type::INT8; case cldnn::data_types::u8: @@ -184,6 +196,10 @@ data_types from_weights_type(kernel_selector::weights_type dt) { switch (dt) { case kernel_selector::weights_type::BINARY: return data_types::u1; + case kernel_selector::weights_type::INT4: + return data_types::i4; + case kernel_selector::weights_type::UINT4: + return data_types::u4; case kernel_selector::weights_type::INT8: return data_types::i8; case kernel_selector::weights_type::UINT8: diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl index f6dacec4a73..dd8adb53e76 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl @@ -38,10 +38,17 @@ # error "fully_connected_gpu_bf_tiled.cl - TILE_K_OFM must be equal to TILE_K * TILE_OFM and at most 8" #endif +#if COMPRESSED_WEIGHTS_INT4 +# if TILE_K_OFM != TILE_K_OFM_PACKED * 2 +# error "fully_connected_gpu_bf_tiled.cl - TILE_K_OFM must be divisible by 2 for 4-bit compressed case" +# endif +#endif + // Macros for vectorized types. #define INPUT_VEC_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, TILE_IFM) #define ACCUMULATOR_VEC_TYPE MAKE_VECTOR_TYPE(ACCUMULATOR_TYPE, TILE_OFM) #define FILTER_VEC_TYPE MAKE_VECTOR_TYPE(ACCUMULATOR_TYPE, TILE_K_OFM) +#define FILTER_PACKED_VEC_TYPE MAKE_VECTOR_TYPE(FILTER_TYPE, TILE_K_OFM_PACKED) #define BIAS_VEC_TYPE MAKE_VECTOR_TYPE(BIAS_TYPE, TILE_OFM) #define OUTPUT_VEC_TYPE MAKE_VECTOR_TYPE(OUTPUT_TYPE, TILE_OFM) #define ACTIVATION_VEC_TYPE MAKE_VECTOR_TYPE(ACTIVATION_TYPE, TILE_OFM) @@ -50,7 +57,7 @@ #define TO_FILTER_VEC_TYPE(x) CAT(convert_, FILTER_VEC_TYPE)(x) #define INPUT_BLOCK_READ(ptr, offset) BLOCK_READN(INPUT0_TYPE, TILE_IFM, ptr, offset) -#define FILTER_BLOCK_READ(ptr, offset) BLOCK_READN(FILTER_TYPE, TILE_K_OFM, ptr, offset) +#define FILTER_BLOCK_READ(ptr, offset) BLOCK_READN(FILTER_TYPE, TILE_K_OFM_PACKED, ptr, offset) #define BIAS_BLOCK_READ(ptr, offset) BLOCK_READN(BIAS_TYPE, TILE_OFM, ptr, offset) #define OUTPUT_BLOCK_WRITE(ptr, offset, val) BLOCK_WRITEN(OUTPUT_TYPE, TILE_OFM, ptr, offset, val) @@ -85,7 +92,7 @@ KERNEL(fc)( #if DECOMPRESSION_SCALE_TERM const __global DECOMPRESSION_SCALE_TYPE* decompression_scale, #endif -#if DECOMPRESSION_ZP_TERM +#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR const __global DECOMPRESSION_ZP_TYPE* decompression_zp, #endif __global OUTPUT_TYPE* output, @@ -118,12 +125,16 @@ KERNEL(fc)( FILTER_VEC_TYPE wei = 0; uint input_offset = out_b * TILE_IN_B_PITCH + INPUT0_OFFSET; +#if COMPRESSED_WEIGHTS_INT4 + uint weights_offset = out_f * (INPUT_ELEMENTS_COUNT / 2); +#else uint weights_offset = out_f * INPUT_ELEMENTS_COUNT; +#endif #if COMPRESSED_WEIGHTS && DECOMPRESSION_SCALE_GROUPS_NUM == 1 - #if DECOMPRESSION_SCALE_LENGTH > 1 && DECOMPRESSION_SCALE_LENGTH % SIMD == 0 + #if DECOMPRESSION_SCALE_LENGTH > 1 && DECOMPRESSION_SCALE_LENGTH % (TILE_OFM * SIMD) == 0 ACCUMULATOR_VEC_TYPE d_scale = BLOCK_READN(ACCUMULATOR_TYPE, TILE_OFM, decompression_scale, out_f); - #elif DECOMPRESSION_SCALE_LENGTH > 1 && DECOMPRESSION_SCALE_LENGTH % SIMD != 0 + #elif DECOMPRESSION_SCALE_LENGTH > 1 && DECOMPRESSION_SCALE_LENGTH % (TILE_OFM * SIMD) != 0 ACCUMULATOR_VEC_TYPE d_scale = 0; unroll_for(uint of = 0; of < TILE_OFM; ++of) { uint offset = out_f + of*SIMD + get_sub_group_local_id(); @@ -137,10 +148,10 @@ KERNEL(fc)( ACCUMULATOR_TYPE* d_scales = (ACCUMULATOR_TYPE*)(&d_scale); #endif -#if COMPRESSED_WEIGHTS && DECOMPRESSION_ZP_TERM && DECOMPRESSION_ZP_GROUPS_NUM == 1 - #if DECOMPRESSION_ZP_LENGTH > 1 && DECOMPRESSION_ZP_LENGTH % SIMD == 0 +#if COMPRESSED_WEIGHTS && DECOMPRESSION_ZP_TERM && DECOMPRESSION_ZP_GROUPS_NUM == 1 && !DECOMPRESSION_ZP_SCALAR + #if DECOMPRESSION_ZP_LENGTH > 1 && DECOMPRESSION_ZP_LENGTH % (TILE_OFM * SIMD) == 0 ACCUMULATOR_VEC_TYPE d_zp = BLOCK_READN(ACCUMULATOR_TYPE, TILE_OFM, decompression_zp, out_f); - #elif DECOMPRESSION_ZP_LENGTH > 1 && DECOMPRESSION_ZP_LENGTH % SIMD != 0 + #elif DECOMPRESSION_ZP_LENGTH > 1 && DECOMPRESSION_ZP_LENGTH % (TILE_OFM * SIMD) != 0 ACCUMULATOR_VEC_TYPE d_zp = 0; unroll_for(uint of = 0; of < TILE_OFM; ++of) { uint offset = out_f + of*SIMD + get_sub_group_local_id(); @@ -187,29 +198,46 @@ KERNEL(fc)( // NOTE: Manually unrolling multiplication loop leads to lower register pressure and allows for bigger block sizes, // but significantly degrades readability and generality of code. // It doesn't also show noticable performance improvement on tested configurations. + #if DECOMPRESSION_SCALE_POST_OP + ACCUMULATOR_VEC_TYPE acc_tmp[TILE_B] = { }; + #endif + unroll_for(uint ki = 0; ki < (TILE_IFM * SIMD) / TILE_K; ++ki) { - wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); + #if COMPRESSED_WEIGHTS_INT4 + FILTER_PACKED_VEC_TYPE wei_packed = FILTER_BLOCK_READ(weights, weights_offset); + wei = UNPACK_INT4x2(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&wei_packed)); + #else + wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); + #endif + #if COMPRESSED_WEIGHTS ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); unroll_for(uint kii = 0; kii < TILE_K; ++kii) { unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { const uint w_idx = kii * TILE_OFM + fi; const uint offset_ofm = out_f + fi*SIMD + sglid; - #if DECOMPRESSION_SCALE_GROUPS_NUM > 1 - const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH + - ((kii + ki*TILE_K + ni*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; - ACCUMULATOR_TYPE ds = decompression_scale[scale_offset]; + #if !DECOMPRESSION_SCALE_POST_OP + // Apply scales before FMA to avoid FP16 overflow in case of INT8 + #if DECOMPRESSION_SCALE_GROUPS_NUM > 1 + const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH + + ((kii + ki*TILE_K + ni*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; + ACCUMULATOR_TYPE ds = decompression_scale[scale_offset]; + #else + ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; + #endif #else - ACCUMULATOR_TYPE ds = d_scales[fi]; + ACCUMULATOR_TYPE ds = ACCUMULATOR_VAL_ONE; #endif #if DECOMPRESSION_ZP_TERM - #if DECOMPRESSION_ZP_GROUPS_NUM > 1 + #if DECOMPRESSION_ZP_SCALAR + ACCUMULATOR_TYPE dzp = DECOMPRESSION_ZP_VALUE; + #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + ((kii + ki*TILE_K + ni*TILE_IFM*SIMD) / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; ACCUMULATOR_TYPE dzp = decompression_zp[zp_offset]; #else - ACCUMULATOR_TYPE dzp = d_zps[fi]; + ACCUMULATOR_TYPE dzp = d_zps[fi % DECOMPRESSION_ZP_LENGTH]; #endif #else ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; @@ -218,18 +246,38 @@ KERNEL(fc)( } } #endif - weights_offset += TILE_K_OFM * SIMD; + weights_offset += TILE_K_OFM_PACKED * SIMD; unroll_for (uint kii = 0; kii < TILE_K; ++kii) { const uint total_k = ki * TILE_K + kii; unroll_for (uint bi = 0; bi < TILE_B; ++bi) { INPUT0_TYPE in_val = _sub_group_shuffle(((INPUT0_TYPE*)(&in_0[bi]))[total_k / SIMD], total_k % SIMD); unroll_for (uint fi = 0; fi < TILE_OFM; ++fi) { +#if DECOMPRESSION_SCALE_POST_OP + ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[kii * TILE_OFM + fi]; +#else ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[kii * TILE_OFM + fi]; +#endif } } } } +#if DECOMPRESSION_SCALE_POST_OP + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + const uint offset_ofm = out_f + fi*SIMD + sglid; + + #if DECOMPRESSION_SCALE_GROUPS_NUM > 1 + const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH + + ((ni*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; + ACCUMULATOR_TYPE ds = decompression_scale[scale_offset]; + #else + ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; + #endif + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; + } + } +#endif } // ===================================================================================================================================== // Leftovers @@ -246,7 +294,13 @@ KERNEL(fc)( #undef LOAD_IN_0 input_offset += TILE_IFM * SIMD - TILE_IN_B_PITCH * TILE_B; unroll_for(uint ki = 0; ki < CEIL_DIV(LEFTOVER_IFM, TILE_K); ++ki) { - wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); + #if COMPRESSED_WEIGHTS_INT4 + FILTER_PACKED_VEC_TYPE wei_packed = FILTER_BLOCK_READ(weights, weights_offset); + wei = UNPACK_INT4x2(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&wei_packed)); + #else + wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); + #endif + #if COMPRESSED_WEIGHTS ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); unroll_for(uint kii = 0; kii < TILE_K; ++kii) { @@ -258,16 +312,18 @@ KERNEL(fc)( ((kii + ki*TILE_K + ni*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; ACCUMULATOR_TYPE ds = decompression_scale[scale_offset]; #else - ACCUMULATOR_TYPE ds = d_scales[fi]; + ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; #endif #if DECOMPRESSION_ZP_TERM - #if DECOMPRESSION_ZP_GROUPS_NUM > 1 + #if DECOMPRESSION_ZP_SCALAR + ACCUMULATOR_TYPE dzp = DECOMPRESSION_ZP_VALUE; + #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + ((kii + ki*TILE_K + ni*TILE_IFM*SIMD) / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; ACCUMULATOR_TYPE dzp = decompression_zp[zp_offset]; #else - ACCUMULATOR_TYPE dzp = d_zps[fi]; + ACCUMULATOR_TYPE dzp = d_zps[fi % DECOMPRESSION_ZP_LENGTH]; #endif #else ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; @@ -276,7 +332,7 @@ KERNEL(fc)( } } #endif - weights_offset += TILE_K_OFM * SIMD; + weights_offset += TILE_K_OFM_PACKED * SIMD; unroll_for (uint kii = 0; kii < TILE_K; ++kii) { unroll_for (uint fi = 0; fi < TILE_OFM; ++fi) { diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bfyx_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bfyx_ref.cl index 6374e65c4f5..b83018e8c9b 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bfyx_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bfyx_ref.cl @@ -4,6 +4,7 @@ #include "include/batch_headers/fetch_data.cl" #include "include/batch_headers/fetch_weights.cl" +#include "include/batch_headers/int4_utils.cl" KERNEL(fc)( OPTIONAL_SHAPE_INFO_ARG @@ -11,7 +12,7 @@ KERNEL(fc)( #if DECOMPRESSION_SCALE_TERM const __global DECOMPRESSION_SCALE_TYPE* decompression_scale, #endif -#if DECOMPRESSION_ZP_TERM +#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR const __global DECOMPRESSION_ZP_TYPE* decompression_zp, #endif __global OUTPUT_TYPE* output, @@ -38,8 +39,12 @@ KERNEL(fc)( const uint input0_idx = INPUT0_GET_INDEX(b, ofm, y, x); #if COMPRESSED_WEIGHTS #if DECOMPRESSION_ZP_TERM - const uint zp_offset = DECOMPRESSION_ZP_GET_INDEX_SAFE(oym, y / DECOMPRESSION_ZP_GROUP_SIZE, 0, 0); - ACCUMULATOR_TYPE zp = TO_ACCUMULATOR_TYPE(decompression_zp[zp_offset]); + #if DECOMPRESSION_ZP_SCALAR + ACCUMULATOR_TYPE zp = DECOMPRESSION_ZP_VALUE; + #else + const uint zp_offset = DECOMPRESSION_ZP_GET_INDEX_SAFE(oym, y / DECOMPRESSION_ZP_GROUP_SIZE, 0, 0); + ACCUMULATOR_TYPE zp = TO_ACCUMULATOR_TYPE(decompression_zp[zp_offset]); + #endif #else ACCUMULATOR_TYPE zp = ACCUMULATOR_VAL_ZERO; #endif @@ -47,13 +52,19 @@ KERNEL(fc)( DECOMPRESSION_SCALE_TYPE scale = decompression_scale[decomp_offset]; #endif - #if COMPRESSED_WEIGHTS_INT8 const uint filter_idx = GET_FILTER_INDEX(FILTER, 0, oym, y, 0, 0); + #if COMPRESSED_WEIGHTS_INT8 ACCUMULATOR_TYPE filter_compressed = TO_ACCUMULATOR_TYPE(weights[filter_idx]); ACCUMULATOR_TYPE filter_val = (filter_compressed - zp) * scale; dotProd += (ACCUMULATOR_TYPE)(input[input0_idx]) * (ACCUMULATOR_TYPE)(filter_val); + #elif COMPRESSED_WEIGHTS_INT4 + FILTER_TYPE filter_packed = weights[filter_idx / 2]; + MAKE_VECTOR_TYPE(ACCUMULATOR_TYPE, 2) filter_unpacked = UNPACK_INT4x2(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&filter_packed)); + + ACCUMULATOR_TYPE filter_compressed = ((ACCUMULATOR_TYPE*)(&filter_unpacked))[filter_idx % 2]; + ACCUMULATOR_TYPE filter_val = (filter_compressed - zp) * scale; + dotProd += (ACCUMULATOR_TYPE)(input[input0_idx]) * filter_val; #else - const uint filter_idx = GET_FILTER_INDEX(FILTER, 0, oym, y, 0, 0); dotProd += (ACCUMULATOR_TYPE)(input[input0_idx]) * (ACCUMULATOR_TYPE)(weights[filter_idx]); #endif } @@ -75,8 +86,12 @@ KERNEL(fc)( const uint input0_idx = INPUT0_GET_INDEX(b, ifm, y, x); #if COMPRESSED_WEIGHTS #if DECOMPRESSION_ZP_TERM - const uint zp_offset = DECOMPRESSION_ZP_GET_INDEX_SAFE(ofm, ifm / DECOMPRESSION_ZP_GROUP_SIZE, 0, 0); - ACCUMULATOR_TYPE zp = TO_ACCUMULATOR_TYPE(decompression_zp[zp_offset]); + #if DECOMPRESSION_ZP_SCALAR + ACCUMULATOR_TYPE zp = DECOMPRESSION_ZP_VALUE; + #else + const uint zp_offset = DECOMPRESSION_ZP_GET_INDEX_SAFE(ofm, ifm / DECOMPRESSION_ZP_GROUP_SIZE, 0, 0); + ACCUMULATOR_TYPE zp = TO_ACCUMULATOR_TYPE(decompression_zp[zp_offset]); + #endif #else ACCUMULATOR_TYPE zp = ACCUMULATOR_VAL_ZERO; #endif @@ -84,14 +99,19 @@ KERNEL(fc)( DECOMPRESSION_SCALE_TYPE scale = decompression_scale[decomp_offset]; #endif - - #if COMPRESSED_WEIGHTS_INT8 const uint filter_idx = GET_FILTER_INDEX(FILTER, 0, ofm, ifm, y, x); + #if COMPRESSED_WEIGHTS_INT8 FILTER_TYPE filter_compressed = weights[filter_idx]; ACCUMULATOR_TYPE filter_val = (TO_ACCUMULATOR_TYPE(filter_compressed) - zp) * scale; dotProd += (ACCUMULATOR_TYPE)(input[input0_idx]) * (ACCUMULATOR_TYPE)(filter_val); + #elif COMPRESSED_WEIGHTS_INT4 + FILTER_TYPE filter_packed = weights[filter_idx / 2]; + MAKE_VECTOR_TYPE(ACCUMULATOR_TYPE, 2) filter_unpacked = UNPACK_INT4x2(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&filter_packed)); + + ACCUMULATOR_TYPE filter_compressed = ((ACCUMULATOR_TYPE*)(&filter_unpacked))[filter_idx % 2]; + ACCUMULATOR_TYPE filter_val = (filter_compressed - zp) * scale; + dotProd += (ACCUMULATOR_TYPE)(input[input0_idx]) * filter_val; #else - const uint filter_idx = GET_FILTER_INDEX(FILTER, 0, ofm, ifm, y, x); dotProd += (ACCUMULATOR_TYPE)(input[input0_idx]) * (ACCUMULATOR_TYPE)(weights[filter_idx]); #endif } diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/int4_utils.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/int4_utils.cl new file mode 100644 index 00000000000..db786b30c15 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/int4_utils.cl @@ -0,0 +1,101 @@ +// Copyright (C) 2023 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +typedef struct __attribute__ ((packed)) int4x2_t { char s0; } int4x2_t; +typedef struct __attribute__ ((packed)) int4x4_t { int4x2_t s0; int4x2_t s1; } int4x4_t; +typedef struct __attribute__ ((packed)) int4x8_t { int4x2_t s0; int4x2_t s1; int4x2_t s2; int4x2_t s3; } int4x8_t; +typedef struct __attribute__ ((packed)) int4x16_t { int4x2_t s0; int4x2_t s1; int4x2_t s2; int4x2_t s3; int4x2_t s4; int4x2_t s5; int4x2_t s6; int4x2_t s7; } int4x16_t; + +typedef struct __attribute__ ((packed)) uint4x2_t { uchar s0; } uint4x2_t; +typedef struct __attribute__ ((packed)) uint4x4_t { uint4x2_t s0; uint4x2_t s1; } uint4x4_t; +typedef struct __attribute__ ((packed)) uint4x8_t { uint4x2_t s0; uint4x2_t s1; uint4x2_t s2; uint4x2_t s3; } uint4x8_t; +typedef struct __attribute__ ((packed)) uint4x16_t { uint4x2_t s0; uint4x2_t s1; uint4x2_t s2; uint4x2_t s3; uint4x2_t s4; uint4x2_t s5; uint4x2_t s6; uint4x2_t s7; } uint4x16_t; + +inline uchar2 cvt_uint4x2_to_uint8x2(uint4x2_t v) __attribute__((overloadable)) { + const uchar v0 = v.s0 & 0x0F; + const uchar v1 = (v.s0 & 0xF0) >> 4; + return (uchar2)(v0, v1); +} + +inline char2 cvt_int4x2_to_int8x2(int4x2_t v) __attribute__((overloadable)) { + const char s_bit = (v.s0 & convert_char(0x08)); + const char mask = s_bit > 0 ? convert_char(0xF0) : convert_char(0x00); + const char v0 = (v.s0 & convert_char(0x0F)) | mask; + const char v1 = v.s0 >> 4; + return (char2)(v0, v1); +} + +inline half2 unpack_to_half(uint4x2_t v) __attribute__((overloadable)) { + return convert_half2(cvt_uint4x2_to_uint8x2(v)); +} + +inline float2 unpack_to_float(uint4x2_t v) __attribute__((overloadable)) { + return convert_float2(cvt_uint4x2_to_uint8x2(v)); +} + +inline half2 unpack_to_half(int4x2_t v) __attribute__((overloadable)) { + return convert_half2(cvt_int4x2_to_int8x2(v)); +} + +inline float2 unpack_to_float(int4x2_t v) __attribute__((overloadable)) { + return convert_float2(cvt_int4x2_to_int8x2(v)); +} + +inline half4 unpack_to_half(uint4x4_t v) __attribute__((overloadable)) { + half2 f0 = unpack_to_half(v.s0); + half2 f1 = unpack_to_half(v.s1); + return (half4)(f0.s0, f0.s1, f1.s0, f1.s1); +} + +inline float4 unpack_to_float(uint4x4_t v) __attribute__((overloadable)) { + float2 f0 = unpack_to_float(v.s0); + float2 f1 = unpack_to_float(v.s1); + return (float4)(f0.s0, f0.s1, f1.s0, f1.s1); +} + +inline half4 unpack_to_half(int4x4_t v) __attribute__((overloadable)) { + half2 f0 = unpack_to_half(v.s0); + half2 f1 = unpack_to_half(v.s1); + return (half4)(f0.s0, f0.s1, f1.s0, f1.s1); +} + +inline float4 unpack_to_float(int4x4_t v) __attribute__((overloadable)) { + float2 f0 = unpack_to_float(v.s0); + float2 f1 = unpack_to_float(v.s1); + return (float4)(f0.s0, f0.s1, f1.s0, f1.s1); +} + +inline half8 unpack_to_half(uint4x8_t v) __attribute__((overloadable)) { + half2 f0 = unpack_to_half(v.s0); + half2 f1 = unpack_to_half(v.s1); + half2 f2 = unpack_to_half(v.s2); + half2 f3 = unpack_to_half(v.s3); + return (half8)(f0.s0, f0.s1, f1.s0, f1.s1, f2.s0, f2.s1, f3.s0, f3.s1); +} + +inline float8 unpack_to_float(uint4x8_t v) __attribute__((overloadable)) { + float2 f0 = unpack_to_float(v.s0); + float2 f1 = unpack_to_float(v.s1); + float2 f2 = unpack_to_float(v.s2); + float2 f3 = unpack_to_float(v.s3); + return (float8)(f0.s0, f0.s1, f1.s0, f1.s1, f2.s0, f2.s1, f3.s0, f3.s1); +} + +inline half8 unpack_to_half(int4x8_t v) __attribute__((overloadable)) { + half2 f0 = unpack_to_half(v.s0); + half2 f1 = unpack_to_half(v.s1); + half2 f2 = unpack_to_half(v.s2); + half2 f3 = unpack_to_half(v.s3); + return (half8)(f0.s0, f0.s1, f1.s0, f1.s1, f2.s0, f2.s1, f3.s0, f3.s1); +} + +inline float8 unpack_to_float(int4x8_t v) __attribute__((overloadable)) { + float2 f0 = unpack_to_float(v.s0); + float2 f1 = unpack_to_float(v.s1); + float2 f2 = unpack_to_float(v.s2); + float2 f3 = unpack_to_float(v.s3); + return (float8)(f0.s0, f0.s1, f1.s0, f1.s1, f2.s0, f2.s1, f3.s0, f3.s1); +} + +#define UNPACK_INT4x2(target_type, value) CAT(unpack_to_, target_type)(value) diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_weights_int4.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_weights_int4.cl new file mode 100644 index 00000000000..46eb8909277 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_weights_int4.cl @@ -0,0 +1,54 @@ +// Copyright (C) 2018-2023 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "include/batch_headers/fetch_weights.cl" + +KERNEL(reorder_weights_int4)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output) { +#if defined(INPUT0_LAYOUT_IOYX) && defined(OUTPUT_LAYOUT_OIYX) + const uint out_byte_offset = get_global_id(0); + + const uint offset0 = out_byte_offset * 2 + 0; + const uint offset1 = out_byte_offset * 2 + 1; + + const uint i0 = offset0 % OUTPUT_IFM_NUM; + const uint i1 = offset1 % OUTPUT_IFM_NUM; + + const uint o0 = offset0 / OUTPUT_IFM_NUM; + const uint o1 = offset1 / OUTPUT_IFM_NUM; + + const uint input0_offset = GET_FILTER_INDEX(INPUT0, 0, o0, i0, 0, 0); + const uint input1_offset = GET_FILTER_INDEX(INPUT0, 0, o1, i1, 0, 0); + + const uint input0_idx = input0_offset % 2; + const uint input1_idx = input1_offset % 2; + + INPUT0_TYPE in0 = (input[input0_offset / 2] >> input0_idx*4) & 0x0F; + INPUT0_TYPE in1 = (input[input1_offset / 2] >> input1_idx*4) & 0x0F; + + OUTPUT_TYPE out = in0 | (in1 << 4); + output[out_byte_offset] = out; +#elif defined(OUTPUT_LAYOUT_OS_IYX_OSV32) + const unsigned o = (uint)get_global_id(0); + const unsigned i = (uint)get_global_id(1); + + const unsigned o0 = (o / 16) * 32 + (o % 16); + const unsigned o1 = (o / 16) * 32 + (o % 16) + 16; + + const uint input0_offset = GET_FILTER_INDEX(INPUT0, 0, o0, i, 0, 0); + const uint input1_offset = GET_FILTER_INDEX(INPUT0, 0, o1, i, 0, 0); + + const uint input0_idx = input0_offset % 2; + const uint input1_idx = input1_offset % 2; + + INPUT0_TYPE in0 = (input[input0_offset / 2] >> input0_idx*4) & 0x0F; + INPUT0_TYPE in1 = (input[input1_offset / 2] >> input1_idx*4) & 0x0F; + + INPUT0_TYPE packed_out_channels = in0 | (in1 << 4); + + const uint output_idx = GET_FILTER_OS_IYX_OSV_INDEX(OUTPUT, o, i, 0, 0, 32 / 2); // Calculate offset as osv16 due to packing + output[output_idx] = packed_out_channels; +#else +#error "reorder_weights_int4: unsupported layouts combination" +#endif +} diff --git a/src/plugins/intel_gpu/src/kernel_selector/common_types.h b/src/plugins/intel_gpu/src/kernel_selector/common_types.h index 1acc0aa89e6..6d84aa65f37 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/common_types.h +++ b/src/plugins/intel_gpu/src/kernel_selector/common_types.h @@ -106,6 +106,8 @@ enum class KernelType { enum class Datatype { UNSUPPORTED, BINARY, + UINT4, + INT4, INT8, UINT8, INT16, @@ -127,6 +129,8 @@ enum class WeightsType { F32, INT8, UINT8, + UINT4, + INT4, INT32 }; diff --git a/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp b/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp index 67c25b89025..ac079ea4487 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp @@ -115,8 +115,10 @@ std::string toCLType(WeightsType wType) { switch (wType) { case WeightsType::BINARY: return GetTypeName(); + case WeightsType::INT4: case WeightsType::INT8: return GetTypeName(); + case WeightsType::UINT4: case WeightsType::UINT8: return GetTypeName(); case WeightsType::F16: @@ -1343,18 +1345,18 @@ JitConstants MakeActivationJitConstants(ActivationFunction activation_function, } JitConstants MakeTypeJitConstants(Datatype dataType, const std::string& macroName) { - std::string type; - std::string max_val; - std::string min_val; - std::string val_one; - std::string val_zero; - std::string to_type; - std::string to_type_sat; - std::string as_type; - std::string max_func; - std::string min_func; - std::string abs_func; - std::string type_size; + std::string type = "undefined"; + std::string max_val = "undefined"; + std::string min_val = "undefined"; + std::string val_one = "undefined"; + std::string val_zero = "undefined"; + std::string to_type = "undefined"; + std::string to_type_sat = "undefined"; + std::string as_type = "undefined"; + std::string max_func = "undefined"; + std::string min_func = "undefined"; + std::string abs_func = "undefined"; + std::string type_size = "undefined"; bool is_fp; switch (dataType) { case Datatype::INT8: @@ -1478,6 +1480,16 @@ JitConstants MakeTypeJitConstants(Datatype dataType, const std::string& macroNam type_size = "2"; is_fp = true; break; + case Datatype::INT4: + type = "char"; + type_size = "0.5f"; + is_fp = false; + break; + case Datatype::UINT4: + type = "uchar"; + type_size = "0.5f"; + is_fp = false; + break; default: type = "float"; max_val = "FLT_MAX"; @@ -1523,6 +1535,10 @@ JitConstants MakeTypeJitConstants(WeightsType weightsType, const std::string& ma return MakeTypeJitConstants(Datatype::INT8, macroName); case WeightsType::UINT8: return MakeTypeJitConstants(Datatype::UINT8, macroName); + case WeightsType::INT4: + return MakeTypeJitConstants(Datatype::INT4, macroName); + case WeightsType::UINT4: + return MakeTypeJitConstants(Datatype::UINT4, macroName); case WeightsType::BINARY: return MakeTypeJitConstants(Datatype::UINT32, macroName); case WeightsType::INT32: @@ -1533,6 +1549,17 @@ JitConstants MakeTypeJitConstants(WeightsType weightsType, const std::string& ma return MakeTypeJitConstants(Datatype::UNSUPPORTED, macroName); } +JitConstants make_int4_packed_type_jit_constant(const std::string& macro_name, WeightsType wt, size_t pack_size) { + OPENVINO_ASSERT(pack_size % 2 == 0 && pack_size != 0 && pack_size <= 16); + std::string type_string = ""; + switch (wt) { + case WeightsType::UINT4: type_string = "uint4x"; break; + case WeightsType::INT4: type_string = "int4x"; break; + default: OPENVINO_THROW("[GPU] Unsupported compressed type"); + } + return { MakeJitConstant(macro_name, type_string + std::to_string(pack_size) + "_t") }; +} + JitConstants MakeActivationJitConstants(const base_activation_params& params, Datatype out_dt, const std::string& suffix, diff --git a/src/plugins/intel_gpu/src/kernel_selector/jitter.h b/src/plugins/intel_gpu/src/kernel_selector/jitter.h index 29dd56efe37..aa1ec7f7d9d 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/jitter.h +++ b/src/plugins/intel_gpu/src/kernel_selector/jitter.h @@ -322,7 +322,7 @@ JitConstants MakeConstantLoopUnrollJitConstants(uint32_t loopCount); JitConstants MakeTypeJitConstants(Datatype dataType, const std::string& macroName); JitConstants MakeTypeJitConstants(WeightsType weightsType, const std::string& macroName); inline JitConstants MakeUnitTypeJitConstants(Datatype dataType) { return MakeTypeJitConstants(dataType, "UNIT"); } - +JitConstants make_int4_packed_type_jit_constant(const std::string& macro_name, WeightsType wt, size_t pack_size); class FusedOpsCodeGenerator { public: diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_common.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_common.cpp index 8c6d2af2fd8..193d589daba 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_common.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_common.cpp @@ -142,6 +142,8 @@ std::string toString(DataLayout l) { std::string toString(Datatype dType) { switch (dType) { case Datatype::BINARY: return "BINARY"; + case Datatype::UINT4: return "UINT4"; + case Datatype::INT4: return "INT4"; case Datatype::INT8: return "INT8"; case Datatype::UINT8: return "UINT8"; case Datatype::INT16: return "INT16"; @@ -160,6 +162,8 @@ std::string toString(WeightsType wType) { case WeightsType::BINARY: return "BINARY"; case WeightsType::F16: return "F16"; case WeightsType::F32: return "F32"; + case WeightsType::UINT4: return "UINT4"; + case WeightsType::INT4: return "INT4"; case WeightsType::INT8: return "INT8"; case WeightsType::UINT8: return "UINT8"; case WeightsType::INT32: return "INT32"; diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.cpp index 32d451704fa..d7b3002bb6f 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.cpp @@ -47,6 +47,12 @@ DeviceFeaturesKey EngineInfo::get_supported_device_features_key() const { //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// void ParamsKey::EnableInputDataType(Datatype dt) { switch (dt) { + case Datatype::INT4: + key.inputType.val.int4 = 1; + break; + case Datatype::UINT4: + key.inputType.val.uint4 = 1; + break; case Datatype::INT8: key.inputType.val.int8 = 1; break; @@ -86,6 +92,12 @@ void ParamsKey::EnableAllInputDataType() { key.inputType.raw = 0xffffffff; } void ParamsKey::EnableOutputDataType(Datatype dt) { switch (dt) { + case Datatype::INT4: + key.outputType.val.int4 = 1; + break; + case Datatype::UINT4: + key.outputType.val.uint4 = 1; + break; case Datatype::INT8: key.outputType.val.int8 = 1; break; @@ -134,6 +146,12 @@ void ParamsKey::EnableInputWeightsType(WeightsType wt) { case WeightsType::INT8: key.inputWeightsType.val.int8 = 1; break; + case WeightsType::INT4: + key.inputWeightsType.val.int4 = 1; + break; + case WeightsType::UINT4: + key.inputWeightsType.val.uint4 = 1; + break; case WeightsType::BINARY: key.inputWeightsType.val.binary = 1; break; @@ -157,6 +175,12 @@ void ParamsKey::EnableOutputWeightsType(WeightsType wt) { case WeightsType::INT8: key.outputWeightsType.val.int8 = 1; break; + case WeightsType::INT4: + key.outputWeightsType.val.int4 = 1; + break; + case WeightsType::UINT4: + key.outputWeightsType.val.uint4 = 1; + break; case WeightsType::BINARY: key.outputWeightsType.val.binary = 1; break; diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.h b/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.h index 604db427f7c..8fb50c17166 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernel_selector_params.h @@ -252,6 +252,8 @@ public: typedef union DataTypesKey_t { struct val_t { + uint32_t int4 : 1; + uint32_t uint4 : 1; uint32_t int8 : 1; uint32_t uint8 : 1; uint32_t int16 : 1; diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_base.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_base.cpp index a75d3546983..2d88dd2f942 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_base.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_base.cpp @@ -26,6 +26,8 @@ JitConstants FullyConnectedKernelBase::GetJitConstants(const fully_connected_par jit.AddConstants({MakeJitConstant("COMPRESSED_WEIGHTS", 1)}); if (params.weights.GetDType() == WeightsType::INT8 || params.weights.GetDType() == WeightsType::UINT8) { jit.AddConstants({MakeJitConstant("COMPRESSED_WEIGHTS_INT8", 1)}); + } else if (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) { + jit.AddConstants({MakeJitConstant("COMPRESSED_WEIGHTS_INT4", 1)}); } const size_t scale_groups_num = params.decompression_scale.Feature().v; @@ -35,12 +37,17 @@ JitConstants FullyConnectedKernelBase::GetJitConstants(const fully_connected_par jit.AddConstants({MakeJitConstant("DECOMPRESSION_SCALE_GROUPS_NUM", scale_groups_num)}); jit.AddConstants({MakeJitConstant("DECOMPRESSION_SCALE_GROUP_SIZE", scale_group_size)}); if (params.has_decompression_zp) { - const size_t zp_groups_num = params.decompression_zero_point.Feature().v; - const size_t zp_group_size = params.weights.IFM().v / params.decompression_zero_point.Feature().v; jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP_TERM", 1)}); - jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP", params.decompression_zero_point)}); - jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP_GROUPS_NUM", zp_groups_num)}); - jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP_GROUP_SIZE", zp_group_size)}); + if (params.scalar_zp) { + jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP_VALUE", params.zp_value)}); + jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP_SCALAR", 1)}); + } else { + const size_t zp_groups_num = params.decompression_zero_point.Feature().v; + const size_t zp_group_size = params.weights.IFM().v / params.decompression_zero_point.Feature().v; + jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP", params.decompression_zero_point)}); + jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP_GROUPS_NUM", zp_groups_num)}); + jit.AddConstants({MakeJitConstant("DECOMPRESSION_ZP_GROUP_SIZE", zp_group_size)}); + } } } @@ -119,7 +126,7 @@ KernelsData FullyConnectedKernelBase::GetCommonKernelsData(const Params ¶ms, int inputs_count = 1; if (newParams.compressed) { inputs_count++; - if (newParams.has_decompression_zp) + if (newParams.has_decompression_zp && !newParams.scalar_zp) inputs_count++; } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp index c272124627d..50a3d8649b1 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp @@ -6,6 +6,7 @@ #include #include +#include "common_types.h" static constexpr size_t simd = 16; @@ -38,6 +39,8 @@ ParamsKey FullyConnected_bf_tiled::GetSupportedKey() const { k.EnableOutputDataType(Datatype::F32); k.EnableOutputDataType(Datatype::INT8); k.EnableOutputDataType(Datatype::UINT8); + k.EnableInputWeightsType(WeightsType::UINT4); + k.EnableInputWeightsType(WeightsType::INT4); k.EnableInputWeightsType(WeightsType::F16); k.EnableInputWeightsType(WeightsType::F32); k.EnableInputLayout(DataLayout::bf); @@ -70,6 +73,7 @@ bool FullyConnected_bf_tiled::Validate(const Params& params, const optional_para auto& fc_params = static_cast(params); auto& input = fc_params.inputs[0]; auto& output = fc_params.outputs[0]; + auto& weights = fc_params.weights; // Block reads must be aligned to 4 bytes, for fp16 we can correct for offset misalignment, // but we need to ensure that batch pitch preserves alignment. @@ -103,6 +107,11 @@ bool FullyConnected_bf_tiled::Validate(const Params& params, const optional_para return false; } + auto wt = weights.GetDType(); + if ((wt == WeightsType::UINT4 || wt == WeightsType::INT4) && (weights.IFM().v % 2 != 0 || weights.OFM().v % 2 != 0)) { + return false; + } + return true; } @@ -150,6 +159,11 @@ bool TuneParamsSelector::VerifyTuneParams(const fully_connected_params& params, output_f = params.outputs[0].Y().v; } + if (params.compressed && + (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && + tparams.tile_ofm != 2) + return false; + auto batch_size = params.is_shape_agnostic ? Align(output_b, tparams.tile_b) : output_b; if (batch_size % (tparams.tile_b * tparams.dispatch_bsv) != 0) return false; @@ -201,7 +215,9 @@ FullyConnected_bf_tiled::GetAutoTuneParams(const fully_connected_params& params, while (max_tile_ofm * 2 * simd <= output_f && max_tile_ofm < 4) max_tile_ofm *= 2; - if (params.compressed && params.engineInfo.supports_immad) { + if (params.weights.GetDType() == WeightsType::UINT4 || params.weights.GetDType() == WeightsType::INT4) { + return selector.Default(tune_params(1, 2, 1, 4, 1, 1, EXE_MODE_DEFAULT)); + } else if (params.compressed && params.engineInfo.supports_immad) { return selector.Default(tune_params(1, 1, 1, 4, 1, 1, EXE_MODE_DEFAULT)); } else if (params.is_shape_agnostic) { // Use special tuning params for Gen12HP dGPUs, since these parameters demonstrate higher performance @@ -314,13 +330,24 @@ KernelsPriority FullyConnected_bf_tiled::GetKernelsPriority(const Params& params JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_params& params, const DispatchData& dispatchData) const { JitConstants jit = Parent::GetJitConstants(params, dispatchData); + size_t tile_k_ofm = dispatchData.tile_nk * dispatchData.tile_n; + size_t tile_k_ofm_packed = tile_k_ofm; + if (params.weights.GetDType() == WeightsType::UINT4 || params.weights.GetDType() == WeightsType::INT4) { + tile_k_ofm_packed /= 2; + + jit.Merge(make_int4_packed_type_jit_constant("INT4_PACKED_TYPE", params.weights.GetDType(), tile_k_ofm)); + const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; + if (scale_group_size % simd == 0) + jit.AddConstant(MakeJitConstant("DECOMPRESSION_SCALE_POST_OP", 1)); + } jit.AddConstant(MakeJitConstant("SIMD", simd)); jit.AddConstant(MakeJitConstant("TILE_B", dispatchData.tile_m)); jit.AddConstant(MakeJitConstant("TILE_OFM", dispatchData.tile_n)); jit.AddConstant(MakeJitConstant("TILE_IFM", dispatchData.tile_mk)); jit.AddConstant(MakeJitConstant("TILE_K", dispatchData.tile_nk)); - jit.AddConstant(MakeJitConstant("TILE_K_OFM", dispatchData.tile_nk * dispatchData.tile_n)); + jit.AddConstant(MakeJitConstant("TILE_K_OFM", tile_k_ofm)); + jit.AddConstant(MakeJitConstant("TILE_K_OFM_PACKED", tile_k_ofm_packed)); jit.AddConstant(MakeJitConstant("DISPATCH_BSV", dispatchData.tile_ms)); jit.AddConstant(MakeJitConstant("DISPATCH_FSV", dispatchData.tile_ns)); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bfyx_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bfyx_ref.cpp index 9b08e672452..d9b1a55e8eb 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bfyx_ref.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bfyx_ref.cpp @@ -4,6 +4,7 @@ #include +#include "common_types.h" #include "fully_connected_kernel_bfyx_ref.h" #include "kernel_selector_utils.h" @@ -22,6 +23,8 @@ ParamsKey FullyConnected_bfyx_Ref::GetSupportedKey() const { k.EnableInputWeightsType(WeightsType::F32); k.EnableInputWeightsType(WeightsType::UINT8); k.EnableInputWeightsType(WeightsType::INT8); + k.EnableInputWeightsType(WeightsType::UINT4); + k.EnableInputWeightsType(WeightsType::INT4); k.EnableAllInputLayout(); k.EnableDifferentInputWeightsTypes(); k.EnableDifferentTypes(); @@ -70,6 +73,11 @@ JitConstants FullyConnected_bfyx_Ref::GetJitConstants(const fully_connected_para jit.Merge(MakeTypeJitConstants(accumulator_dt, "ACCUMULATOR")); jit.Merge(MakeActivationJitConstants(params.activations, activation_dt, "_TYPED")); + auto wt = params.weights.GetDType(); + if (wt == WeightsType::UINT4 || wt == WeightsType::INT4) { + jit.Merge(make_int4_packed_type_jit_constant("INT4_PACKED_TYPE", wt, 2)); + } + if (!params.fused_ops.empty()) { std::vector idx_order = { "b", "ofm", "0", "0" }; if (params.outputs[0].GetLayout() == DataLayout::bfyx) diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_int4.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_int4.cpp new file mode 100644 index 00000000000..d443e2888c0 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_int4.cpp @@ -0,0 +1,68 @@ +// Copyright (C) 2023 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "reorder_weights_int4.h" +#include "kernel_selector_common.h" +#include "kernel_selector_params.h" +#include "kernel_selector_utils.h" + +namespace kernel_selector { + +ParamsKey ReorderWeightsKernelInt4::GetSupportedKey() const { + ParamsKey k; + k.EnableInputWeightsType(WeightsType::INT4); + k.EnableInputWeightsType(WeightsType::UINT4); + k.EnableOutputWeightsType(WeightsType::UINT4); + k.EnableOutputWeightsType(WeightsType::INT4); + k.EnableInputWeightsLayout(WeightsLayout::oiyx); + k.EnableInputWeightsLayout(WeightsLayout::ioyx); + k.EnableOutputWeightsLayout(WeightsLayout::os_iyx_osv32); + k.EnableOutputWeightsLayout(WeightsLayout::oiyx); + k.EnableTensorOffset(); + k.EnableTensorPitches(); + return k; +} + +KernelsData ReorderWeightsKernelInt4::GetKernelsData(const Params& params, const optional_params& options) const { + const reorder_weights_params& orgParams = static_cast(params); + return GetCommonKernelsData(orgParams, options); +} + +ReorderWeightsKernelInt4::DispatchData ReorderWeightsKernelInt4::SetDefault(const reorder_weights_params& params) const { + DispatchData dispatchData; + + const auto& output = params.output; + + // Divide one of the dimensions by 2 to save with byte granularity + if (output.GetLayout() == WeightsLayout::os_iyx_osv32) { + dispatchData.gws = { Align(output.OFM().v, 32) / 2, output.IFM().v, 1 }; + } else { + dispatchData.gws = { CeilDiv(output.LogicalSize(), 2), 1, 1 }; + } + dispatchData.lws = GetOptimalLocalWorkGroupSizes(dispatchData.gws, params.engineInfo); + + return dispatchData; +} + +bool ReorderWeightsKernelInt4::Validate(const Params& params, const optional_params& /*options*/) const { + const auto& p = static_cast(params); + const auto& input = p.input; + const auto& output = p.output; + + if (input.LogicalSize() != input.OFM().v * input.IFM().v || + output.LogicalSize() != output.OFM().v * output.IFM().v) { + return false; + } + + bool supported_case = input.GetLayout() == WeightsLayout::oiyx && output.GetLayout() == WeightsLayout::os_iyx_osv32; + supported_case |= input.GetLayout() == WeightsLayout::ioyx && output.GetLayout() == WeightsLayout::oiyx; + supported_case |= input.GetLayout() == WeightsLayout::ioyx && output.GetLayout() == WeightsLayout::os_iyx_osv32; + + return supported_case; +} + +KernelsPriority ReorderWeightsKernelInt4::GetKernelsPriority(const Params& /*params*/, const optional_params& /*options*/) const { + return DONT_USE_IF_HAVE_SOMETHING_ELSE; +} +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_int4.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_int4.h new file mode 100644 index 00000000000..d98fc9adee6 --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_int4.h @@ -0,0 +1,23 @@ +// Copyright (C) 2018-2023 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#pragma once + +#include "reorder_kernel_base.h" + +namespace kernel_selector { +class ReorderWeightsKernelInt4 : public ReorderKernelBase { +public: + ReorderWeightsKernelInt4() : ReorderKernelBase("reorder_weights_int4") {} + virtual ~ReorderWeightsKernelInt4() {} + + KernelsData GetKernelsData(const Params& params, const optional_params& options) const override; + KernelsPriority GetKernelsPriority(const Params& params, const optional_params& options) const override; + ParamsKey GetSupportedKey() const override; + DispatchData SetDefault(const reorder_weights_params& arg) const override; + +protected: + bool Validate(const Params& params, const optional_params& options) const override; +}; +} // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_kernel_selector.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_kernel_selector.cpp index 3006063045f..fbe010f111b 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_kernel_selector.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_weights_kernel_selector.cpp @@ -10,6 +10,7 @@ #include "reorder_weights_image_winograd_6x3_kernel.h" #include "reorder_weights_binary_kernel.h" #include "reorder_weights_opt.h" +#include "reorder_weights_int4.h" namespace kernel_selector { @@ -21,6 +22,7 @@ ReorderWeightsKernelSelector::ReorderWeightsKernelSelector() { Attach(); Attach(); Attach(); + Attach(); } KernelsData ReorderWeightsKernelSelector::GetBestKernels(const Params& params, const optional_params& options) const { diff --git a/src/plugins/intel_gpu/src/kernel_selector/weight_bias_params.h b/src/plugins/intel_gpu/src/kernel_selector/weight_bias_params.h index ad1c03cb8a7..236c1afeff3 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/weight_bias_params.h +++ b/src/plugins/intel_gpu/src/kernel_selector/weight_bias_params.h @@ -19,6 +19,8 @@ struct weight_bias_params : public base_params { bool compressed = false; bool has_decompression_zp = false; + bool scalar_zp = false; + float zp_value = 0.0f; DataTensor decompression_scale; DataTensor decompression_zero_point; diff --git a/src/plugins/intel_gpu/src/plugin/ops/fully_connected.cpp b/src/plugins/intel_gpu/src/plugin/ops/fully_connected.cpp index 2c2d42b842d..8a628809266 100644 --- a/src/plugins/intel_gpu/src/plugin/ops/fully_connected.cpp +++ b/src/plugins/intel_gpu/src/plugin/ops/fully_connected.cpp @@ -7,6 +7,7 @@ #include "intel_gpu/op/fully_connected.hpp" #include "intel_gpu/op/fully_connected_compressed.hpp" +#include "openvino/op/constant.hpp" #include "intel_gpu/primitives/fully_connected.hpp" #include "intel_gpu/primitives/reshape.hpp" @@ -34,17 +35,30 @@ static void CreateFullyConnectedCompressedOp(ProgramBuilder& p, const std::share auto scale_name = inputs[2].pid; auto zp_name = inputs.size() == 4 ? inputs[3].pid : ""; + float zp_value = 0.0f; + bool has_scalar_zp = false; + if (op->get_input_size() == 4) { + auto zp_const = std::dynamic_pointer_cast(op->get_input_node_shared_ptr(3)); + if (zp_const && ov::shape_size(zp_const->get_output_shape(0)) == 1) { + has_scalar_zp = true; + zp_value = zp_const->cast_vector()[0]; + } + } auto fc = cldnn::fully_connected(primitive_name, cldnn::input_info(input_name), weights_name, "", scale_name, - zp_name, + has_scalar_zp ? "" : zp_name, cldnn::element_type_to_data_type(op->get_output_element_type(0)), cldnn::padding(), op->get_input_partial_shape(0).size(), op->get_input_partial_shape(1).size()); + if (has_scalar_zp) { + fc.decompression_zero_point_scalar = zp_value; + } + p.add_primitive(*op, fc); } diff --git a/src/plugins/intel_gpu/src/plugin/plugin.cpp b/src/plugins/intel_gpu/src/plugin/plugin.cpp index 388269ddbb4..60d2077495f 100644 --- a/src/plugins/intel_gpu/src/plugin/plugin.cpp +++ b/src/plugins/intel_gpu/src/plugin/plugin.cpp @@ -125,7 +125,7 @@ std::shared_ptr Plugin::clone_and_transform_model(const std::shared_p GPU_DEBUG_IF(!debug_config->dump_graphs.empty()) { auto path_base = debug_config->dump_graphs + "/" + cloned_model->get_name(); ov::pass::Serialize(path_base + ".xml", path_base + ".bin").run_on_model(cloned_model); - ov::pass::VisualizeTree(path_base + ".dot").run_on_model(cloned_model); + ov::pass::VisualizeTree(path_base + ".svg").run_on_model(cloned_model); } transform_model(cloned_model, config); @@ -146,7 +146,7 @@ std::shared_ptr Plugin::clone_and_transform_model(const std::shared_p GPU_DEBUG_IF(!debug_config->dump_graphs.empty()) { auto path_base = debug_config->dump_graphs + "/" + cloned_model->get_name() + "_" + "transformed_func"; ov::pass::Serialize(path_base + ".xml", path_base + ".bin").run_on_model(cloned_model); - ov::pass::VisualizeTree(path_base + "_transformed.dot").run_on_model(cloned_model); + ov::pass::VisualizeTree(path_base + ".svg").run_on_model(cloned_model); } return cloned_model; } diff --git a/src/plugins/intel_gpu/src/plugin/transformations/convert_fc_to_compressed.cpp b/src/plugins/intel_gpu/src/plugin/transformations/convert_fc_to_compressed.cpp index 0ff0e1fd0bf..c4ae333f207 100644 --- a/src/plugins/intel_gpu/src/plugin/transformations/convert_fc_to_compressed.cpp +++ b/src/plugins/intel_gpu/src/plugin/transformations/convert_fc_to_compressed.cpp @@ -28,7 +28,9 @@ ConvertFullyConnectedToFullyConnectedCompressed::ConvertFullyConnectedToFullyCon auto compressed_constant = [](const ov::Output& output) { return (output.get_element_type() == ov::element::u8 || - output.get_element_type() == ov::element::i8) && + output.get_element_type() == ov::element::i8 || + output.get_element_type() == ov::element::u4 || + output.get_element_type() == ov::element::i4) && output.get_target_inputs().size() == 1; }; @@ -101,6 +103,7 @@ ConvertFullyConnectedToFullyConnectedCompressed::ConvertFullyConnectedToFullyCon std::shared_ptr fc_input_b = reshape_const_to_2d(pattern_map.at(weights_m).get_node_shared_ptr()); std::shared_ptr fc_input_scale = scale; std::shared_ptr fc_input_zp = optional_zero_point; + std::vector> result_nodes = {}; if (has_transpose) { const auto& transpose = pattern_map.at(transpose_m).get_node_shared_ptr(); std::shared_ptr transpose_const = pattern_map.at(transpose_const_m).get_node_shared_ptr(); @@ -112,9 +115,13 @@ ConvertFullyConnectedToFullyConnectedCompressed::ConvertFullyConnectedToFullyCon } fc_input_b = transpose->clone_with_new_inputs({ fc_input_b->output(0), transpose_const }); + result_nodes.push_back(fc_input_b); fc_input_scale = transpose->clone_with_new_inputs({ scale->output(0), transpose_const }); - if (with_zero_point) + result_nodes.push_back(fc_input_scale); + if (with_zero_point && ov::shape_size(optional_zero_point->output(0).get_shape()) > 1) { fc_input_zp = transpose->clone_with_new_inputs({ optional_zero_point->output(0), transpose_const }); + result_nodes.push_back(fc_input_zp); + } } std::shared_ptr new_fc = nullptr; @@ -131,8 +138,9 @@ ConvertFullyConnectedToFullyConnectedCompressed::ConvertFullyConnectedToFullyCon fc->get_output_type()); } + result_nodes.push_back(new_fc); new_fc->set_friendly_name(fc->get_friendly_name()); - ov::copy_runtime_info(m.get_matched_nodes(), new_fc); + ov::copy_runtime_info(m.get_matched_nodes(), result_nodes); ov::replace_node(fc, new_fc); return true; }; diff --git a/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp b/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp index ac567cd998f..99623a72c32 100644 --- a/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp +++ b/src/plugins/intel_gpu/src/plugin/transformations_pipeline.cpp @@ -233,7 +233,7 @@ void TransformationsPipeline::apply(std::shared_ptr func) { return !is_type(next_node); }); - manager.register_pass(ov::element::TypeVector{ov::element::u8}, true); + manager.register_pass(ov::element::TypeVector{ov::element::u8, ov::element::u4, ov::element::i4}, true); const bool keep_precision_sensitive_in_fp32_1 = true; const bool convert_input_output_precision = false; @@ -635,7 +635,6 @@ void TransformationsPipeline::apply(std::shared_ptr func) { return num_iter != 1; return num_iter >= 16; }); - manager.register_pass(true); manager.run_passes(func); } @@ -647,6 +646,10 @@ void TransformationsPipeline::apply(std::shared_ptr func) { manager.register_pass(); manager.register_pass(); + // This is supposed to be the last pass to ensure that we don't have name collisions until + // GPU plugin stops using friendly names for program creation + manager.register_pass(true); + manager.run_passes(func); } } diff --git a/src/plugins/intel_gpu/tests/functional/subgraph_tests/dynamic/matmul_weights_decompression.cpp b/src/plugins/intel_gpu/tests/functional/subgraph_tests/dynamic/matmul_weights_decompression.cpp index 75bdb9f0ec7..1830e07cda9 100644 --- a/src/plugins/intel_gpu/tests/functional/subgraph_tests/dynamic/matmul_weights_decompression.cpp +++ b/src/plugins/intel_gpu/tests/functional/subgraph_tests/dynamic/matmul_weights_decompression.cpp @@ -14,10 +14,10 @@ using namespace ov::test; namespace SubgraphTestsDefinitions { /* - * Subtract_const(U8/NF4/U4) - * / - * Weights(U8/NF4/U4) Convert(F32) - * | / + * Subtract_const(U8/NF4/U4/I4) + * / + * Weights(U8/NF4/U4/I4) Convert(F32) + * | / * Convert(F32) Reshape(optional) * \ / Multiply_const(F32) * Subtract(optional) / @@ -50,6 +50,7 @@ using MatmulWeightsDecompressionParams = std::tuple>; // additional config class MatmulWeightsDecompression : public testing::WithParamInterface, public SubgraphBaseTest { @@ -61,6 +62,7 @@ public: bool transpose; bool decompression_sub; bool reshape_on_decompression; + bool per_tensor_zp; std::map additional_config; std::tie(shape_params, @@ -69,6 +71,7 @@ public: transpose, decompression_sub, reshape_on_decompression, + per_tensor_zp, additional_config) = obj.param; std::ostringstream result; @@ -80,6 +83,7 @@ public: result << "transpose_weights=" << transpose << "_"; result << "decompression_subtract=" << decompression_sub << "_"; result << "reshape_on_decompression=" << reshape_on_decompression << "_"; + result << "per_tensor_zp=" << per_tensor_zp << "_"; result << "config=("; for (const auto& configEntry : additional_config) { @@ -98,7 +102,8 @@ protected: const ov::element::Type weights_precision, const bool transpose_weights, const bool add_subtract, - const bool reshape_on_decompression) { + const bool reshape_on_decompression, + const bool per_tensor_zp) { ov::ParameterVector params{std::make_shared(data_precision, data_shape)}; const auto weights_subgraph = init_compressed_weights_subgraph(weights_shape, group_size, @@ -106,7 +111,8 @@ protected: weights_precision, transpose_weights, add_subtract, - reshape_on_decompression); + reshape_on_decompression, + per_tensor_zp); auto mat_mul = std::make_shared(params[0], weights_subgraph); return std::make_shared(NodeVector{mat_mul}, params, "MatmulWeightsDecompression"); @@ -118,7 +124,8 @@ protected: const ov::element::Type weights_precision, const bool transpose_weights, const bool add_subtract, - const bool reshape_on_decompression_constant) { + const bool reshape_on_decompression_constant, + const bool per_tensor_zp) { auto transpose_if_necessary = [&](const ov::Shape& shape) { auto result_shape = shape; if (transpose_weights) @@ -168,10 +175,14 @@ protected: if (reshape_on_decompression_constant) scaleshift_const_shape.erase(std::remove(scaleshift_const_shape.begin(), scaleshift_const_shape.end(), 1), scaleshift_const_shape.end()); if (add_subtract) { - auto shift_tensor = ov::test::utils::create_and_fill_tensor(weights_precision, scaleshift_const_shape); + auto shift_tensor_shape = per_tensor_zp ? ov::Shape{1} : scaleshift_const_shape; + auto shift_tensor = ov::test::utils::create_and_fill_tensor(weights_precision, shift_tensor_shape); + if (per_tensor_zp && weights_precision.bitwidth() == 4) { + static_cast(shift_tensor.data())[0] = 0x88; + } auto shift_const = std::make_shared(shift_tensor); std::shared_ptr shift_convert = std::make_shared(shift_const, data_precision); - if (reshape_on_decompression_constant) { + if (reshape_on_decompression_constant && !per_tensor_zp) { auto shift_reshape_const = ov::opset10::Constant::create(ov::element::i32, {scaleshift_target_shape.size()}, scaleshift_target_shape); auto shift_reshape = std::make_shared(shift_convert, shift_reshape_const, false); shift_convert = shift_reshape; @@ -179,7 +190,13 @@ protected: mul_parent = std::make_shared(weights_convert, shift_convert); } - auto scale_tensor = ov::test::utils::create_and_fill_tensor(data_precision, scaleshift_const_shape, 1, -0.5, 10000); + auto scale_tensor = ov::test::utils::create_and_fill_tensor(data_precision, scaleshift_const_shape, 1, -0.5, 30000); + for (size_t i = 0; i < scale_tensor.get_size(); i++) { + if (data_precision == ov::element::f16) + scale_tensor.data()[i] /= ov::float16(16.f); + else if (data_precision == ov::element::f32) + scale_tensor.data()[i] /= 16.f; + } std::shared_ptr scale_const = std::make_shared(scale_tensor); if (reshape_on_decompression_constant) { auto scale_reshape_const = ov::opset10::Constant::create(ov::element::i32, {scaleshift_target_shape.size()}, scaleshift_target_shape); @@ -214,6 +231,7 @@ protected: bool transpose_weights; bool decompression_sub; bool reshape_on_decompression; + bool per_tensor_zp; std::map additional_config; std::tie(shape_params, @@ -222,6 +240,7 @@ protected: transpose_weights, decompression_sub, reshape_on_decompression, + per_tensor_zp, additional_config) = GetParam(); configuration.insert(additional_config.begin(), additional_config.end()); @@ -236,18 +255,14 @@ protected: weights_precision, transpose_weights, decompression_sub, - reshape_on_decompression); + reshape_on_decompression, + per_tensor_zp); if (activations_precision == ov::element::f16) { - auto weights_size = ov::shape_size(shape_params.weights_shape); - auto weights_input_channels = weights_size / (transpose_weights ? shape_params.weights_shape[0] : shape_params.weights_shape.back()); - // Absolute values range during accumulation may be quite big ( > 200) so fp16 representation & math error is larger than default threshold - if (weights_input_channels > 2048) { - abs_threshold = 4.0f; - } else { - abs_threshold = 1.0f; - } + abs_threshold = 1.0f; + } else { + abs_threshold = 1e-4f; } } @@ -285,12 +300,12 @@ TEST_P(MatmulWeightsDecompression, CompareWithRefs) { namespace { const std::vector activations_precisions = {ov::element::f32, ov::element::f16}; -const std::vector weights_precisions = {ov::element::u8}; +const std::vector weights_precisions = {ov::element::u8, ov::element::u4, ov::element::i4}; +const std::vector transpose_weights = {true, false}; const std::vector input_shapes_basic = { {{{-1, -1, -1}, {{1, 4, 16}, {10, 16, 16}}}, {16, 32}}, {{{}, {{1, 4, 16}}}, {16, 32}, 2ul}, {{{}, {{1, 4, 16}}}, {1, 16, 32}}, - {{{}, {{10, 40, 496}}}, {1, 496, 240}}, {{{}, {{1, 4, 48}}}, {48, 256}}, {{{}, {{11, 339, 377}}}, {377, 335}} }; @@ -300,9 +315,10 @@ INSTANTIATE_TEST_SUITE_P(smoke_MatMulCompressedWeights_basic, ::testing::Combine(::testing::ValuesIn(input_shapes_basic), ::testing::ValuesIn(weights_precisions), ::testing::ValuesIn(activations_precisions), + ::testing::ValuesIn(transpose_weights), ::testing::Values(true), ::testing::Values(true), - ::testing::Values(true), + ::testing::Values(false), ::testing::Values(std::map())), MatmulWeightsDecompression::get_test_case_name); @@ -310,6 +326,7 @@ const std::vector input_shapes_corner_cases_basic = { {{{-1, -1, -1}, {{1, 4, 16}}}, {1, 16, 32}}, {{{-1, -1, -1}, {{1, 4, 16}}}, {16, 32}}, {{{-1, -1, 16}, {{1, 4, 16}}}, {16, 32}, 4}, + {{{-1, 16}, {{4, 16}}}, {16, 32}, 4}, }; const std::vector input_shapes_corner_cases_big = { {{{-1, -1, -1}, {{10, 40, 480}, {11, 40, 480}}}, {1, 480, 256}}, @@ -318,9 +335,9 @@ const std::vector input_shapes_corner_cases_big = { {{{-1, 4096}, {{1, 4096}}}, {4096, 4096}, 128}, }; -const std::vector transpose_weights = {true, false}; const std::vector add_decompression_sub = {true, false}; const std::vector reshape_on_decompression = {true, false}; +const std::vector per_tensor_zp = {true, false}; INSTANTIATE_TEST_SUITE_P(smoke_MatMulCompressedWeights_corner_cases_basic, MatmulWeightsDecompression, @@ -330,6 +347,7 @@ INSTANTIATE_TEST_SUITE_P(smoke_MatMulCompressedWeights_corner_cases_basic, ::testing::ValuesIn(transpose_weights), ::testing::ValuesIn(add_decompression_sub), ::testing::ValuesIn(reshape_on_decompression), + ::testing::ValuesIn(per_tensor_zp), ::testing::Values(std::map{})), MatmulWeightsDecompression::get_test_case_name); @@ -341,6 +359,7 @@ INSTANTIATE_TEST_SUITE_P(MatMulCompressedWeights_corner_cases_big, ::testing::ValuesIn(transpose_weights), ::testing::ValuesIn(add_decompression_sub), ::testing::ValuesIn(reshape_on_decompression), + ::testing::ValuesIn(per_tensor_zp), ::testing::Values(std::map{})), MatmulWeightsDecompression::get_test_case_name); } // namespace diff --git a/src/plugins/intel_gpu/tests/unit/test_cases/hash_key_gpu_test.cpp b/src/plugins/intel_gpu/tests/unit/test_cases/hash_key_gpu_test.cpp index 68cfc542377..17447be4266 100644 --- a/src/plugins/intel_gpu/tests/unit/test_cases/hash_key_gpu_test.cpp +++ b/src/plugins/intel_gpu/tests/unit/test_cases/hash_key_gpu_test.cpp @@ -71,11 +71,11 @@ public: const auto primitive_hash = primitve->hash(); const auto params_hash = primitve->type->get_fake_aligned_params(*prim_inst->get_impl_params()).hash(); if (!engine.get_device_info().supports_immad) { - ASSERT_EQ(primitive_hash, 6924775129729406941UL); - ASSERT_EQ(params_hash, 8142839956977133460UL); + ASSERT_EQ(primitive_hash, 14259723886449306729UL); + ASSERT_EQ(params_hash, 1637150664489130388UL); } else { - ASSERT_EQ(primitive_hash, 6924775129729406941UL); - ASSERT_EQ(params_hash, 9266224209991282259UL); + ASSERT_EQ(primitive_hash, 14259723886449306729UL); + ASSERT_EQ(params_hash, 6343702278017463925UL); } } diff --git a/src/plugins/intel_gpu/tests/unit/transformations/convert_fc_to_compressed_test.cpp b/src/plugins/intel_gpu/tests/unit/transformations/convert_fc_to_compressed_test.cpp new file mode 100644 index 00000000000..bfed37f571d --- /dev/null +++ b/src/plugins/intel_gpu/tests/unit/transformations/convert_fc_to_compressed_test.cpp @@ -0,0 +1,237 @@ +// Copyright (C) 2023 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +#include "common_test_utils/ov_test_utils.hpp" + +#include "openvino/core/model.hpp" +#include "openvino/pass/manager.hpp" +#include "openvino/op/transpose.hpp" +#include "openvino/op/convert.hpp" +#include "openvino/op/multiply.hpp" +#include "openvino/op/parameter.hpp" +#include "openvino/op/constant.hpp" +#include "openvino/op/reshape.hpp" +#include "openvino/op/result.hpp" +#include "openvino/op/subtract.hpp" +#include "intel_gpu/op/fully_connected.hpp" +#include "intel_gpu/op/fully_connected_compressed.hpp" + +#include "plugin/transformations/convert_fc_to_compressed.hpp" + +#include + +using namespace testing; +using namespace ov::intel_gpu; + +namespace ov { +namespace test { +namespace intel_gpu { + +TEST_F(TransformationTestsF, ConvertFCToCompressed1) { + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u8, ov::Shape{ 32, 16 }, { 1 }); + auto convert = std::make_shared(weights_const, ov::element::f32); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 1 }, { 1 }); + auto scale = std::make_shared(convert, scale_const); + auto fc = std::make_shared(input1, scale); + + model = std::make_shared(ov::NodeVector{ fc }, ov::ParameterVector{ input1 }); + manager.register_pass(); + } + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u8, ov::Shape{ 32, 16 }, { 1 }); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 1 }, { 1 }); + auto fc_compressed = std::make_shared(input1, weights_const, scale_const); + + model_ref = std::make_shared(ov::NodeVector{ fc_compressed }, ov::ParameterVector{ input1 }); + } +} + +TEST_F(TransformationTestsF, ConvertFCToCompressed2) { + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u8, ov::Shape{ 32, 16 }, { 1 }); + auto convert = std::make_shared(weights_const, ov::element::f32); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 1 }, { 1 }); + auto sub = std::make_shared(convert, zp_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 1 }, { 1 }); + auto scale = std::make_shared(sub, scale_const); + auto fc = std::make_shared(input1, scale); + + model = std::make_shared(ov::NodeVector{ fc }, ov::ParameterVector{ input1 }); + manager.register_pass(); + } + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u8, ov::Shape{ 32, 16 }, { 1 }); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 1 }, { 1 }); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 1 }, { 1 }); + auto fc_compressed = std::make_shared(input1, weights_const, scale_const, zp_const); + + model_ref = std::make_shared(ov::NodeVector{ fc_compressed }, ov::ParameterVector{ input1 }); + } +} + +TEST_F(TransformationTestsF, ConvertFCToCompressed3) { + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u8, ov::Shape{ 32, 4, 4 }, { 1 }); + auto convert = std::make_shared(weights_const, ov::element::f32); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 4, 1 }, { 1 }); + auto sub = std::make_shared(convert, zp_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 4, 1 }, { 1 }); + auto scale = std::make_shared(sub, scale_const); + auto reshape_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { -1, 16 }); + auto reshape = std::make_shared(scale, reshape_const, false); + auto fc = std::make_shared(input1, reshape); + + model = std::make_shared(ov::NodeVector{ fc }, ov::ParameterVector{ input1 }); + manager.register_pass(); + } + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u8, ov::Shape{ 32, 16 }, { 1 }); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 4 }, { 1 }); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 4 }, { 1 }); + auto fc_compressed = std::make_shared(input1, weights_const, scale_const, zp_const); + + model_ref = std::make_shared(ov::NodeVector{ fc_compressed }, ov::ParameterVector{ input1 }); + } +} + +TEST_F(TransformationTestsF, ConvertFCToCompressed4) { + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 32, 4, 4 }, { 1 }); + auto convert = std::make_shared(weights_const, ov::element::f32); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 1, 1, 1 }, { 1 }); + auto sub = std::make_shared(convert, zp_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 4, 1 }, { 1 }); + auto scale = std::make_shared(sub, scale_const); + auto reshape_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { -1, 16 }); + auto reshape = std::make_shared(scale, reshape_const, false); + auto fc = std::make_shared(input1, reshape); + + model = std::make_shared(ov::NodeVector{ fc }, ov::ParameterVector{ input1 }); + manager.register_pass(); + } + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 32, 16 }, { 1 }); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 32, 4 }, { 1 }); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 1, 1 }, { 1 }); + auto fc_compressed = std::make_shared(input1, weights_const, scale_const, zp_const); + + model_ref = std::make_shared(ov::NodeVector{ fc_compressed }, ov::ParameterVector{ input1 }); + } +} + +TEST_F(TransformationTestsF, ConvertFCToCompressed5) { + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 4, 4, 32 }, { 1 }); + auto convert = std::make_shared(weights_const, ov::element::f32); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 1, 1, 1 }, { 1 }); + auto sub = std::make_shared(convert, zp_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 4, 1, 32 }, { 1 }); + auto scale = std::make_shared(sub, scale_const); + auto reshape_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 16, -1 }); + auto reshape = std::make_shared(scale, reshape_const, false); + auto transpose_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose = std::make_shared(reshape, transpose_const); + auto fc = std::make_shared(input1, transpose); + + model = std::make_shared(ov::NodeVector{ fc }, ov::ParameterVector{ input1 }); + manager.register_pass(); + } + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 16, 32 }, { 1 }); + auto transpose_weights_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_weights = std::make_shared(weights_const, transpose_weights_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 4, 32 }, { 1 }); + auto transpose_scale_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_scale = std::make_shared(scale_const, transpose_scale_const); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 1, 1 }, { 1 }); + auto fc_compressed = std::make_shared(input1, transpose_weights, transpose_scale, zp_const); + + model_ref = std::make_shared(ov::NodeVector{ fc_compressed }, ov::ParameterVector{ input1 }); + } +} + +TEST_F(TransformationTestsF, ConvertFCToCompressed6) { + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 4, 4, 32 }, { 1 }); + auto convert = std::make_shared(weights_const, ov::element::f32); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 4, 1, 32 }, { 1 }); + auto sub = std::make_shared(convert, zp_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 4, 1, 32 }, { 1 }); + auto scale = std::make_shared(sub, scale_const); + auto reshape_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 16, -1 }); + auto reshape = std::make_shared(scale, reshape_const, false); + auto transpose_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose = std::make_shared(reshape, transpose_const); + auto fc = std::make_shared(input1, transpose); + + model = std::make_shared(ov::NodeVector{ fc }, ov::ParameterVector{ input1 }); + manager.register_pass(); + } + { + auto input1 = std::make_shared(ov::element::f32, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 16, 32 }, { 1 }); + auto transpose_weights_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_weights = std::make_shared(weights_const, transpose_weights_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 4, 32 }, { 1 }); + auto transpose_scale_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_scale = std::make_shared(scale_const, transpose_scale_const); + auto zp_const = ov::op::v0::Constant::create(ov::element::f32, ov::Shape{ 4, 32 }, { 1 }); + auto transpose_zp_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_zp = std::make_shared(zp_const, transpose_zp_const); + auto fc_compressed = std::make_shared(input1, transpose_weights, transpose_scale, transpose_zp); + + model_ref = std::make_shared(ov::NodeVector{ fc_compressed }, ov::ParameterVector{ input1 }); + } +} + +TEST_F(TransformationTestsF, ConvertFCToCompressed7) { + { + auto input1 = std::make_shared(ov::element::f16, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 4, 4, 32 }, { 1 }); + auto convert = std::make_shared(weights_const, ov::element::f16); + auto zp_const = ov::op::v0::Constant::create(ov::element::f16, ov::Shape{ 4, 1, 32 }, { 1 }); + auto sub = std::make_shared(convert, zp_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f16, ov::Shape{ 4, 1, 32 }, { 1 }); + auto scale = std::make_shared(sub, scale_const); + auto reshape_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 16, -1 }); + auto reshape = std::make_shared(scale, reshape_const, false); + auto transpose_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose = std::make_shared(reshape, transpose_const); + auto fc = std::make_shared(input1, transpose); + + model = std::make_shared(ov::NodeVector{ fc }, ov::ParameterVector{ input1 }); + manager.register_pass(); + } + { + auto input1 = std::make_shared(ov::element::f16, ov::PartialShape{ -1, 16 }); + auto weights_const = ov::op::v0::Constant::create(ov::element::u4, ov::Shape{ 16, 32 }, { 1 }); + auto transpose_weights_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_weights = std::make_shared(weights_const, transpose_weights_const); + auto scale_const = ov::op::v0::Constant::create(ov::element::f16, ov::Shape{ 4, 32 }, { 1 }); + auto transpose_scale_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_scale = std::make_shared(scale_const, transpose_scale_const); + auto zp_const = ov::op::v0::Constant::create(ov::element::f16, ov::Shape{ 4, 32 }, { 1 }); + auto transpose_zp_const = ov::op::v0::Constant::create(ov::element::i32, ov::Shape{ 2 }, { 1, 0 }); + auto transpose_zp = std::make_shared(zp_const, transpose_zp_const); + auto fc_compressed = std::make_shared(input1, transpose_weights, transpose_scale, transpose_zp); + + model_ref = std::make_shared(ov::NodeVector{ fc_compressed }, ov::ParameterVector{ input1 }); + } +} + +} // namespace intel_gpu +} // namespace test +} // namespace ov