diff --git a/src/common/low_precision_transformations/include/low_precision/layer_transformation.hpp b/src/common/low_precision_transformations/include/low_precision/layer_transformation.hpp index 788e61e11c88..950be3320678 100644 --- a/src/common/low_precision_transformations/include/low_precision/layer_transformation.hpp +++ b/src/common/low_precision_transformations/include/low_precision/layer_transformation.hpp @@ -27,7 +27,7 @@ namespace pass { namespace low_precision { namespace precision_set { LP_TRANSFORMATIONS_API const std::vector& get_int8_support(); - LP_TRANSFORMATIONS_API const std::vector& get_fp8_support(); + LP_TRANSFORMATIONS_API const std::vector& get_low_bit_float_support(); LP_TRANSFORMATIONS_API const std::vector& get_int8_int16_int32_support(); } // namespace precision_set @@ -59,7 +59,7 @@ class LP_TRANSFORMATIONS_API DataPrecision { element::i8, element::u8, element::i16, element::u16, element::i32, element::u32, - element::f8e4m3, element::f8e5m2, + element::f8e4m3, element::f8e5m2, element::f4e2m1, }; return lowPrecision.find(precision) != lowPrecision.end(); } diff --git a/src/common/low_precision_transformations/src/layer_transformation.cpp b/src/common/low_precision_transformations/src/layer_transformation.cpp index 17e6c66791fd..dc189a7a03f6 100644 --- a/src/common/low_precision_transformations/src/layer_transformation.cpp +++ b/src/common/low_precision_transformations/src/layer_transformation.cpp @@ -29,11 +29,11 @@ const std::vector& precision_set::get_int8_support() { return int8_support; } -const std::vector& precision_set::get_fp8_support() { - static const std::vector fp8_support = { - ov::element::f8e4m3, ov::element::f8e5m2, +const std::vector& precision_set::get_low_bit_float_support() { + static const std::vector low_bit_float_support = { + ov::element::f8e4m3, ov::element::f8e5m2, ov::element::f4e2m1, }; - return fp8_support; + return low_bit_float_support; } const std::vector& precision_set::get_int8_int16_int32_support() { diff --git a/src/common/transformations/src/transformations/fp16_compression/mark_decompression_convert_constant_folding.cpp b/src/common/transformations/src/transformations/fp16_compression/mark_decompression_convert_constant_folding.cpp index a04dfdfefc6c..714f4fb6122b 100644 --- a/src/common/transformations/src/transformations/fp16_compression/mark_decompression_convert_constant_folding.cpp +++ b/src/common/transformations/src/transformations/fp16_compression/mark_decompression_convert_constant_folding.cpp @@ -143,6 +143,7 @@ MarkCompressedFloatConstants::MarkCompressedFloatConstants() { return false; if (const_node->get_output_element_type(0) != element::f16 && const_node->get_output_element_type(0) != element::bf16 && + const_node->get_output_element_type(0) != element::f4e2m1 && const_node->get_output_element_type(0) != element::f8e4m3 && const_node->get_output_element_type(0) != element::f8e5m2 && const_node->get_output_element_type(0) != element::f8e8m0) diff --git a/src/plugins/intel_gpu/src/graph/debug_helper.cpp b/src/plugins/intel_gpu/src/graph/debug_helper.cpp index 86a0a377e46f..308157829387 100644 --- a/src/plugins/intel_gpu/src/graph/debug_helper.cpp +++ b/src/plugins/intel_gpu/src/graph/debug_helper.cpp @@ -284,6 +284,8 @@ void log_memory_to_file(memory::ptr mem, layout data_layout, stream& stream, std dump(actual_mem, stream, file_stream, dump_raw); else if (mem_dt == cldnn::data_types::f8e4m3) dump(actual_mem, stream, file_stream, dump_raw); + else if (mem_dt == cldnn::data_types::f4e2m1) + dump(actual_mem, stream, file_stream, dump_raw); else if (mem_dt == cldnn::data_types::f8e8m0) dump(actual_mem, stream, file_stream, dump_raw); else if (mem_dt == cldnn::data_types::boolean) diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl/dynamic_quantize.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl/dynamic_quantize.cpp index 826f8924100b..e5e3574dbd0e 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl/dynamic_quantize.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl/dynamic_quantize.cpp @@ -80,6 +80,7 @@ attach_dynamic_quantize_impl::attach_dynamic_quantize_impl() { data_types::f16, data_types::i8, data_types::u8, + data_types::f4e2m1, data_types::f8e4m3, data_types::f8e5m2, data_types::f8e8m0, 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 7d4c259a40eb..da797cfa0cd0 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 @@ -266,6 +266,8 @@ kernel_selector::data_type to_data_type(data_types dt) { return kernel_selector::data_type::F32; case cldnn::data_types::bf16: return kernel_selector::data_type::BF16; + case cldnn::data_types::f4e2m1: + return kernel_selector::data_type::F4E2M1; case cldnn::data_types::f8e4m3: return kernel_selector::data_type::F8E4M3; case cldnn::data_types::f8e5m2: @@ -301,6 +303,8 @@ data_types from_data_type(kernel_selector::data_type dt) { return cldnn::data_types::f16; case kernel_selector::data_type::F32: return cldnn::data_types::f32; + case kernel_selector::data_type::F4E2M1: + return cldnn::data_types::f4e2m1; case kernel_selector::data_type::F8E4M3: return cldnn::data_types::f8e4m3; case kernel_selector::data_type::F8E5M2: @@ -330,6 +334,8 @@ kernel_selector::weights_type to_weights_type(data_types dt) { return kernel_selector::weights_type::INT32; case cldnn::data_types::bf16: return kernel_selector::weights_type::BF16; + case cldnn::data_types::f4e2m1: + return kernel_selector::weights_type::F4E2M1; case cldnn::data_types::f8e4m3: return kernel_selector::weights_type::F8E4M3; case cldnn::data_types::f8e5m2: @@ -357,6 +363,8 @@ data_types from_weights_type(kernel_selector::weights_type dt) { return data_types::f32; case kernel_selector::weights_type::INT32: return data_types::i32; + case kernel_selector::weights_type::F4E2M1: + return data_types::f4e2m1; case kernel_selector::weights_type::F8E4M3: return data_types::f8e4m3; case kernel_selector::weights_type::F8E5M2: diff --git a/src/plugins/intel_gpu/src/graph/impls/ocl_v2/utils/jitter.cpp b/src/plugins/intel_gpu/src/graph/impls/ocl_v2/utils/jitter.cpp index f5f708dc21df..0e0c80195ab1 100644 --- a/src/plugins/intel_gpu/src/graph/impls/ocl_v2/utils/jitter.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/ocl_v2/utils/jitter.cpp @@ -320,6 +320,18 @@ JitConstants make_type_jit_constants(const std::string& name, const ov::element: type_size = "4"; is_fp = true; break; + case ov::element::f4e2m1: + type = "fp4e2m1_t"; + max_val = "(fp4e2m1_t){as_uchar((uchar)0x7)}"; // 6.0 + min_val = "(fp4e2m1_t){as_uchar((uchar)0xF)}"; // -6.0 + val_one = "(fp4e2m1_t){as_uchar((uchar)0x2)}"; + val_zero = "(fp4e2m1_t){as_uchar((uchar)0x0)}"; + to_type = "_convert_fp4e2m1_t(v)"; + to_type_sat = "_convert_fp4e2m1_t_sat(v)"; + as_type = "as_fp4e2m1_t(v)"; + type_size = "0.5f"; + is_fp = true; + break; case ov::element::f8e4m3: type = "fp8e4m3_t"; max_val = "(fp8e4m3_t){as_char((char)0x7E)}"; // 448.0 diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp index 031feb602530..20f014c98405 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp @@ -71,7 +71,7 @@ struct fully_connected_onednn : typed_primitive_onednn_impl { } const auto input_dt = instance.get_input_layout(0).data_type; - const bool is_dyn_quan_input = cldnn::one_of(input_dt, {data_types::i8, data_types::u8, data_types::f8e4m3, data_types::f8e5m2}); + const bool is_dyn_quan_input = cldnn::one_of(input_dt, {data_types::i8, data_types::u8, data_types::f4e2m1, data_types::f8e4m3, data_types::f8e5m2}); if (is_dyn_quan_input && prim->activation_scale.is_valid()) { const auto activation_scale_idx = idx++; @@ -312,7 +312,7 @@ struct fully_connected_onednn : typed_primitive_onednn_impl { } const auto input_dt = impl_params->get_input_layout(0).data_type; - const bool is_dyn_quan_input = cldnn::one_of(input_dt, {data_types::i8, data_types::u8, data_types::f8e4m3, data_types::f8e5m2}); + const bool is_dyn_quan_input = cldnn::one_of(input_dt, {data_types::i8, data_types::u8, data_types::f4e2m1, data_types::f8e4m3, data_types::f8e5m2}); if (is_dyn_quan_input && dynamic_quantized_activation) { auto src_scale_idx = ++idx; auto partial_shape = impl_params->get_input_layout(0).get_partial_shape(); @@ -360,7 +360,7 @@ struct fully_connected_onednn : typed_primitive_onednn_impl { if (prim->compressed_weights) { const auto input_dt = impl_params.get_input_layout(0).data_type; - const bool is_dyn_quan_input = cldnn::one_of(input_dt, {data_types::i8, data_types::u8, data_types::f8e4m3, data_types::f8e5m2}); + const bool is_dyn_quan_input = cldnn::one_of(input_dt, {data_types::i8, data_types::u8, data_types::f4e2m1, data_types::f8e4m3, data_types::f8e5m2}); if (is_dyn_quan_input) { OPENVINO_ASSERT(prim->input_size <= 3, "[GPU] Dynamic quantization for 4D matmul is not implemented"); } else { diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.hpp b/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.hpp index 8052aabb07f5..d856703b83a9 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.hpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.hpp @@ -62,8 +62,8 @@ struct FullyConnectedImplementationManager : public ImplementationManager { one_of(in0_dt, {data_types::f16, data_types::bf16, data_types::f32, data_types::i8, data_types::u8}) && one_of(wei_dt, {data_types::u8, data_types::i8, data_types::u4, data_types::i4}) && one_of(out_dt, {data_types::f16, data_types::bf16, data_types::f32, data_types::u8, data_types::i8}); - const bool fp_compressed_case = fc_prim->compressed_weights && one_of(in0_dt, {data_types::f8e4m3, data_types::f8e5m2}) && - one_of(wei_dt, {data_types::f8e4m3, data_types::f8e5m2}) && one_of(out_dt, {data_types::f16, data_types::f32}); + const bool fp_compressed_case = fc_prim->compressed_weights && one_of(in0_dt, {data_types::f8e4m3, data_types::f4e2m1, data_types::f8e5m2}) && + one_of(wei_dt, {data_types::f8e4m3, data_types::f4e2m1, data_types::f8e5m2}) && one_of(out_dt, {data_types::f16, data_types::f32}); if (!f16f16_case && !bf16bf16_case && !f32f32_case && !u8s8_case && !compressed_case && !fp_compressed_case) LOG_AND_RETURN_FALSE(node); diff --git a/src/plugins/intel_gpu/src/graph/impls/onednn/utils.cpp b/src/plugins/intel_gpu/src/graph/impls/onednn/utils.cpp index 6d830c3c9e4a..e24e5de2df0e 100644 --- a/src/plugins/intel_gpu/src/graph/impls/onednn/utils.cpp +++ b/src/plugins/intel_gpu/src/graph/impls/onednn/utils.cpp @@ -142,6 +142,7 @@ dnnl::memory::data_type convert_data_type(cldnn::data_types dt) { case cldnn::data_types::i32: return dnnl::memory::data_type::s32; case cldnn::data_types::i4: return dnnl::memory::data_type::s4; case cldnn::data_types::u4: return dnnl::memory::data_type::u4; + case cldnn::data_types::f4e2m1: return dnnl::memory::data_type::f4_e2m1; case cldnn::data_types::f8e4m3: return dnnl::memory::data_type::f8_e4m3; case cldnn::data_types::f8e5m2: return dnnl::memory::data_type::f8_e5m2; case cldnn::data_types::f8e8m0: return dnnl::memory::data_type::e8m0; @@ -258,6 +259,7 @@ int64_t get_offset(const cldnn::layout& l, dnnl::memory::desc&& desc) { switch (desc.get_data_type()) { case dnnl::memory::data_type::s4: case dnnl::memory::data_type::u4: + case dnnl::memory::data_type::f4_e2m1: return offset / 2; case dnnl::memory::data_type::s8: case dnnl::memory::data_type::u8: diff --git a/src/plugins/intel_gpu/src/graph/primitive_inst.cpp b/src/plugins/intel_gpu/src/graph/primitive_inst.cpp index 01d06fc58679..5943b4837c9c 100644 --- a/src/plugins/intel_gpu/src/graph/primitive_inst.cpp +++ b/src/plugins/intel_gpu/src/graph/primitive_inst.cpp @@ -544,7 +544,7 @@ void primitive_inst::update_shape() { if (get_node().is_type() && get_flag(ExecutionFlags::SHAPE_CHANGED)) { auto &layout = _impl_params->get_output_layout(0); - OPENVINO_ASSERT(one_of(layout.data_type, {data_types::f16, data_types::i8, data_types::u8, data_types::f8e4m3, data_types::f8e5m2}), + OPENVINO_ASSERT(one_of(layout.data_type, {data_types::f16, data_types::i8, data_types::u8, data_types::f4e2m1, data_types::f8e4m3, data_types::f8e5m2}), "[GPU] Unsupported data type of dynamic_quantize: ", layout.data_type); if (layout.data_type == data_types::f16) set_can_be_optimized(true); diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_opt.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_opt.cl index 80bed3e059b0..c1a23c1c68be 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_opt.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_opt.cl @@ -3,17 +3,26 @@ // #define IS_F8 (F8E5M2_OUTPUT || F8E4M3_OUTPUT) +#define IS_F8_F4 (IS_F8 || F4E2M1_OUTPUT) #include "include/batch_headers/fetch_data.cl" -#if IS_F8 +#if IS_F8_F4 +#include "include/batch_headers/common.cl" #include "include/f8_utils.cl" +#include "include/f4_utils.cl" +#endif + +#if F4E2M1_OUTPUT +#define ELEMENTS_PER_BYTE 2 +#else +#define ELEMENTS_PER_BYTE 1 #endif #if OUTPUT_DIMS != 4 && OUTPUT_DIMS != 2 #error "dynamic_quantize_gpu_opt.cl: Unsupported output dimension" #endif -#if IS_F8 +#if IS_F8_F4 #define SCALE_TYPE float #define TO_SCALE_TYPE(x) _convert_float(x) #define ACT_MIN_VAL 0.000000059604645h // min half dtype val @@ -38,6 +47,18 @@ #define AS_TYPE_N(type, n, x) AS_TYPE_N_(type, n, x) #define AS_INPUT_TYPE_N(x) AS_TYPE_N(INPUT0_TYPE, VEC_SIZE, x) +#if VEC_SIZE == 2 +#define VSTORE_F4(vec, off, ptr) (*((ptr) + (off)) = (vec)) +#elif VEC_SIZE == 4 +#define VSTORE_F4(vec, off, ptr) vstore2(vec, off, ptr) +#elif VEC_SIZE == 8 +#define VSTORE_F4(vec, off, ptr) vstore4(vec, off, ptr) +#elif VEC_SIZE == 16 +#define VSTORE_F4(vec, off, ptr) vstore8(vec, off, ptr) +#else +#error "Unsupported VEC_SIZE for F4 packing" +#endif + #if GENERATE_PRECOMPUTED_REDUCTION #define FOR_PRECOMPUTED_REDUCTION(x) x #else @@ -68,14 +89,14 @@ KERNEL(dynamic_quantize_gpu_opt)( const uint b = get_global_id(0); const uint f_grp = get_global_id(1); const uint input_offset = INPUT0_GET_INDEX(b, f_grp * QUANTIZE_GROUP_SIZE, 0, 0); - const uint output_offset = OUTPUT_GET_INDEX(b, f_grp * QUANTIZE_GROUP_SIZE, 0, 0); + const uint output_offset = (OUTPUT_GET_INDEX(b, f_grp * QUANTIZE_GROUP_SIZE, 0, 0)) / ELEMENTS_PER_BYTE; #else const uint bf = get_global_id(0); const uint b = bf / INPUT0_FEATURE_NUM; const uint f = bf % INPUT0_FEATURE_NUM; const uint y_grp = get_global_id(1); const uint input_offset = INPUT0_GET_INDEX(b, f, y_grp * QUANTIZE_GROUP_SIZE, 0); - const uint output_offset = OUTPUT_GET_INDEX(b, f, y_grp * QUANTIZE_GROUP_SIZE, 0); + const uint output_offset = (OUTPUT_GET_INDEX(b, f, y_grp * QUANTIZE_GROUP_SIZE, 0)) / ELEMENTS_PER_BYTE; #endif const uint quantize_block = QUANTIZE_GROUP_SIZE / 4; @@ -101,7 +122,12 @@ KERNEL(dynamic_quantize_gpu_opt)( #endif // MXFP unroll_for (uint i = 0 ; i < quantize_block; ++i) { -#if IS_F8 +#if F4E2M1_OUTPUT + float4 val_f = convert_float4(input_0[i]) * (MAKE_VECTOR_TYPE(SCALE_TYPE, 4))quan_scale; + val_f = clamp(val_f, -_convert_float(OUTPUT_VAL_MAX), _convert_float(OUTPUT_VAL_MAX)); + quantized_value[i] = TO_TYPE_N_SAT(OUTPUT_TYPE, 4, val_f); + vstore2(quantized_value[i].data, 0, (uchar*)(&output[output_offset + i * 2])); +#elif IS_F8 quantized_value[i] = TO_TYPE_N_SAT(OUTPUT_TYPE, 4, convert_float4(input_0[i]) * (MAKE_VECTOR_TYPE(SCALE_TYPE, 4))quan_scale); vstore4(quantized_value[i].data, 0, (char*)(&output[output_offset + i * 4])); #else @@ -116,7 +142,7 @@ KERNEL(dynamic_quantize_gpu_opt)( #else const uint output_idx = OUTPUT1_GET_INDEX(b, f, y_grp, 0); #endif - output_scale[output_idx] = TO_OUTPUT1_TYPE(1.0h / quan_scale); + output_scale[output_idx] = TO_OUTPUT1_TYPE(1.0f / quan_scale); #if !(IS_MXFP) FOR_PRECOMPUTED_REDUCTION(output_precomputed_reduction[output_idx] = precomputed_reduction); @@ -151,10 +177,10 @@ KERNEL(dynamic_quantize_gpu_opt)( const uint blockid = (uint)get_global_id(1) % (QUANTIZE_GROUP_SIZE / VEC_SIZE / SIMD); #if OUTPUT_DIMS == 2 const uint input_offset = INPUT0_GET_INDEX (b, f_grp * QUANTIZE_GROUP_SIZE + VEC_SIZE * sglid, 0, 0); - const uint output_offset = OUTPUT_GET_INDEX(b, f_grp * QUANTIZE_GROUP_SIZE + VEC_SIZE * sglid, 0, 0); + const uint output_offset = (OUTPUT_GET_INDEX(b, f_grp * QUANTIZE_GROUP_SIZE + VEC_SIZE * sglid, 0, 0)) / ELEMENTS_PER_BYTE; #else const uint input_offset = INPUT0_GET_INDEX (0, b, f_grp * QUANTIZE_GROUP_SIZE + VEC_SIZE * sglid, 0); - const uint output_offset = OUTPUT_GET_INDEX(0, b, f_grp * QUANTIZE_GROUP_SIZE + VEC_SIZE * sglid, 0); + const uint output_offset = (OUTPUT_GET_INDEX(0, b, f_grp * QUANTIZE_GROUP_SIZE + VEC_SIZE * sglid, 0)) / ELEMENTS_PER_BYTE; #endif const uint block_size = SIMD * VEC_SIZE; @@ -233,15 +259,22 @@ KERNEL(dynamic_quantize_gpu_opt)( SCALE_TYPE scale = TO_SCALE_TYPE(OUTPUT_VAL_MAX) / max_value; #endif -#if IS_F8 - val = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, TO_TYPE_N(SCALE_TYPE, VEC_SIZE, val) * (MAKE_VECTOR_TYPE(SCALE_TYPE, VEC_SIZE))scale); + MAKE_VECTOR_TYPE(SCALE_TYPE, VEC_SIZE) val_scaled = TO_TYPE_N(SCALE_TYPE, VEC_SIZE, val) * (MAKE_VECTOR_TYPE(SCALE_TYPE, VEC_SIZE))scale; +#if F4E2M1_OUTPUT + val_scaled = clamp(val_scaled, -TO_SCALE_TYPE(OUTPUT_VAL_MAX), TO_SCALE_TYPE(OUTPUT_VAL_MAX)); + MAKE_VECTOR_TYPE(OUTPUT_TYPE, VEC_SIZE) out_f4 = TO_TYPE_N_SAT(OUTPUT_TYPE, VEC_SIZE, val_scaled); + VSTORE_F4(out_f4.data, 0, (uchar*)(&output[output_offset + (blockid * block_size) / ELEMENTS_PER_BYTE])); +#elif IS_F8 + val = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, val_scaled); MAKE_VECTOR_TYPE(OUTPUT_TYPE, VEC_SIZE) out = TO_TYPE_N_SAT(OUTPUT_TYPE, VEC_SIZE, val); VSTORE_N(out.data, 0, (char*)(&output[output_offset + (blockid * block_size)])); #elif ASYMMETRIC_QUANTIZATION + val = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, val_scaled); val *= scale; val += zp; VSTORE_N(CAT(CONVERT_UCHAR_N, _rte)(val), 0, output + output_offset + (blockid * block_size)); #else // i8 symmetric + val = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, val_scaled); val *= scale; VSTORE_N(CAT(CONVERT_CHAR_N, _rte)(val), 0, output + output_offset + (blockid * block_size)); #endif @@ -322,6 +355,7 @@ KERNEL(dynamic_quantize_gpu_opt)( const uint b_offset = bf * INPUT0_FEATURE_PITCH; #endif const uint offset = b_offset + VEC_SIZE * sglid; + const uint output_byte_offset = (b_offset + VEC_SIZE * sglid) / ELEMENTS_PER_BYTE; const uint iteration = ALIGNED_BLOCK_NUM / BLOCK_NUM; @@ -390,14 +424,21 @@ KERNEL(dynamic_quantize_gpu_opt)( if ((local_id * iteration + i) >= TOTAL_BLOCK_NUM) continue; - val[i] = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, TO_TYPE_N(SCALE_TYPE, VEC_SIZE, val[i]) * (MAKE_VECTOR_TYPE(SCALE_TYPE, VEC_SIZE))scale); -#if IS_F8 + MAKE_VECTOR_TYPE(SCALE_TYPE, VEC_SIZE) val_scaled = TO_TYPE_N(SCALE_TYPE, VEC_SIZE, val[i]) * (MAKE_VECTOR_TYPE(SCALE_TYPE, VEC_SIZE))scale; +#if F4E2M1_OUTPUT + val_scaled = clamp(val_scaled, -TO_SCALE_TYPE(OUTPUT_VAL_MAX), TO_SCALE_TYPE(OUTPUT_VAL_MAX)); + MAKE_VECTOR_TYPE(OUTPUT_TYPE, VEC_SIZE) out_f4 = TO_TYPE_N_SAT(OUTPUT_TYPE, VEC_SIZE, val_scaled); + VSTORE_F4(out_f4.data, 0, (uchar*)(&output[output_byte_offset + ((local_id * iteration + i) * block_size) / ELEMENTS_PER_BYTE])); +#elif IS_F8 + val[i] = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, val_scaled); MAKE_VECTOR_TYPE(OUTPUT_TYPE, VEC_SIZE) out = TO_TYPE_N_SAT(OUTPUT_TYPE, VEC_SIZE, val[i]); VSTORE_N(out.data, 0, (char*)(&output[offset + ((local_id * iteration + i) * block_size)])); #elif ASYMMETRIC_QUANTIZATION + val[i] = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, val_scaled); val[i] += zp; VSTORE_N(CAT(CONVERT_UCHAR_N, _rte)(val[i]), 0, output + offset + ((local_id * iteration + i) * block_size)); #else // i8 symmetric + val[i] = TO_TYPE_N(INPUT0_TYPE, VEC_SIZE, val_scaled); VSTORE_N(CAT(CONVERT_CHAR_N, _rte)(val[i]), 0, output + offset + ((local_id * iteration + i) * block_size)); #endif } diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_ref.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_ref.cl index dee2976bd61d..129a92d93180 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_ref.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/dynamic_quantize_gpu_ref.cl @@ -3,15 +3,17 @@ // #define IS_F8 (F8E5M2_OUTPUT || F8E4M3_OUTPUT) +#define IS_F8_F4 (IS_F8 || F4E2M1_OUTPUT) #include "include/batch_headers/fetch_data.cl" -#if IS_F8 +#if IS_F8_F4 #include "include/f8_utils.cl" +#include "include/f4_utils.cl" #endif #define UINT64_MAX 0xFFFFFFFFFFFFFFFF -#if IS_F8 +#if IS_F8_F4 #define SCALE_TYPE float #define TO_SCALE_TYPE(x) _convert_float(x) #define TO_SCALE_TYPE_8(x) convert_float8(x) @@ -29,6 +31,9 @@ #elif F8E4M3_OUTPUT #define TO_OUTPUT_TYPE_CUSTOM(val) _convert_fp8e4m3_t_sat(val) #define TO_OUTPUT_VEC_TYPE_CUSTOM(val) _convert_fp8e4m3_t8_sat(val) +#elif F4E2M1_OUTPUT + #define TO_OUTPUT_TYPE_CUSTOM(val) _convert_fp4e2m1_t_sat(val) + #define TO_OUTPUT_VEC_TYPE_CUSTOM(val) _convert_fp4e2m1_t8_sat(val) #elif (ASYMMETRIC_QUANTIZATION && UNSIGNED_OUTPUT) #define TO_OUTPUT_TYPE_CUSTOM(val) convert_uchar_rte(val) #define TO_OUTPUT_VEC_TYPE_CUSTOM(val) convert_uchar8_rte(val) @@ -43,6 +48,12 @@ #define FOR_PRECOMPUTED_REDUCTION(x) #endif +#if F4E2M1_OUTPUT +#define ELEMENTS_PER_BYTE 2 +#else +#define ELEMENTS_PER_BYTE 1 +#endif + #if OUTPUT_DIMS != 4 #error "dynamic_quantize_gpu_ref.cl: Unsupported output dimension" #endif @@ -161,6 +172,7 @@ KERNEL(dynamic_quantize_gpu_ref)( #if GROUP_SIZE_DIM3 == 1 const uint in_offset = INPUT0_GET_INDEX(b + b_off, f + f_off, y + y_off, x); const uint out_offset = OUTPUT_GET_INDEX(b + b_off, f + f_off, y + y_off, x); + const uint byte_offset = out_offset / ELEMENTS_PER_BYTE; half val = input[in_offset]; val *= scale; @@ -168,11 +180,24 @@ KERNEL(dynamic_quantize_gpu_ref)( val += zp; #endif OUTPUT_TYPE ival = TO_OUTPUT_TYPE_CUSTOM(val); - output[out_offset] = ival; +#if F4E2M1_OUTPUT + { + volatile __global uint* output_u32 = (volatile __global uint*)output; + uint main_idx = out_offset / 8; + uint sub_idx = out_offset % 8; + uint shift = sub_idx * 4; + uint val_u32 = (uint)(ival.data & 0x0F); + atomic_and(&output_u32[main_idx], ~(0x0F << shift)); + atomic_or (&output_u32[main_idx], (val_u32 << shift)); + } +#else + output[out_offset] = ival; +#endif FOR_PRECOMPUTED_REDUCTION(precomputed_reduction += ival); #else // GROUP_SIZE_DIM3 != 1 const uint in_offset = INPUT0_GET_INDEX(b + b_off, f + f_off, y + y_off, 0); const uint out_offset = OUTPUT_GET_INDEX(b + b_off, f + f_off, y + y_off, 0); + const uint byte_offset = out_offset / ELEMENTS_PER_BYTE; int x; for (x = 0; x < INPUT0_SIZE_X / 8; x++) { half8 val = as_half8(vload8(0, (ushort*)input + in_offset + x * 8)); @@ -180,11 +205,13 @@ KERNEL(dynamic_quantize_gpu_ref)( #if ASYMMETRIC_QUANTIZATION val += zp; #endif -#if IS_F8 - vstore8(TO_OUTPUT_VEC_TYPE_CUSTOM(val).data, 0, (char*)(&output[out_offset + x * 8])); +#if F4E2M1_OUTPUT + vstore4(TO_OUTPUT_VEC_TYPE_CUSTOM(val).data, 0, (uchar*)(&output[byte_offset + x * 4])); +#elif IS_F8 + vstore8(TO_OUTPUT_VEC_TYPE_CUSTOM(val).data, 0, (char*)(&output[byte_offset + x * 8])); #else MAKE_VECTOR_TYPE(OUTPUT_TYPE, 8) ival = TO_OUTPUT_VEC_TYPE_CUSTOM(val); - vstore8(ival, 0, output + out_offset + x * 8); + vstore8(ival, 0, output + byte_offset + x * 8); FOR_PRECOMPUTED_REDUCTION(precomputed_reduction += ((int)ival[0]) + ival[1] + ival[2] + ival[3] + ival[4] + ival[5] + ival[6] + ival[7]); #endif } @@ -196,7 +223,20 @@ KERNEL(dynamic_quantize_gpu_ref)( val += zp; #endif OUTPUT_TYPE ival = TO_OUTPUT_TYPE_CUSTOM(val); - output[out_offset + x] = ival; + uint out_idx = out_offset + x; +#if F4E2M1_OUTPUT + { + volatile __global uint* output_u32 = (volatile __global uint*)output; + uint main_idx = out_idx / 8; + uint sub_idx = out_idx % 8; + uint shift = sub_idx * 4; + uint val_u32 = (uint)(ival.data & 0x0F); + atomic_and(&output_u32[main_idx], ~(0x0F << shift)); + atomic_or (&output_u32[main_idx], (val_u32 << shift)); + } +#else + output[out_idx] = ival; +#endif FOR_PRECOMPUTED_REDUCTION(precomputed_reduction += ival); } #endif @@ -204,7 +244,7 @@ KERNEL(dynamic_quantize_gpu_ref)( } } - output_scale[scale_idx] = TO_OUTPUT1_TYPE(1.0h / scale); + output_scale[scale_idx] = TO_OUTPUT1_TYPE(1.0f / scale); FOR_PRECOMPUTED_REDUCTION(output_precomputed_reduction[scale_idx] = precomputed_reduction); #if ASYMMETRIC_QUANTIZATION && GROUP_SCALES_WITH_ZP output_scale[scale_idx + 1] = zp; diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/common.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/common.cl index 5bcd72ef1ff8..589d26e14ed0 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/common.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/common.cl @@ -40,6 +40,7 @@ #define AS_TYPE_PREFIX_uchar as_ #define AS_TYPE_PREFIX_char as_ +#define AS_TYPE_PREFIX_fp4e2m1_t _as_ #define AS_TYPE_PREFIX_fp8e5m2_t _as_ #define AS_TYPE_PREFIX_fp8e4m3_t _as_ #define AS_TYPE_PREFIX_fp8e8m0_t _as_ diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/f4_utils.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/f4_utils.cl new file mode 100644 index 000000000000..b1cabced5b0a --- /dev/null +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/f4_utils.cl @@ -0,0 +1,398 @@ +// Copyright (C) 2026 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 +// + +// TODO: Replace `_intel_convert*` with builtins when ready, current implementations are copied from XeTLA: +#ifndef OV_GPU_OCL_F4_UTILS_H +#define OV_GPU_OCL_F4_UTILS_H + +#define FP4_MASK 0x0F +#define FP4_SHIFT 4 +#define FP4_HIGH_MASK 0xF0 + +uchar _intel_convert_f16_to_f4(half val) { + half val_fp16 = val; + ushort *p = (ushort *)&val_fp16; + ushort src = p[0]; + const ushort src_exp_size = 5; + const ushort src_mant_size = 10; + const ushort src_exp_bias = (1 << (src_exp_size - 1)) - 1; + const ushort src_exp_mask = (1 << src_exp_size) - 1; + const ushort src_mant_mask = (1 << src_mant_size) - 1; + const short max_exp_unbiased = 2; + const short min_exp_unbiased = 0; + const short exp_bias = 1; + const ushort exp_size = 2; + const ushort mant_size = 1; + const uchar max_val = 0x7; + + ushort src_sign = src >> (src_exp_size + src_mant_size); + ushort src_exp = (src >> src_mant_size) & src_exp_mask; + short src_exp_unbiased = src_exp - src_exp_bias; + ushort src_mant = src & src_mant_mask; + + bool is_src_inf_nan = src_exp == 0x1f; + bool is_overflow = (src_exp_unbiased > max_exp_unbiased) + || ((src_exp_unbiased == max_exp_unbiased) && (src_mant > 0x0340)); + bool is_zero = (src_exp_unbiased < (min_exp_unbiased - mant_size)); + bool is_denorm = (src_exp_unbiased < min_exp_unbiased) && (!is_zero); + + uchar dst_val; + if (is_src_inf_nan || is_overflow) { + dst_val = max_val; + } else if (is_zero) { + dst_val = 0; + if (src_exp == 0xD && src_mant & 0x3FF) // if larger then 0.25 round to 0.5 + dst_val = 0x1; + } else if (is_denorm) { + dst_val = 0x1; + if (src_exp == 0xE && src_mant & 0x200) // if larger then 0.75 round to 1.0 + dst_val = 0x2; + } else { + ushort tail_size = src_mant_size - mant_size; + bool sticky_flag = (src_mant & ((1 << (tail_size - 1)) - 1)) != 0; + bool lsb_bit = src_mant & (1 << tail_size); + bool rnd_bit = src_mant & (1 << (tail_size - 1)); + bool carry = (lsb_bit && rnd_bit) || (rnd_bit && sticky_flag); + ushort src_m = (src_mant >> tail_size) + carry; + ushort src_e = src_exp_unbiased + exp_bias; + dst_val = (src_e << mant_size) + src_m; + if(dst_val > max_val){ + dst_val = max_val; + } + } + return (src_sign << (exp_size + mant_size) | dst_val) & FP4_MASK; +} + +half _intel_convert_fp4_to_f16(uchar val){ + static const ushort LUT[16] = { + 0x0000, 0x3800, 0x3c00, 0x3e00, 0x4000, 0x4200, 0x4400, 0x4600, + 0x8000, 0xb800, 0xbc00, 0xbe00, 0xc000, 0xc200, 0xc400, 0xc600 + }; + + ushort idx = val & FP4_MASK; + return as_half(LUT[idx]); +} + +static inline uchar pack_nibbles(uchar low, uchar high) { + return ((high << FP4_SHIFT) & FP4_HIGH_MASK) | (low & FP4_MASK); +} + +static inline uchar unpack_nibble_low(uchar packed) { + return packed & FP4_MASK; +} + +static inline uchar unpack_nibble_high(uchar packed) { + return (packed >> FP4_SHIFT) & FP4_MASK; +} + +static inline uchar2 unpack_nibbles(uchar packed) { + return (uchar2)(unpack_nibble_low(packed), unpack_nibble_high(packed)); +} + +static inline uchar pack_half_pair_to_fp4(half low, half high) { + return pack_nibbles(_intel_convert_f16_to_f4(low), _intel_convert_f16_to_f4(high)); +} + +static inline uchar pack_float_pair_to_fp4(float low, float high) { + return pack_half_pair_to_fp4((half)low, (half)high); +} + +static inline half unpack_fp4_low_to_half(uchar packed) { + return _intel_convert_fp4_to_f16(unpack_nibble_low(packed)); +} + +static inline half unpack_fp4_high_to_half(uchar packed) { + return _intel_convert_fp4_to_f16(unpack_nibble_high(packed)); +} + +static inline half2 unpack_fp4_to_half_pair(uchar packed) { + return (half2)(unpack_fp4_low_to_half(packed), unpack_fp4_high_to_half(packed)); +} + +typedef struct fp4e2m1_t { uchar data; } fp4e2m1_t; // f4 +typedef struct fp4e2m1_t1 { uchar data; } fp4e2m1_t1; +typedef struct fp4e2m1_t2 { uchar data; } fp4e2m1_t2; +typedef struct fp4e2m1_t3 { uchar2 data; } fp4e2m1_t3; +typedef struct fp4e2m1_t4 { uchar2 data; } fp4e2m1_t4; +typedef struct fp4e2m1_t8 { uchar4 data; } fp4e2m1_t8; +typedef struct fp4e2m1_t16 { uchar8 data; } fp4e2m1_t16; + +half __attribute__((overloadable)) _convert_half(fp4e2m1_t val) { + return unpack_fp4_low_to_half(val.data); +} +half __attribute__((overloadable)) _convert_half(fp4e2m1_t1 val) { + return unpack_fp4_low_to_half(val.data); +} +half2 __attribute__((overloadable)) _convert_half2(fp4e2m1_t2 val) { + return unpack_fp4_to_half_pair(val.data); +} +half3 __attribute__((overloadable)) _convert_half3(fp4e2m1_t3 val) { + return (half3)(unpack_fp4_low_to_half(val.data.s0), unpack_fp4_high_to_half(val.data.s0), + unpack_fp4_low_to_half(val.data.s1)); +} +half4 __attribute__((overloadable)) _convert_half4(fp4e2m1_t4 val) { + return (half4)( + unpack_fp4_to_half_pair(val.data.s0), + unpack_fp4_to_half_pair(val.data.s1) + ); +} +half8 __attribute__((overloadable)) _convert_half8(fp4e2m1_t8 val) { + return (half8)( + unpack_fp4_to_half_pair(val.data.s0), + unpack_fp4_to_half_pair(val.data.s1), + unpack_fp4_to_half_pair(val.data.s2), + unpack_fp4_to_half_pair(val.data.s3) + ); +} +half16 __attribute__((overloadable)) _convert_half16(fp4e2m1_t16 val) { + return (half16)( + unpack_fp4_to_half_pair(val.data.s0), unpack_fp4_to_half_pair(val.data.s1), + unpack_fp4_to_half_pair(val.data.s2), unpack_fp4_to_half_pair(val.data.s3), + unpack_fp4_to_half_pair(val.data.s4), unpack_fp4_to_half_pair(val.data.s5), + unpack_fp4_to_half_pair(val.data.s6), unpack_fp4_to_half_pair(val.data.s7) + ); +} + +float __attribute__((overloadable)) _convert_float(fp4e2m1_t val) { + return (float)_convert_half(val); +} +float __attribute__((overloadable)) _convert_float(fp4e2m1_t1 val) { + return (float)_convert_half(val); +} +float2 __attribute__((overloadable)) _convert_float2(fp4e2m1_t2 val) { + return convert_float2(_convert_half2(val)); +} +float3 __attribute__((overloadable)) _convert_float3(fp4e2m1_t3 val) { + return convert_float3(_convert_half3(val)); +} +float4 __attribute__((overloadable)) _convert_float4(fp4e2m1_t4 val) { + return convert_float4(_convert_half4(val)); +} +float8 __attribute__((overloadable)) _convert_float8(fp4e2m1_t8 val) { + return convert_float8(_convert_half8(val)); +} +float16 __attribute__((overloadable)) _convert_float16(fp4e2m1_t16 val) { + return convert_float16(_convert_half16(val)); +} + + +fp4e2m1_t __attribute__((overloadable)) _convert_fp4e2m1_t(half val) { + fp4e2m1_t res; + res.data = _intel_convert_f16_to_f4(val); + return res; +} +fp4e2m1_t1 __attribute__((overloadable)) _convert_fp4e2m1_t1(half val[1]) { + fp4e2m1_t1 res; + res.data = _intel_convert_f16_to_f4(val[0]); + return res; +} +fp4e2m1_t2 __attribute__((overloadable)) _convert_fp4e2m1_t2(half2 val) { + fp4e2m1_t2 res; + res.data = pack_half_pair_to_fp4(val.x, val.y); + return res; +} +fp4e2m1_t3 __attribute__((overloadable)) _convert_fp4e2m1_t3(half3 val) { + fp4e2m1_t3 res; + res.data.s0 = pack_half_pair_to_fp4(val.x, val.y); + res.data.s1 = _intel_convert_f16_to_f4(val.z); + return res; +} +fp4e2m1_t4 __attribute__((overloadable)) _convert_fp4e2m1_t4(half4 val) { + fp4e2m1_t4 res; + res.data.s0 = pack_half_pair_to_fp4(val.x, val.y); + res.data.s1 = pack_half_pair_to_fp4(val.z, val.w); + return res; +} +fp4e2m1_t8 __attribute__((overloadable)) _convert_fp4e2m1_t8(half8 val) { + fp4e2m1_t8 res; + res.data.s0 = pack_half_pair_to_fp4(val.s0, val.s1); + res.data.s1 = pack_half_pair_to_fp4(val.s2, val.s3); + res.data.s2 = pack_half_pair_to_fp4(val.s4, val.s5); + res.data.s3 = pack_half_pair_to_fp4(val.s6, val.s7); + return res; +} +fp4e2m1_t16 __attribute__((overloadable)) _convert_fp4e2m1_t16(half16 val) { + fp4e2m1_t16 res; + res.data.s0 = pack_half_pair_to_fp4(val.s0, val.s1); + res.data.s1 = pack_half_pair_to_fp4(val.s2, val.s3); + res.data.s2 = pack_half_pair_to_fp4(val.s4, val.s5); + res.data.s3 = pack_half_pair_to_fp4(val.s6, val.s7); + res.data.s4 = pack_half_pair_to_fp4(val.s8, val.s9); + res.data.s5 = pack_half_pair_to_fp4(val.sA, val.sB); + res.data.s6 = pack_half_pair_to_fp4(val.sC, val.sD); + res.data.s7 = pack_half_pair_to_fp4(val.sE, val.sF); + return res; +} + +fp4e2m1_t __attribute__((overloadable)) _convert_fp4e2m1_t(float val) { + fp4e2m1_t res; + res.data = _intel_convert_f16_to_f4((half)val); + return res; +} +fp4e2m1_t1 __attribute__((overloadable)) _convert_fp4e2m1_t1(float val[1]) { + fp4e2m1_t1 res; + res.data = _intel_convert_f16_to_f4((half)val[0]); + return res; +} +fp4e2m1_t2 __attribute__((overloadable)) _convert_fp4e2m1_t2(float2 val) { + fp4e2m1_t2 res; + res.data = pack_float_pair_to_fp4(val.x, val.y); + return res; +} +fp4e2m1_t3 __attribute__((overloadable)) _convert_fp4e2m1_t3(float3 val) { + fp4e2m1_t3 res; + res.data.s0 = pack_float_pair_to_fp4(val.x, val.y); + res.data.s1 = _intel_convert_f16_to_f4((half)val.z); + return res; +} +fp4e2m1_t4 __attribute__((overloadable)) _convert_fp4e2m1_t4(float4 val) { + fp4e2m1_t4 res; + res.data.s0 = pack_float_pair_to_fp4(val.x, val.y); + res.data.s1 = pack_float_pair_to_fp4(val.z, val.w); + return res; +} +fp4e2m1_t8 __attribute__((overloadable)) _convert_fp4e2m1_t8(float8 val) { + fp4e2m1_t8 res; + res.data.s0 = pack_float_pair_to_fp4(val.s0, val.s1); + res.data.s1 = pack_float_pair_to_fp4(val.s2, val.s3); + res.data.s2 = pack_float_pair_to_fp4(val.s4, val.s5); + res.data.s3 = pack_float_pair_to_fp4(val.s6, val.s7); + return res; +} +fp4e2m1_t16 __attribute__((overloadable)) _convert_fp4e2m1_t16(float16 val) { + fp4e2m1_t16 res; + res.data.s0 = pack_float_pair_to_fp4(val.s0, val.s1); + res.data.s1 = pack_float_pair_to_fp4(val.s2, val.s3); + res.data.s2 = pack_float_pair_to_fp4(val.s4, val.s5); + res.data.s3 = pack_float_pair_to_fp4(val.s6, val.s7); + res.data.s4 = pack_float_pair_to_fp4(val.s8, val.s9); + res.data.s5 = pack_float_pair_to_fp4(val.sA, val.sB); + res.data.s6 = pack_float_pair_to_fp4(val.sC, val.sD); + res.data.s7 = pack_float_pair_to_fp4(val.sE, val.sF); + return res; +} + +fp4e2m1_t __attribute__((overloadable)) _convert_fp4e2m1_t_sat(half val) { + return _convert_fp4e2m1_t(val); +} +fp4e2m1_t1 __attribute__((overloadable)) _convert_fp4e2m1_t1_sat(half val[1]) { + return _convert_fp4e2m1_t1(val); +} +fp4e2m1_t2 __attribute__((overloadable)) _convert_fp4e2m1_t2_sat(half2 val) { + return _convert_fp4e2m1_t2(val); +} +fp4e2m1_t3 __attribute__((overloadable)) _convert_fp4e2m1_t3_sat(half3 val) { + return _convert_fp4e2m1_t3(val); +} +fp4e2m1_t4 __attribute__((overloadable)) _convert_fp4e2m1_t4_sat(half4 val) { + return _convert_fp4e2m1_t4(val); +} +fp4e2m1_t8 __attribute__((overloadable)) _convert_fp4e2m1_t8_sat(half8 val) { + return _convert_fp4e2m1_t8(val); +} +fp4e2m1_t16 __attribute__((overloadable)) _convert_fp4e2m1_t16_sat(half16 val) { + return _convert_fp4e2m1_t16(val); +} + +fp4e2m1_t __attribute__((overloadable)) _convert_fp4e2m1_t_sat(float val) { + return _convert_fp4e2m1_t(val); +} +fp4e2m1_t1 __attribute__((overloadable)) _convert_fp4e2m1_t1_sat(float val[1]) { + return _convert_fp4e2m1_t1(val); +} +fp4e2m1_t2 __attribute__((overloadable)) _convert_fp4e2m1_t2_sat(float2 val) { + return _convert_fp4e2m1_t2(val); +} +fp4e2m1_t3 __attribute__((overloadable)) _convert_fp4e2m1_t3_sat(float3 val) { + return _convert_fp4e2m1_t3(val); +} +fp4e2m1_t4 __attribute__((overloadable)) _convert_fp4e2m1_t4_sat(float4 val) { + return _convert_fp4e2m1_t4(val); +} +fp4e2m1_t8 __attribute__((overloadable)) _convert_fp4e2m1_t8_sat(float8 val) { + return _convert_fp4e2m1_t8(val); +} +fp4e2m1_t16 __attribute__((overloadable)) _convert_fp4e2m1_t16_sat(float16 val) { + return _convert_fp4e2m1_t16(val); +} + +fp4e2m1_t __attribute__((overloadable)) as_fp4e2m1_t(uchar val) { + fp4e2m1_t res; + res.data = val; + return res; +} +fp4e2m1_t1 __attribute__((overloadable)) as_fp4e2m1_t1(uchar val[1]) { + fp4e2m1_t1 res; + res.data = val[0]; + return res; +} +fp4e2m1_t2 __attribute__((overloadable)) as_fp4e2m1_t2(uchar2 val) { + fp4e2m1_t2 res; + res.data = pack_nibbles(val.x, val.y); + return res; +} +fp4e2m1_t3 __attribute__((overloadable)) as_fp4e2m1_t3(uchar3 val) { + fp4e2m1_t3 res; + res.data.s0 = pack_nibbles(val.x, val.y); + res.data.s1 = val.z; + return res; +} +fp4e2m1_t4 __attribute__((overloadable)) as_fp4e2m1_t4(uchar4 val) { + fp4e2m1_t4 res; + res.data.s0 = pack_nibbles(val.x, val.y); + res.data.s1 = pack_nibbles(val.z, val.w); + return res; +} +fp4e2m1_t8 __attribute__((overloadable)) as_fp4e2m1_t8(uchar8 val) { + fp4e2m1_t8 res; + res.data.s0 = pack_nibbles(val.s0, val.s1); + res.data.s1 = pack_nibbles(val.s2, val.s3); + res.data.s2 = pack_nibbles(val.s4, val.s5); + res.data.s3 = pack_nibbles(val.s6, val.s7); + return res; +} +fp4e2m1_t16 __attribute__((overloadable)) as_fp4e2m1_t16(uchar16 val) { + fp4e2m1_t16 res; + res.data.s0 = pack_nibbles(val.s0, val.s1); + res.data.s1 = pack_nibbles(val.s2, val.s3); + res.data.s2 = pack_nibbles(val.s4, val.s5); + res.data.s3 = pack_nibbles(val.s6, val.s7); + res.data.s4 = pack_nibbles(val.s8, val.s9); + res.data.s5 = pack_nibbles(val.sA, val.sB); + res.data.s6 = pack_nibbles(val.sC, val.sD); + res.data.s7 = pack_nibbles(val.sE, val.sF); + return res; +} + +uchar __attribute__((overloadable)) _as_uchar(fp4e2m1_t val) { + return val.data; +} +uchar __attribute__((overloadable)) _as_uchar(fp4e2m1_t1 val) { + return val.data; +} +uchar2 __attribute__((overloadable)) _as_uchar2(fp4e2m1_t2 val) { + return unpack_nibbles(val.data); +} +uchar3 __attribute__((overloadable)) _as_uchar3(fp4e2m1_t3 val) { + return (uchar3)(unpack_nibbles(val.data.s0), unpack_nibble_low(val.data.s1)); +} +uchar4 __attribute__((overloadable)) _as_uchar4(fp4e2m1_t4 val) { + return (uchar4)(unpack_nibbles(val.data.s0), + unpack_nibbles(val.data.s1)); +} +uchar8 __attribute__((overloadable)) _as_uchar8(fp4e2m1_t8 val) { + return (uchar8)(unpack_nibbles(val.data.s0), + unpack_nibbles(val.data.s1), + unpack_nibbles(val.data.s2), + unpack_nibbles(val.data.s3)); +} +uchar16 __attribute__((overloadable)) _as_uchar16(fp4e2m1_t16 val) { + return (uchar16)(unpack_nibbles(val.data.s0), unpack_nibbles(val.data.s1), + unpack_nibbles(val.data.s2), unpack_nibbles(val.data.s3), + unpack_nibbles(val.data.s4), unpack_nibbles(val.data.s5), + unpack_nibbles(val.data.s6), unpack_nibbles(val.data.s7)); +} + +#endif + diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data.cl index d3bc565d1977..33b19934daa6 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data.cl +++ b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/reorder_data.cl @@ -2,11 +2,12 @@ // SPDX-License-Identifier: Apache-2.0 // -#define IS_F8 (F8E5M2_INPUT || F8E4M3_INPUT || F8E8M0_INPUT || F8E5M2_OUTPUT || F8E4M3_OUTPUT || F8E8M0_OUTPUT) +#define IS_LOW_BIT_FP (F4E2M1_INPUT || F8E5M2_INPUT || F8E4M3_INPUT || F8E8M0_INPUT || F4E2M1_OUTPUT || F8E5M2_OUTPUT || F8E4M3_OUTPUT || F8E8M0_OUTPUT) -#if IS_F8 +#if IS_LOW_BIT_FP #include "include/batch_headers/common.cl" #include "include/f8_utils.cl" +#include "include/f4_utils.cl" #endif #include "include/reshape_dims.cl" @@ -162,6 +163,16 @@ KERNEL (reorder_data)( OUTPUT_TYPE res = TO_OUTPUT_REORDER_TYPE(convert_as_uint4_float(input[uint4_idx], input_idx)); #elif (F8E5M2_INPUT || F8E4M3_INPUT || F8E8M0_INPUT) OUTPUT_TYPE res = TO_OUTPUT_REORDER_TYPE(_convert_float(input[input_idx])); + #elif defined F4E2M1_INPUT + // FP4 unpacking: 2 elements per byte + const uint byte_idx = input_idx / 2; + const uint sub_idx = input_idx % 2; + const uint shift = sub_idx * 4; + + uchar packed_byte = ((const __global uchar*)input)[byte_idx]; + uchar val_u8 = (packed_byte >> shift) & 0x0F; + + OUTPUT_TYPE res = TO_OUTPUT_REORDER_TYPE(_convert_float(as_fp4e2m1_t(val_u8))); #else CALC_TYPE res = TO_CALC_TYPE(input[input_idx]); #endif @@ -240,6 +251,25 @@ KERNEL (reorder_data)( res = __TO_OUTPUT_REORDER_TYPE(res); FUSED_OPS; output[output_idx] = FUSED_OPS_RESULT; + #elif defined(F8E5M2_OUTPUT) || defined(F8E4M3_OUTPUT) || defined(F8E8M0_OUTPUT) + // FP8 output: 1 element per byte, requires conversion function + OUTPUT_TYPE val_fp8_out = ACTIVATION_TYPED(OUTPUT_REORDER, __TO_OUTPUT_REORDER_TYPE(res), ACTIVATION_PARAMS_TYPED); + output[output_idx] = val_fp8_out; + #elif defined(F4E2M1_OUTPUT) + // FP4 packing: 2 elements per byte + OUTPUT_TYPE val_fp_out = ACTIVATION_TYPED(OUTPUT_REORDER, __TO_OUTPUT_REORDER_TYPE(res), ACTIVATION_PARAMS_TYPED); + half val_half_out = (half)val_fp_out; + fp4e2m1_t val_fp4_out = _convert_fp4e2m1_t(val_half_out); + uchar val_u8_out = val_fp4_out.data; + + volatile __global uint* output_u32 = (volatile __global uint*)output; + uint main_idx_out = output_idx / 8; + uint sub_idx_u32_out = output_idx % 8; + uint shift_u32_out = sub_idx_u32_out * 4; + uint val_u32_out = (uint)(val_u8_out & 0x0F); + + atomic_and(&output_u32[main_idx_out], ~(0x0F << shift_u32_out)); + atomic_or(&output_u32[main_idx_out], (val_u32_out << shift_u32_out)); #elif defined(INT4_OUTPUT) || defined(UINT4_OUTPUT) OUTPUT_TYPE val_char = __TO_OUTPUT_REORDER_TYPE(res); int val_i32 = convert_int(val_char); 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 37a34f83a80f..4258f916b4df 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/common_types.h +++ b/src/plugins/intel_gpu/src/kernel_selector/common_types.h @@ -131,6 +131,7 @@ enum class Datatype { F16, F32, BF16, + F4E2M1, F8E4M3, F8E5M2, F8E8M0, @@ -149,6 +150,7 @@ enum class WeightsType { INT4, INT32, BF16, + F4E2M1, F8E4M3, F8E5M2, F8E8M0, diff --git a/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp b/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp index f639bffd5a0f..a760137707dc 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/jitter.cpp @@ -1596,6 +1596,18 @@ JitConstants MakeTypeJitConstants(Datatype dataType, const std::string& macroNam type_size = "2"; is_fp = false; break; + case Datatype::F4E2M1: + type = "fp4e2m1_t"; + max_val = "(fp4e2m1_t){as_uchar((uchar)0x7)}"; // 6.0 + min_val = "(fp4e2m1_t){as_uchar((uchar)0xF)}"; // -6.0 + val_one = "(fp4e2m1_t){as_uchar((uchar)0x2)}"; + val_zero = "(fp4e2m1_t){as_uchar((uchar)0x0)}"; + to_type = "_convert_fp4e2m1_t(v)"; + to_type_sat = "_convert_fp4e2m1_t_sat(v)"; + as_type = "as_fp4e2m1_t(v)"; + type_size = "0.5f"; + is_fp = true; + break; case Datatype::F8E4M3: type = "fp8e4m3_t"; max_val = "(fp8e4m3_t){as_char((char)0x7E)}"; // 448.0 @@ -1685,6 +1697,8 @@ JitConstants MakeTypeJitConstants(WeightsType weightsType, const std::string& ma return MakeTypeJitConstants(Datatype::INT32, macroName); case WeightsType::BF16: return MakeTypeJitConstants(Datatype::BF16, macroName); + case WeightsType::F4E2M1: + return MakeTypeJitConstants(Datatype::F4E2M1, macroName); case WeightsType::F8E4M3: return MakeTypeJitConstants(Datatype::F8E4M3, macroName); case WeightsType::F8E5M2: 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 54fa4b16f446..416d0a13adc8 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 @@ -157,6 +157,7 @@ std::string toString(Datatype dType) { case Datatype::INT64: return "INT64"; case Datatype::F16: return "F16"; case Datatype::F32: return "F32"; + case Datatype::F4E2M1: return "F4E2M1"; case Datatype::F8E4M3: return "F8E4M3"; case Datatype::F8E5M2: return "F8E5M2"; case Datatype::F8E8M0: return "F8E8M0"; 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 af31dcd0539e..eeab8c597d3a 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 @@ -83,6 +83,9 @@ void ParamsKey::EnableInputDataType(Datatype dt) { case Datatype::BF16: key.inputType.val.BF16 = 1; break; + case Datatype::F4E2M1: + key.inputType.val.F4E2M1 = 1; + break; case Datatype::F8E4M3: key.inputType.val.F8E4M3 = 1; break; @@ -137,6 +140,9 @@ void ParamsKey::EnableOutputDataType(Datatype dt) { case Datatype::BF16: key.outputType.val.BF16 = 1; break; + case Datatype::F4E2M1: + key.outputType.val.F4E2M1 = 1; + break; case Datatype::F8E4M3: key.outputType.val.F8E4M3 = 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 c14b55efa803..d9737a27255b 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 @@ -256,6 +256,7 @@ class ParamsKey { uint32_t F16 : 1; uint32_t F32 : 1; uint32_t BF16 : 1; + uint32_t F4E2M1 : 1; uint32_t F8E4M3 : 1; uint32_t F8E5M2 : 1; uint32_t F8E8M0 : 1; diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_opt.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_opt.cpp index 987ae930c4af..53265770dade 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_opt.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_opt.cpp @@ -63,6 +63,7 @@ ParamsKey DynamicQuantizeKernelOpt::GetSupportedKey() const { k.EnableInputDataType(Datatype::F16); k.EnableOutputDataType(Datatype::UINT8); k.EnableOutputDataType(Datatype::INT8); + k.EnableOutputDataType(Datatype::F4E2M1); k.EnableOutputDataType(Datatype::F8E4M3); k.EnableOutputDataType(Datatype::F8E5M2); k.EnableOutputDataType(Datatype::F8E8M0); @@ -95,6 +96,7 @@ JitConstants DynamicQuantizeKernelOpt::GetJitConstants(const dynamic_quantize_pa jit.AddConstant(MakeJitConstant("MODE_SMALL_GS", static_cast(DynQuanMode::SMALL_GS))); jit.AddConstant(MakeJitConstant("MODE_LARGE_GS", static_cast(DynQuanMode::LARGE_GS))); jit.AddConstant(MakeJitConstant("MODE_PER_TOKEN", static_cast(DynQuanMode::PER_TOKEN))); + jit.AddConstant(MakeJitConstant("F4E2M1_OUTPUT", params.outputs[0].GetDType() == Datatype::F4E2M1 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E5M2_OUTPUT", params.outputs[0].GetDType() == Datatype::F8E5M2 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E4M3_OUTPUT", params.outputs[0].GetDType() == Datatype::F8E4M3 ? 1 : 0)); jit.AddConstant(MakeJitConstant("IS_MXFP", params.outputs[1].GetDType() == Datatype::F8E8M0 ? 1 : 0)); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_ref.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_ref.cpp index 4b36c2eded23..f59bb1c8385d 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_ref.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/dynamic_quantize/dynamic_quantize_kernel_ref.cpp @@ -12,6 +12,7 @@ ParamsKey DynamicQuantizeKernelRef::GetSupportedKey() const { k.EnableInputDataType(Datatype::F16); k.EnableOutputDataType(Datatype::INT8); k.EnableOutputDataType(Datatype::UINT8); + k.EnableOutputDataType(Datatype::F4E2M1); k.EnableOutputDataType(Datatype::F8E4M3); k.EnableOutputDataType(Datatype::F8E5M2); k.EnableOutputDataType(Datatype::F8E8M0); @@ -59,6 +60,7 @@ JitConstants DynamicQuantizeKernelRef::GetJitConstants(const dynamic_quantize_pa jit.AddConstant(MakeJitConstant("GENERATE_PRECOMPUTED_REDUCTION", params.generate_precomputed_reduction)); jit.AddConstant(MakeJitConstant("GROUP_SCALES_WITH_ZP", params.combine_scales_and_zp)); jit.AddConstant(MakeJitConstant("UNSIGNED_OUTPUT", params.outputs[0].GetDType() == Datatype::UINT8 ? 1 : 0)); + jit.AddConstant(MakeJitConstant("F4E2M1_OUTPUT", params.outputs[0].GetDType() == Datatype::F4E2M1 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E5M2_OUTPUT", params.outputs[0].GetDType() == Datatype::F8E5M2 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E4M3_OUTPUT", params.outputs[0].GetDType() == Datatype::F8E4M3 ? 1 : 0)); jit.AddConstant(MakeJitConstant("IS_MXFP", params.outputs[1].GetDType() == Datatype::F8E8M0 ? 1 : 0)); @@ -163,7 +165,7 @@ bool DynamicQuantizeKernelRef::Validate(const Params& params) const { DO_NOT_USE_THIS_KERNEL(params.layerID); const auto& dq_params = static_cast(params); - if (dq_params.generate_precomputed_reduction && cldnn::one_of(dq_params.outputs[0].GetDType(), {Datatype::F8E4M3, Datatype::F8E5M2})) { + if (dq_params.generate_precomputed_reduction && cldnn::one_of(dq_params.outputs[0].GetDType(), {Datatype::F4E2M1, Datatype::F8E4M3, Datatype::F8E5M2})) { DO_NOT_USE_THIS_KERNEL(params.layerID); } diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel.cpp index e41c6824ce7b..4b286ed986bb 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel.cpp @@ -20,6 +20,7 @@ ParamsKey ReorderKernelRef::GetSupportedKey() const { k.EnableInputDataType(Datatype::INT64); k.EnableInputDataType(Datatype::F16); k.EnableInputDataType(Datatype::F32); + k.EnableInputDataType(Datatype::F4E2M1); k.EnableInputDataType(Datatype::F8E4M3); k.EnableInputDataType(Datatype::F8E5M2); k.EnableInputDataType(Datatype::F8E8M0); @@ -35,6 +36,7 @@ ParamsKey ReorderKernelRef::GetSupportedKey() const { k.EnableOutputDataType(Datatype::UINT4); k.EnableOutputDataType(Datatype::INT4); k.EnableOutputDataType(Datatype::BF16); + k.EnableOutputDataType(Datatype::F4E2M1); k.EnableOutputDataType(Datatype::F8E4M3); k.EnableOutputDataType(Datatype::F8E5M2); k.EnableOutputDataType(Datatype::F8E8M0); @@ -90,9 +92,11 @@ JitConstants ReorderKernelRef::GetJitConstants(const reorder_params& params) con jit.AddConstant(MakeJitConstant("INT4_OUTPUT", true)); } + jit.AddConstant(MakeJitConstant("F4E2M1_INPUT", params.inputs[0].GetDType() == Datatype::F4E2M1 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E5M2_INPUT", params.inputs[0].GetDType() == Datatype::F8E5M2 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E4M3_INPUT", params.inputs[0].GetDType() == Datatype::F8E4M3 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E8M0_INPUT", params.inputs[0].GetDType() == Datatype::F8E8M0 ? 1 : 0)); + jit.AddConstant(MakeJitConstant("F4E2M1_OUTPUT", params.outputs[0].GetDType() == Datatype::F4E2M1 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E5M2_OUTPUT", params.outputs[0].GetDType() == Datatype::F8E5M2 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E4M3_OUTPUT", params.outputs[0].GetDType() == Datatype::F8E4M3 ? 1 : 0)); jit.AddConstant(MakeJitConstant("F8E8M0_OUTPUT", params.outputs[0].GetDType() == Datatype::F8E8M0 ? 1 : 0)); diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel_base.cpp b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel_base.cpp index 136994090dc5..9705d4d8e6ba 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel_base.cpp +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/reorder/reorder_kernel_base.cpp @@ -122,7 +122,7 @@ JitConstants ReorderKernelBase::GetJitConstants(const reorder_params& params) co if (params.inputs[0].GetDType() == Datatype::BF16 || params.inputs[0].GetDType() == Datatype::F8E8M0) { calc_type = Datatype::F32; } - if (params.inputs[0].GetDType() == Datatype::F8E4M3 || params.inputs[0].GetDType() == Datatype::F8E5M2) { + if (params.inputs[0].GetDType() == Datatype::F8E4M3 || params.inputs[0].GetDType() == Datatype::F8E5M2 || params.inputs[0].GetDType() == Datatype::F4E2M1) { calc_type = Datatype::F16; } Datatype output_reorder_type = useUshort ? Datatype::UINT16 : params.outputs[0].GetDType(); diff --git a/src/plugins/intel_gpu/src/plugin/common_utils.cpp b/src/plugins/intel_gpu/src/plugin/common_utils.cpp index 854c2fb08a5e..8341cabb905e 100644 --- a/src/plugins/intel_gpu/src/plugin/common_utils.cpp +++ b/src/plugins/intel_gpu/src/plugin/common_utils.cpp @@ -201,6 +201,7 @@ bool is_supported(ov::element::Type_t et) { case ov::element::Type_t::u32: return true; // converted to i32 case ov::element::Type_t::u64: return true; // converted to i32 case ov::element::Type_t::nf4: return false; + case ov::element::Type_t::f4e2m1: return true; case ov::element::Type_t::f8e4m3: return true; case ov::element::Type_t::f8e5m2: return true; case ov::element::Type_t::string: return false; diff --git a/src/plugins/intel_gpu/src/plugin/transformations/compressed_weights_pattern.hpp b/src/plugins/intel_gpu/src/plugin/transformations/compressed_weights_pattern.hpp index ed37b0b64606..65455df6210e 100644 --- a/src/plugins/intel_gpu/src/plugin/transformations/compressed_weights_pattern.hpp +++ b/src/plugins/intel_gpu/src/plugin/transformations/compressed_weights_pattern.hpp @@ -11,7 +11,8 @@ using namespace ov::pass::pattern; 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::u4 || output.get_element_type() == ov::element::i4 ||\ - output.get_element_type() == ov::element::f8e4m3 || output.get_element_type() == ov::element::f8e5m2);\ + output.get_element_type() == ov::element::f8e4m3 || output.get_element_type() == ov::element::f8e5m2 ||\ + output.get_element_type() == ov::element::f8e8m0 || output.get_element_type() == ov::element::f4e2m1);\ };\ \ auto reshape_squeeze = [](const ov::Output& output) {\ 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 9b987d5626c0..bc753b723cd8 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 @@ -66,7 +66,7 @@ using MatmulWeightsDecompressionParams = std::tuple; class MatmulWeightsDecompression : public testing::WithParamInterface, @@ -83,7 +83,7 @@ class MatmulWeightsDecompression : public testing::WithParamInterface