[GPU] Enable dynamic quantization for MXFP4 & regular FP4 dtypes

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<element::Type>& get_int8_support();
-    LP_TRANSFORMATIONS_API const std::vector<element::Type>& get_fp8_support();
+    LP_TRANSFORMATIONS_API const std::vector<element::Type>& get_low_bit_float_support();
     LP_TRANSFORMATIONS_API const std::vector<element::Type>& 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();
     }

src/common/low_precision_transformations/src/layer_transformation.cpp

@@ -29,11 +29,11 @@ const std::vector<element::Type>& precision_set::get_int8_support() {
     return int8_support;
 }
 
-const std::vector<element::Type>& precision_set::get_fp8_support() {
-    static const std::vector<element::Type> fp8_support = {
-        ov::element::f8e4m3,  ov::element::f8e5m2,
+const std::vector<element::Type>& precision_set::get_low_bit_float_support() {
+    static const std::vector<element::Type> low_bit_float_support = {
+        ov::element::f8e4m3,  ov::element::f8e5m2, ov::element::f4e2m1,
     };
-    return fp8_support;
+    return low_bit_float_support;
 }
 
 const std::vector<element::Type>& precision_set::get_int8_int16_int32_support() {

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)

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<ov::float8_e5m2>(actual_mem, stream, file_stream, dump_raw);
     else if (mem_dt == cldnn::data_types::f8e4m3)
         dump<ov::float8_e4m3>(actual_mem, stream, file_stream, dump_raw);
+         else if (mem_dt == cldnn::data_types::f4e2m1)
+        dump<ov::float4_e2m1>(actual_mem, stream, file_stream, dump_raw);
     else if (mem_dt == cldnn::data_types::f8e8m0)
         dump<ov::float8_e8m0>(actual_mem, stream, file_stream, dump_raw);
     else if (mem_dt == cldnn::data_types::boolean)

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,

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:

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

src/plugins/intel_gpu/src/graph/impls/onednn/fully_connected_onednn.cpp

@@ -71,7 +71,7 @@ struct fully_connected_onednn : typed_primitive_onednn_impl<fully_connected> {
             }
 
             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<fully_connected> {
         }
 
         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<fully_connected> {
 
         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 {

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);
 

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:

src/plugins/intel_gpu/src/graph/primitive_inst.cpp

@@ -544,7 +544,7 @@ void primitive_inst::update_shape() {
 
     if (get_node().is_type<dynamic_quantize>() && 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);

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
     }