diff --git a/src/plugins/intel_gpu/include/intel_gpu/runtime/debug_configuration.hpp b/src/plugins/intel_gpu/include/intel_gpu/runtime/debug_configuration.hpp index d86bdab7578..b3220136c6c 100644 --- a/src/plugins/intel_gpu/include/intel_gpu/runtime/debug_configuration.hpp +++ b/src/plugins/intel_gpu/include/intel_gpu/runtime/debug_configuration.hpp @@ -140,6 +140,7 @@ public: int disable_runtime_skip_reorder; // Disable runtime skip reorder int disable_primitive_fusing; // Disable primitive fusing int disable_fake_alignment; // Disable fake alignment + int enable_dynamic_quantize; // Enable Dynamic quantization for fully connected primitive std::set dump_iteration; // Dump n-th execution of network. std::vector load_layers_raw_dump; // List of layers to load dumped raw binary and filenames static const debug_configuration *get_instance(); 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 2d62e67c918..96bafc7e431 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 @@ -188,6 +188,8 @@ public: params.quantization = kernel_selector::QuantizationType::NONE; } + params.dynamic_quantization_group_size = impl_param.get_program().get_config().get_property(ov::hint::dynamic_quantization_group_size); + return params; } 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 9263421ccee..f22c1ee136c 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 @@ -17,6 +17,39 @@ // DISPATCH_FSV - output coordinates for each sub-group are calculated from linearized coordinates // DISPATCH_BSV as if they laid in bs_fs_bsv_fsv format, these macros describe fsv and bsv factors; +#define INPUT_LOAD_SIZE 4 + +#if FC_KERNEL_DYNAMIC_QUANTIZE +KERNEL(quantize_input)( + const __global INPUT0_TYPE* input, + __global char* quantized_input, + __global INPUT0_TYPE* de_quan_scale) { + const uint offset = get_global_id(0); + + uint input_offset = offset * QUANTIZE_GROUP_SIZE; + half4 input_0[8]; + char4 quantized_value[8]; + half max[8]; + + unroll_for (uint i = 0 ; i < 8 ; ++i) { + input_0[i] = vload4(0, &input[input_offset + i * 4]); + max[i] = fmax(fmax(fabs(input_0[i][0]), fabs(input_0[i][1])), fmax(fabs(input_0[i][2]), fabs(input_0[i][3]))); + } + + half max_value = fmax(fmax(fmax(max[0], max[1]), fmax(max[2], max[3])), + fmax(fmax(max[4], max[5]), fmax(max[6], max[7]))); + + half quan_scale = max_value / 128; + + unroll_for (uint i = 0 ; i < 8 ; ++i) { + quantized_value[i] = CAT(convert_, MAKE_VECTOR_TYPE(char, INPUT_LOAD_SIZE))(input_0[i] / (half4)quan_scale); + vstore4(quantized_value[i], 0, &quantized_input[input_offset + i * 4]); + } + + de_quan_scale[offset] = quan_scale; +} +#else // !FC_KERNEL_DYNAMIC_QUANTIZE + // Verify JIT parameters. #if SIMD != 8 && SIMD != 16 # error "fully_connected_gpu_bf_tiled.cl - SIMD must be one of {8, 16}" @@ -50,8 +83,10 @@ // Data stored in memory : f0k0k1|f16k0k1|f0k2k3|f16k2k3 // => unpack as f0k0k1|f0k2k3|f16k0k1|f16k2k3 so that the weight access order is preserved #define UNPACK_INT4 UNPACK_INT4x2_OSV32_ISV2 +#define UNPACK_TRANSPOSED_INT4 UNPACK_INT4x2_OSV32_ISV2 #else #define UNPACK_INT4 UNPACK_INT4x2 +#define UNPACK_TRANSPOSED_INT4 UNPACK_TRANSPOSED_INT4x2 #endif // Macros for vectorized types. #define INPUT_VEC_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, TILE_IFM) @@ -79,6 +114,7 @@ // Check alignment restrictions for using block writes on output. #define USE_BLOCK_WRITE ((OUTPUT_TYPE_SIZE * TILE_OUT_B_PITCH) % 16 == 0 && (OUTPUT_TYPE_SIZE * OUTPUT_OFFSET) % 16 == 0) + #if !REALIGN_FP16_OFFSET # if OUTPUT_3D # define MAIN_LOOP_ELEMENTS_COUNT INPUT0_SIZE_Y @@ -645,28 +681,6 @@ inline void FUNC(fc_bf_tiled_kernel_default)( #undef WRITE_OUTPUT } else { output_offset += sglid; - - // TODO: Investigate why below code doesn't compile and check how it affects performance. - //#define WRITE_OUTPUT_FEATURE(fi) do { \ - // const bool should_write = \ - // TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0 || \ - // out_f + (fi) * SIMD + sglid < TILE_OUT_F_NUM; \ - // if (should_write) { \ - // output[output_offset] = result[out_bi][fi]; \ - // } \ - // output_offset += SIMD; \ - // } while (false) - // - //#define WRITE_OUTPUT(bi) do { \ - // const uint out_bi = bi; \ - // CONST_LOOP(TILE_OFM, WRITE_OUTPUT_FEATURE); \ - // output_offset += TILE_OUT_B_PITCH - TILE_OFM * SIMD; \ - // } while (false) - // - //CONST_LOOP(TILE_B, WRITE_OUTPUT); - //#undef WRITE_OUTPUT - //#undef WRITE_OUTPUT_FEATURE - for (uint bi = 0; bi < TILE_B; ++bi) { for (uint fi = 0; fi < TILE_OFM; ++fi) { const bool should_write = @@ -686,6 +700,355 @@ inline void FUNC(fc_bf_tiled_kernel_default)( // ===================================================================================================================================== } +// Dyc Quantize +#if USE_SLM && DYNAMIC_QUANTIZE +#define PACKED_DQ_TYPE int +#define DQ_VEC_TYPE MAKE_VECTOR_TYPE(DQ_TYPE, TILE_IFM) +#define DQ_SLM_FILTER_VEC MAKE_VECTOR_TYPE(DQ_TYPE, 4) +#define DQ_SLM_FILTER_PACKED_VEC MAKE_VECTOR_TYPE(FILTER_TYPE, FILTER_LOAD_BLOCK_SIZE) +#define DQ_SLM_FILTER_UNPACKED_VEC MAKE_VECTOR_TYPE(DQ_TYPE, FILTER_ELEMENTS_PER_LOAD) +#define DQ_FILTER_VEC_TYPE MAKE_VECTOR_TYPE(DQ_TYPE, TILE_K_OFM) + +#define TO_DQ_TYPE(x) CAT(CAT(convert_, DQ_TYPE),_sat)(x) +#define TO_DQ_VEC_TYPE(x) CAT(convert_, DQ_VEC_TYPE)(x) +#define TO_DQ_SLM_FILTER_UNPACKED_VEC(x) CAT(convert_, DQ_SLM_FILTER_UNPACKED_VEC)(x) +#define TO_DQ_FILTER_VEC_TYPE(x) CAT(convert_, DQ_FILTER_VEC_TYPE)(x) + +#define AS_TYPE_N_(type, n, x) as_##type##n(x) +#define AS_TYPE_N(type, n, x) AS_TYPE_N_(type, n, x) +#define AS_DQ_TYPE_4(x) AS_TYPE_N(DQ_TYPE, INPUT_LOAD_SIZE, x) + +inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( + OPTIONAL_SHAPE_INFO_ARG + const __global INPUT0_TYPE* input, + __global char* quantized_input, + __global INPUT0_TYPE* scale, +#if DECOMPRESSION_SCALE_TERM + const __global DECOMPRESSION_SCALE_TYPE* decompression_scale, +#endif +#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR + const __global DECOMPRESSION_ZP_TYPE* decompression_zp, +#endif + __global OUTPUT_TYPE* output, + const __global FILTER_TYPE* weights + , __local int* wei_local_mem +#if BIAS_TERM + , const __global BIAS_TYPE* biases +#endif +#if HAS_FUSED_OPS_DECLS + , FUSED_OPS_DECLS +#endif +) { + uint gid = (uint)get_group_id(0); + uint local_id = (uint)get_local_id(2); + uint sglid = (uint)get_sub_group_local_id(); + + // Dispatch as bs_fs_bsv_fsv, where bsv = DISPATCH_BSV and fsv = DISPATCH_FSV. + // This allows more fine grained control over dispatch order than using work-groups and + // avoids requirement of threads being available for whole work-group. + // It could hovewer have some drawbacks like not providing physical locality or not using + // full dispatch pipeline. + uint feature_mini_block = gid % DISPATCH_FSV; + uint batch_mini_block = gid / DISPATCH_FSV % DISPATCH_BSV; + uint feature_mega_block = gid / (DISPATCH_FSV * DISPATCH_BSV) % (CEIL_DIV(TILE_OUT_F_NUM, TILE_OFM * SIMD) / DISPATCH_FSV); + uint batch_mega_block = gid / (DISPATCH_FSV * DISPATCH_BSV * CEIL_DIV(TILE_OUT_F_NUM, TILE_OFM * SIMD) / DISPATCH_FSV); + + FILTER_VEC_TYPE wei = 0; + + uint out_f = gid * (TILE_OFM * SIMD); + uint out_b = LWS_BATCHES * TILE_B * (uint)get_group_id(2) + local_id * TILE_B; + +#if OUTPUT_3D + uint out_b0 = out_b / OUTPUT_FEATURE_NUM; + uint out_b1 = out_b % OUTPUT_FEATURE_NUM; + uint input_offset = out_b0 * INPUT0_BATCH_PITCH + out_b1 * INPUT0_FEATURE_PITCH + INPUT0_OFFSET; +#else + uint input_offset = out_b * TILE_IN_B_PITCH + INPUT0_OFFSET; +#endif + + uint weights_offset = out_f * (INPUT_ELEMENTS_COUNT / 2); + + ACCUMULATOR_VEC_TYPE acc[TILE_B] = { }; + + // Dynamic Quantize + MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE) tiled_input_0[HALF_TILE_B] = { }; // Load 4 linear inputs for packing + PACKED_DQ_TYPE packed_in_0[HALF_TILE_B] = { }; // Packing char4 inputs to 1 integer + INPUT0_TYPE de_quantize_scale[TILE_B]; + +#if COMPRESSED_WEIGHTS && DECOMPRESSION_SCALE_GROUPS_NUM == 1 + #if DECOMPRESSION_SCALE_LENGTH > 1 && DECOMPRESSION_SCALE_LENGTH % (TILE_OFM * SIMD) == 0 + ACCUMULATOR_VEC_TYPE d_scale = TO_ACCUMULATOR_VEC_TYPE(BLOCK_READN(DECOMPRESSION_SCALE_TYPE, TILE_OFM, decompression_scale, out_f)); + #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(); + if (offset < DECOMPRESSION_SCALE_LENGTH) + ((ACCUMULATOR_TYPE*)(&d_scale))[of] = decompression_scale[offset]; + } + #else + ACCUMULATOR_VEC_TYPE d_scale = decompression_scale[0]; + #endif + + ACCUMULATOR_TYPE* d_scales = (ACCUMULATOR_TYPE*)(&d_scale); +#endif + +#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 = TO_ACCUMULATOR_VEC_TYPE(BLOCK_READN(DECOMPRESSION_ZP_TYPE, TILE_OFM, decompression_zp, out_f)); + #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(); + if (offset < DECOMPRESSION_ZP_LENGTH) + ((ACCUMULATOR_TYPE*)(&d_zp))[of] = decompression_zp[offset]; + } + #else + ACCUMULATOR_VEC_TYPE d_zp = decompression_zp[0]; + #endif + ACCUMULATOR_TYPE* d_zps = (ACCUMULATOR_TYPE*)(&d_zp); +#endif + + // ===================================================================================================================================== + // Main computation loop + const uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD); + // Each sub-group loads 2 Batch + uint idx_sglid = (sglid * TILE_K) % QUANTIZE_GROUP_SIZE; // same index for sglid 0~7 : to tile_k direction + uint batch_sglid = (sglid * TILE_K) / QUANTIZE_GROUP_SIZE; // 0 to 1 : to batch direction + + __attribute__((opencl_unroll_hint(1))) + for (uint ni = 0; ni < iterations; ++ni) { + uint in_offset = input_offset + (idx_sglid + batch_sglid * TILE_IN_B_PITCH); + uint scale_offset = input_offset / QUANTIZE_GROUP_SIZE; + for (uint bi = 0; bi < HALF_TILE_B; ++bi) { + // Load quantizing info from pre-quantizing kernel + tiled_input_0[bi] = vload4(0, &quantized_input[in_offset]); + de_quantize_scale[bi * 2] = scale[scale_offset]; + de_quantize_scale[bi * 2 + 1] = scale[scale_offset+ (TILE_IN_B_PITCH/QUANTIZE_GROUP_SIZE)]; + + // Packing : Get 4(B)x4(K) integer vector (packing to 4x1 vector) + packed_in_0[bi] = as_int(tiled_input_0[bi]); + + // Next batch + in_offset += (TILE_IN_B_PITCH * 2); + scale_offset += (TILE_IN_B_PITCH/QUANTIZE_GROUP_SIZE * 2); + } + + input_offset += TILE_IFM * SIMD; + + // Packing + MAKE_VECTOR_TYPE(int, TILE_B) acc_tmp[TILE_OFM] = { }; + + #if TILE_OFM != 2 + #error "FC bf_tiled kernel: can't use SLM optimization with TILE_OFM != 2" + #endif + + // Skip first barrier synchronization if there is only single outer loop iteration. + #if MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD) > 1 + barrier(CLK_LOCAL_MEM_FENCE); + #endif + + __local int* char_slm_weight = (__local int*)wei_local_mem; + + uint weights_idx = weights_offset + local_id * SIMD * FILTER_LOAD_ITERS * FILTER_LOAD_BLOCK_SIZE; + uint wei_local_idx = local_id * SIMD * FILTER_LOAD_ITERS * (FILTER_LOAD_BLOCK_SIZE/2) + sglid * 2; + + // DECOMPRESSION_SCALE_POST_OP SHOULD be enabled for dynamic quantize FC : scale is ACCUMULATOR_VAL_ONE + unroll_for(uint load_iter = 0; load_iter < FILTER_LOAD_ITERS; ++load_iter) { + SLM_FILTER_PACKED_VEC wei_packed = BLOCK_READN(FILTER_TYPE, FILTER_LOAD_BLOCK_SIZE, weights, weights_idx); + DQ_SLM_FILTER_UNPACKED_VEC dq_wei_unpacked = UNPACK_TRANSPOSED_INT4(DQ_TYPE, *((uint4x8_t *)&wei_packed)); + + // Calculate zero-point and scale only for DECOMPRESSION_SCALE_POST_OP enabled + #if DECOMPRESSION_ZP_TERM + #if DECOMPRESSION_ZP_SCALAR + DQ_SLM_FILTER_UNPACKED_VEC dzp = (DQ_SLM_FILTER_UNPACKED_VEC)(DECOMPRESSION_ZP_VALUE); + #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 + DQ_SLM_FILTER_UNPACKED_VEC dzp; + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + unroll_for(uint kii = 0; kii < FILTER_LOAD_BLOCK_SIZE; ++kii) { + const uint offset_ofm = out_f + fi*SIMD + sglid; + const uint offset_ifm = ni * TILE_IFM * SIMD + local_id * FILTER_LOAD_ITERS * FILTER_LOAD_BLOCK_SIZE + load_iter * FILTER_LOAD_BLOCK_SIZE + kii; + const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + + (offset_ifm / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; + dzp[W_IDX] = decompression_zp[zp_offset]; + } + } + #else + DQ_SLM_FILTER_UNPACKED_VEC dzp = (DQ_SLM_FILTER_UNPACKED_VEC)(d_zps[0]); + #endif + #else + DQ_SLM_FILTER_UNPACKED_VEC dzp = (DQ_SLM_FILTER_UNPACKED_VEC)(ACCUMULATOR_VAL_ZERO); + #endif + + // Calculate weight : w = (w - dzp) * ds + dq_wei_unpacked -= dzp; + + #if FILTER_LOAD_BLOCK_SIZE == 2 + DQ_SLM_FILTER_VEC wei_1 = {dq_wei_unpacked.s01, dq_wei_unpacked.s23}; + char_slm_weight[wei_local_idx] = as_int(wei_1); + #elif FILTER_LOAD_BLOCK_SIZE == 4 + DQ_SLM_FILTER_VEC wei_1 = {dq_wei_unpacked.s01, dq_wei_unpacked.s23}; + char_slm_weight[wei_local_idx] = as_int(wei_1); + DQ_SLM_FILTER_VEC wei_2 = {dq_wei_unpacked.s45, dq_wei_unpacked.s67}; + char_slm_weight[wei_local_idx+1] = as_int(wei_2); + #elif FILTER_LOAD_BLOCK_SIZE == 8 + DQ_SLM_FILTER_VEC wei_1 = {dq_wei_unpacked.s01, dq_wei_unpacked.s23}; + char_slm_weight[wei_local_idx] = as_int(wei_1); + DQ_SLM_FILTER_VEC wei_2 = {dq_wei_unpacked.s45, dq_wei_unpacked.s67}; + char_slm_weight[wei_local_idx+1] = as_int(wei_2); + DQ_SLM_FILTER_VEC wei_3 = {dq_wei_unpacked.s89, dq_wei_unpacked.sab}; + char_slm_weight[wei_local_idx+2] = as_int(wei_3); + DQ_SLM_FILTER_VEC wei_4 = {dq_wei_unpacked.scd, dq_wei_unpacked.sef}; + char_slm_weight[wei_local_idx+3] = as_int(wei_4); + #else + #error "FC bf_tiled kernel: unsupported FILTER_LOAD_BLOCK_SIZE for SLM kernel" + #endif + + wei_local_idx += SIMD * (FILTER_LOAD_BLOCK_SIZE/2); + weights_idx += SIMD * FILTER_LOAD_BLOCK_SIZE; + } + + wei_local_idx = sglid * 2; + + barrier(CLK_LOCAL_MEM_FENCE); + + unroll_for(uint ki = 0; ki < (TILE_IFM * SIMD) / TILE_K; ++ki) { + #if TILE_K != 4 + #error "FC bf_tiled kernel: unsupported TILE_K size for SLM kernel" + #endif + + // Compute input * weight : packed char4 type + char8 weight = vload8(0, (__local char *)(&char_slm_weight[wei_local_idx + 16*2*ki])); + char4 first_weight = weight.s0123; + char4 second_weight = weight.s4567; + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + char4 input_val = as_char4(_sub_group_shuffle(packed_in_0[bi / 2], (bi % 2) * 8 + ki)); + acc_tmp[0][bi] = imad_SW(acc_tmp[0][bi], input_val, first_weight); + acc_tmp[1][bi] = imad_SW(acc_tmp[1][bi], input_val, second_weight); + } + + weights_offset += TILE_K_OFM_PACKED * SIMD; + + #if DECOMPRESSION_SCALE_POST_OP && (TILE_IFM * SIMD > DECOMPRESSION_SCALE_GROUP_SIZE) + 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 + ki*TILE_K) / 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] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; + acc_tmp[fi][bi] = 0; + } + } + #endif + } // Whole tile_k elements of each iteration : ki + + #if DECOMPRESSION_SCALE_POST_OP && (TILE_IFM * SIMD <= DECOMPRESSION_SCALE_GROUP_SIZE) + const uint ni_offset = ((ni*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; + 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_offset; + ACCUMULATOR_TYPE ds = decompression_scale[scale_offset]; + #else + ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; + #endif + + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; + } + } + #endif + } // Main compute loop : ni + + // ===================================================================================================================================== + // Post-processing: bias, activation, fused-ops + ACTIVATION_VEC_TYPE activated[TILE_B] = { }; + for (uint bi = 0; bi < TILE_B; ++bi) { + activated[bi] = TO_ACTIVATION_VEC_TYPE(acc[bi]); + } + +#if BIAS_TERM + #if TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0 + BIAS_VEC_TYPE bias = BIAS_BLOCK_READ(biases, out_f); + #else + BIAS_VEC_TYPE bias = 0; + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + ((BIAS_TYPE*)(&bias))[fi] = biases[out_f + sglid + fi * SIMD]; + } + #endif + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + activated[bi] += TO_ACTIVATION_VEC_TYPE(bias); + } +#endif + + OUTPUT_VEC_TYPE result[TILE_B] = { }; +#if HAS_FUSED_OPS + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + #if TILE_OFM > 1 + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + FUSED_OPS_VEC; + result[bi][fi] = FUSED_OPS_RESULT_VEC; + } + #else + FUSED_OPS_SCALAR; + result[bi] = FUSED_OPS_RESULT_SCALAR; + #endif // TILE_OFM > 1 + } +#else + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + result[bi] = TO_OUTPUT_VEC_TYPE(ACTIVATION_TYPED(activated[bi], ACTIVATION_PARAMS_TYPED)); + } +#endif + + // ===================================================================================================================================== + // Write results + uint output_offset = out_f * TILE_OUT_F_PITCH + out_b * TILE_OUT_B_PITCH + OUTPUT_OFFSET; + + if (USE_BLOCK_WRITE && (TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0 || out_f + (TILE_OFM * SIMD) <= TILE_OUT_F_NUM)) { +#if IS_DYNAMIC + #define WRITE_OUTPUT(bi) do { \ + if (bi + out_b < BATCH_SIZE) \ + OUTPUT_BLOCK_WRITE(output, output_offset, result[bi]); \ + output_offset += TILE_OUT_B_PITCH; \ + } while (false) +#else + #define WRITE_OUTPUT(bi) do { \ + OUTPUT_BLOCK_WRITE(output, output_offset, result[bi]); \ + output_offset += TILE_OUT_B_PITCH; \ + } while (false) +#endif + CONST_LOOP(TILE_B, WRITE_OUTPUT); + #undef WRITE_OUTPUT + } else { + output_offset += sglid; + + for (uint bi = 0; bi < TILE_B; ++bi) { + for (uint fi = 0; fi < TILE_OFM; ++fi) { + const bool should_write = +#if IS_DYNAMIC + bi + out_b < BATCH_SIZE && +#endif + (TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0 || + out_f + fi * SIMD + sglid < TILE_OUT_F_NUM); + if (should_write) { + output[output_offset] = ((OUTPUT_TYPE*)(&result[bi]))[fi]; + } + output_offset += SIMD; + } + output_offset += TILE_OUT_B_PITCH - TILE_OFM * SIMD; + } + } + // ===================================================================================================================================== +} +#endif + REQD_SUB_GROUP_SIZE(SIMD) KERNEL(fc)( OPTIONAL_SHAPE_INFO_ARG @@ -704,9 +1067,17 @@ KERNEL(fc)( #if HAS_FUSED_OPS_DECLS , FUSED_OPS_DECLS #endif +#if DYNAMIC_QUANTIZE + , __global char* quantized_input + , __global INPUT0_TYPE* de_quan_scale +#endif ) { #if USE_SLM - __local ACCUMULATOR_TYPE wei_local_mem[TILE_IFM * SIMD * TILE_OFM * SIMD]; + #if DYNAMIC_QUANTIZE + __local int dq_wei_local_mem[SIMD * TILE_OFM * SIMD]; + #else + __local ACCUMULATOR_TYPE wei_local_mem[TILE_IFM * SIMD * TILE_OFM * SIMD]; + #endif #endif #if IS_DYNAMIC && COMPRESSED_WEIGHTS_INT4 const int batch_size = BATCH_SIZE; @@ -844,6 +1215,76 @@ KERNEL(fc)( #endif ); } else { + #if USE_SLM && DYNAMIC_QUANTIZE + FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( + OPTIONAL_SHAPE_INFO_TENSOR + input, + quantized_input, + de_quan_scale, + #if DECOMPRESSION_SCALE_TERM + decompression_scale, + #endif + #if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR + decompression_zp, + #endif + output, + weights + , dq_wei_local_mem + #if BIAS_TERM + , biases + #endif + #if HAS_FUSED_OPS_DECLS + , FUSED_OPS_ARGS + #endif + ); + #else + FUNC_CALL(fc_bf_tiled_kernel_default)( + OPTIONAL_SHAPE_INFO_TENSOR + input, + #if DECOMPRESSION_SCALE_TERM + decompression_scale, + #endif + #if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR + decompression_zp, + #endif + output, + weights + #if USE_SLM + , wei_local_mem + #endif + #if BIAS_TERM + , biases + #endif + #if HAS_FUSED_OPS_DECLS + , FUSED_OPS_ARGS + #endif + ); + #endif + } +#else + #if USE_SLM && DYNAMIC_QUANTIZE + FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( + OPTIONAL_SHAPE_INFO_TENSOR + input, + quantized_input, + de_quan_scale, + #if DECOMPRESSION_SCALE_TERM + decompression_scale, + #endif + #if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR + decompression_zp, + #endif + output, + weights + , dq_wei_local_mem + #if BIAS_TERM + , biases + #endif + #if HAS_FUSED_OPS_DECLS + , FUSED_OPS_ARGS + #endif + ); + #else FUNC_CALL(fc_bf_tiled_kernel_default)( OPTIONAL_SHAPE_INFO_TENSOR input, @@ -865,31 +1306,10 @@ KERNEL(fc)( , FUSED_OPS_ARGS #endif ); - } -#else - FUNC_CALL(fc_bf_tiled_kernel_default)( - OPTIONAL_SHAPE_INFO_TENSOR - input, - #if DECOMPRESSION_SCALE_TERM - decompression_scale, #endif - #if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR - decompression_zp, - #endif - output, - weights - #if USE_SLM - , wei_local_mem - #endif - #if BIAS_TERM - , biases - #endif - #if HAS_FUSED_OPS_DECLS - , FUSED_OPS_ARGS - #endif - ); #endif } +#endif // !FC_KERNEL_DYNAMIC_QUANTIZE #undef INPUT_VEC_TYPE #undef ACCUMULATOR_VEC_TYPE 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 index 00bebcb1162..0bdb3b3498a 100644 --- 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 @@ -20,6 +20,12 @@ inline uchar2 cvt_uint4x2_to_uint8x2(uint4x2_t v) __attribute__((overloadable)) return (uchar2)(v0, v1); } +inline char2 cvt_uint4x2_to_int8x2(uint4x2_t v) __attribute__((overloadable)) { + const char v0 = convert_char(v.s0 & 0x0F); + const char v1 = convert_char((v.s0 & 0xF0) >> 4); + return (char2)(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); @@ -28,6 +34,68 @@ inline char2 cvt_int4x2_to_int8x2(int4x2_t v) __attribute__((overloadable)) { return (char2)(v0, v1); } +inline uchar2 unpack_to_uchar(uint4x2_t v) __attribute__((overloadable)) { + return cvt_uint4x2_to_uint8x2(v); +} + +inline uchar8 unpack_to_uchar(uint4x8_t v) __attribute__((overloadable)) { + uchar2 v0 = unpack_to_uchar(v.s0); + uchar2 v1 = unpack_to_uchar(v.s1); + uchar2 v2 = unpack_to_uchar(v.s2); + uchar2 v3 = unpack_to_uchar(v.s3); + return (uchar8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1); +} + +inline char2 unpack_to_char(uint4x2_t v) __attribute__((overloadable)) { + return cvt_uint4x2_to_int8x2(v); +} + +inline char4 unpack_to_char(uint4x4_t v) __attribute__((overloadable)) { + char2 v0 = unpack_to_char(v.s0); + char2 v1 = unpack_to_char(v.s1); + return (char4)(v0.s0, v0.s1, v1.s0, v1.s1); +} + +inline char8 unpack_to_char(uint4x8_t v) __attribute__((overloadable)) { + char2 v0 = unpack_to_char(v.s0); + char2 v1 = unpack_to_char(v.s1); + char2 v2 = unpack_to_char(v.s2); + char2 v3 = unpack_to_char(v.s3); + return (char8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1); +} + +inline char8 unpack_transposed_to_char(uint4x8_t v) __attribute__((overloadable)) { + char2 v0 = unpack_to_char(v.s0); + char2 v1 = unpack_to_char(v.s1); + char2 v2 = unpack_to_char(v.s2); + char2 v3 = unpack_to_char(v.s3); + return (char8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); +} + +inline uchar8 unpack_transposed_to_uchar(uint4x8_t v) __attribute__((overloadable)) { + uchar2 v0 = unpack_to_uchar(v.s0); + uchar2 v1 = unpack_to_uchar(v.s1); + uchar2 v2 = unpack_to_uchar(v.s2); + uchar2 v3 = unpack_to_uchar(v.s3); + return (uchar8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); +} + +inline char8 unpack_transposed_to_char_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) { + char2 v0 = unpack_to_char(v.s0); + char2 v1 = unpack_to_char(v.s2); + char2 v2 = unpack_to_char(v.s1); + char2 v3 = unpack_to_char(v.s3); + return (char8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); +} + +inline uchar8 unpack_transposed_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) { + uchar2 v0 = unpack_to_uchar(v.s0); + uchar2 v1 = unpack_to_uchar(v.s2); + uchar2 v2 = unpack_to_uchar(v.s1); + uchar2 v3 = unpack_to_uchar(v.s3); + return (uchar8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); +} + inline float2 unpack_to_float(uint4x2_t v) __attribute__((overloadable)) { return convert_float2(cvt_uint4x2_to_uint8x2(v)); } @@ -116,7 +184,28 @@ inline half8 unpack_to_half_osv32_isv2(int4x8_t v) __attribute__((overloadable)) 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 char8 unpack_to_char_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) { + char2 v0 = unpack_to_char(v.s0); + char2 v1 = unpack_to_char(v.s2); + char2 v2 = unpack_to_char(v.s1); + char2 v3 = unpack_to_char(v.s3); + return (char8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1); +} + +inline uchar8 unpack_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) { + uchar2 v0 = unpack_to_uchar(v.s0); + uchar2 v1 = unpack_to_uchar(v.s2); + uchar2 v2 = unpack_to_uchar(v.s1); + uchar2 v3 = unpack_to_uchar(v.s3); + return (uchar8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1); +} + + #endif // defined(cl_khr_fp16) + #define UNPACK_INT4x2(target_type, value) CAT(unpack_to_, target_type)(value) #define UNPACK_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_to_, target_type), _osv32_isv2)(value) +#define UNPACK_TRANSPOSED_INT4x2(target_type, value) CAT(unpack_transposed_to_, target_type)(value) +#define UNPACK_TRANSPOSED_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_transposed_to_, target_type), _osv32_isv2)(value) 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 4e43398e391..c6b0acda06c 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 @@ -3,15 +3,79 @@ // #include "fully_connected_kernel_bf_tiled.h" - +#include "kernel_selector_utils.h" #include #include #include "common_types.h" static constexpr size_t simd = 16; +static constexpr size_t quantize_grp_size = 32; +static constexpr size_t min_slm_size = 256; namespace kernel_selector { +static std::pair get_input_bf_size(const fully_connected_params& params) { + size_t input_f = params.inputs[0].Feature().v; + size_t input_batch = params.inputs[0].Batch().v; + // 3D input + if (params.outputs[0].GetLayout() == DataLayout::bfyx) { + input_f = params.inputs[0].Y().v; + input_batch = params.inputs[0].Batch().v * params.inputs[0].Feature().v; + } + + return {input_batch, input_f}; +} + +static std::pair get_output_aligned_bf_size(const fully_connected_params& params, + bool needs_align, + uint32_t align_b = 1, + int32_t align_f = 1) { + size_t output_f = (needs_align == true) ? CeilDiv(params.outputs[0].Feature().v, align_f) : params.outputs[0].Feature().v; + size_t output_b = params.outputs[0].Batch().v; + // 3D output + if (params.outputs[0].GetLayout() == DataLayout::bfyx) { + output_f = (needs_align == true) ? CeilDiv(params.outputs[0].Y().v, align_f) : params.outputs[0].Y().v; + output_b = params.outputs[0].Batch().v * params.outputs[0].Feature().v; + } + + output_b = (needs_align == true) ? CeilDiv(output_b, align_b) : output_b; + + return {output_b, output_f}; +} + +// DYNAMIC_QUANTIZE +static bool should_dynamic_quantize(const fully_connected_params& params) { + auto dynamic_quantization_group_size = params.dynamic_quantization_group_size; + GPU_DEBUG_GET_INSTANCE(debug_config); + GPU_DEBUG_IF(debug_config->enable_dynamic_quantize) { + dynamic_quantization_group_size = quantize_grp_size; + } + + if (params.inputs[0].GetFirstElementOffset() != 0) + return false; + + if (dynamic_quantization_group_size < quantize_grp_size) + return false; + + auto threads = get_input_bf_size(params); + auto input_b = threads.first; + auto input_f = threads.second; + + const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; + if ((scale_group_size % simd == 0) && (input_f % quantize_grp_size == 0) && + (params.is_shape_agnostic || (params.inputs[0].Batch().v > 1 && input_b > min_slm_size)) && + params.inputs[0].GetDType() == Datatype::F16 && + (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && + (params.decompression_zero_point.Feature().v == 1)) { + GPU_DEBUG_TRACE_DETAIL << " Dynamic quantizing for FC : scale_group_size " << scale_group_size << ", Input (" << + kernel_selector::toString(params.inputs[0].GetDType()) << ", " << kernel_selector::toString(params.outputs[0].GetLayout()) << + ") B: " << params.inputs[0].Batch().v << ", F: " << params.inputs[0].Feature().v << ", Y: " << params.inputs[0].Y().v << std ::endl; + return true; + } + + return false; +} + FullyConnected_bf_tiled::FullyConnected_bf_tiled() : FullyConnectedKernelBase("fully_connected_gpu_bf_tiled") { for (unsigned tile_b = 1; tile_b <= 32; ++tile_b) for (unsigned tile_ofm = 1; tile_ofm <= 4; tile_ofm *= 2) @@ -154,12 +218,9 @@ struct TuneParamsSelector { bool TuneParamsSelector::VerifyTuneParams(const fully_connected_params& params, const tune_params& tparams) { // Check divisibility by dispatch tile sizes. - size_t output_f = params.outputs[0].Feature().v; - size_t output_b = params.outputs[0].Batch().v; - if (params.outputs[0].GetLayout() == DataLayout::bfyx) { - output_b *= params.outputs[0].Feature().v; - output_f = params.outputs[0].Y().v; - } + auto bf_size = get_output_aligned_bf_size(params, false); + size_t output_b = bf_size.first; + size_t output_f = bf_size.second; if (params.compressed && (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && @@ -182,7 +243,7 @@ bool TuneParamsSelector::VerifyTuneParams(const fully_connected_params& params, if (tparams.kernel_type == FullyConnected_bf_tiled::KernelType::SLM) { bool is_i4_u4 = (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4); const auto required_batch_alignment = 64; - if (!params.is_shape_agnostic && (!IsAligned(output_b, required_batch_alignment) || output_b < 256)) + if (!params.is_shape_agnostic && (!IsAligned(output_b, required_batch_alignment) || output_b < min_slm_size)) return false; const auto required_tile_b = 8; @@ -228,14 +289,10 @@ FullyConnected_bf_tiled::GetAutoTuneParams(const fully_connected_params& params, && TuneParamsSelector::VerifyTuneParams(params, auto_tune_params[idx])) return auto_tune_params[idx]; - size_t batch = params.outputs[0].Batch().v; - size_t output_f = params.outputs[0].Feature().v; + auto bf_size = get_output_aligned_bf_size(params, false); + size_t batch = bf_size.first; + size_t output_f = bf_size.second; - // 3d output - if (params.outputs[0].GetLayout() == DataLayout::bfyx) { - batch *= params.outputs[0].Feature().v; - output_f = params.outputs[0].Y().v; - } Datatype dtype = params.inputs[0].GetDType(); auto selector = TuneParamsSelector(params); @@ -259,7 +316,7 @@ FullyConnected_bf_tiled::GetAutoTuneParams(const fully_connected_params& params, } else { // Try to use SLM kernels if possible if (preferred_kernel_type != KernelType::DEFAULT) { - if (params.is_shape_agnostic) { + if (params.is_shape_agnostic && !should_dynamic_quantize(params)) { selector.Case(tune_params(16, 2, 2, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM)) .Case(tune_params(16, 2, 1, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM)); } @@ -344,14 +401,9 @@ FullyConnected_bf_tiled::SetDefault(const fully_connected_params& params, int au auto tparams = GetAutoTuneParams(params, kernel_type, autoTuneIndex); - size_t feature_threads = CeilDiv(params.outputs[0].Feature().v, tparams.tile_ofm * simd); - size_t batch_threads = params.outputs[0].Batch().v; - if (params.outputs[0].GetLayout() == DataLayout::bfyx) { - feature_threads = CeilDiv(params.outputs[0].Y().v, tparams.tile_ofm * simd); - batch_threads = params.outputs[0].Batch().v * params.outputs[0].Feature().v; - } - - batch_threads = CeilDiv(batch_threads, tparams.tile_b); + auto threads = get_output_aligned_bf_size(params, true, tparams.tile_b, tparams.tile_ofm * simd); + auto batch_threads = threads.first; + auto feature_threads = threads.second; const size_t lws_batches = 8; const size_t aligned_batch = Align(batch_threads, lws_batches); // Each WG calculates 8x8 batches (TILE_B x LWS[2] size) @@ -380,9 +432,7 @@ FullyConnected_bf_tiled::SetDefault(const fully_connected_params& params, int au KernelsPriority FullyConnected_bf_tiled::GetKernelsPriority(const Params& params) const { const auto& fc_params = static_cast(params); - size_t output_b = fc_params.outputs[0].Batch().v; - if (fc_params.outputs[0].GetLayout() == DataLayout::bfyx) - output_b *= fc_params.outputs[0].Feature().v; + size_t output_b = get_output_aligned_bf_size(fc_params, false).first; float estimated_time = FORCE_PRIORITY_9; if (output_b > 1 && fc_params.inputs[0].GetDType() == Datatype::F32) @@ -445,10 +495,23 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para jit.AddConstant(MakeJitConstant("FILTER_LOAD_BLOCK_SIZE", block_read_size)); jit.AddConstant(MakeJitConstant("FILTER_ELEMENTS_PER_LOAD", weights_elements_per_load)); jit.Merge(make_int4_packed_type_jit_constant("INT4_PACKED_TYPE_PRELOAD", params.weights.GetDType(), weights_elements_per_load)); + } else { + jit.AddConstant(MakeJitConstant("USE_SLM", 0)); + } + + // Validated perf gain, Dynamic quantize force enable SCALE_POST_OP for char type multiplication + if (should_dynamic_quantize(params) && dispatchData.tile_m > 1 && dispatchData.tile_n == 2) { + jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 1)); + jit.AddConstant(MakeJitConstant("DECOMPRESSION_SCALE_POST_OP", 1)); + jit.AddConstant(MakeJitConstant("DQ_TYPE", "char")); + jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", quantize_grp_size)); + } else { + jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 0)); } jit.AddConstant(MakeJitConstant("SIMD", simd)); jit.AddConstant(MakeJitConstant("TILE_B", dispatchData.tile_m)); + jit.AddConstant(MakeJitConstant("HALF_TILE_B", dispatchData.tile_m/2)); jit.AddConstant(MakeJitConstant("TILE_OFM", dispatchData.tile_n)); jit.AddConstant(MakeJitConstant("TILE_IFM", dispatchData.tile_mk)); jit.AddConstant(MakeJitConstant("TILE_K", dispatchData.tile_nk)); @@ -523,29 +586,53 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { kd.update_dispatch_data_func = [this](const Params& params, KernelData& kd) { const auto& prim_params = static_cast(params); - OPENVINO_ASSERT(kd.kernels.size() == 2, "[GPU] Invalid kernels size for update dispatch data func, expected 2, got ", kd.kernels.size()); + size_t output_batch = get_output_aligned_bf_size(prim_params, false).first; - size_t output_batch = prim_params.outputs[0].Batch().v; - if (prim_params.outputs[0].GetLayout() == DataLayout::bfyx) - output_batch *= prim_params.outputs[0].Feature().v; + // Get index of the added shape-agnostic kernel + int kernel_offset = 0; + if (kd.kernels.size() == 3) + kernel_offset = 1; // quantize kernel exists - // Choose one of the two shape agnostic kernels: - // - kd.kernels[0] for batches <= 240 (default version) - // - kd.kernels[1] for batches >= 256 (slm version) + // Choose one of the two shape agnostic kernels: N == added kernel number + // - kd.kernels[N-1] for batches <= 240 (default version) + // - kd.kernels[N] for batches >= 256 (slm version) const auto default_alignment = 16; - // We can use SLM version if `output_batch + default_alignment > 256` because memory and batch are aligned (whether 16 or 64 elements) - const auto skip_kernel_idx = output_batch + default_alignment > 256 ? 0 : 1; - const auto execute_kernel_idx = 1 - skip_kernel_idx; + // We can use SLM version if `output_batch + default_alignment > min_slm_size(256)` because memory and batch are aligned (whether 16 or 64 elements) + const auto execute_type = (output_batch + default_alignment > min_slm_size) ? KernelType::SLM : KernelType::DEFAULT; + const auto execute_kernel_idx = ((execute_type == KernelType::SLM) ? 1 : 0) + kernel_offset; + const auto skip_kernel_idx = ((execute_type == KernelType::SLM) ? 0 : 1) + kernel_offset; + + // Check default or SLM version FC, and disable remain version kd.kernels[skip_kernel_idx].skip_execution = true; - GPU_DEBUG_TRACE_DETAIL << "FC bf tiled: " << (execute_kernel_idx == 1 ? "SLM" : "Default") << " shape-agnostic kernel version " + GPU_DEBUG_TRACE_DETAIL << "FC bf tiled: " << (execute_type == KernelType::SLM ? "SLM" : "Default") << " shape-agnostic kernel version " << "will be used for batch size = " << output_batch << "\n"; - auto dispatchData = SetDefault(prim_params, -1, execute_kernel_idx); + auto dispatchData = SetDefault(prim_params, -1, static_cast(execute_type)); kd.kernels[execute_kernel_idx].params.workGroups.global = dispatchData.gws; kd.kernels[execute_kernel_idx].params.workGroups.local = dispatchData.lws; kd.kernels[execute_kernel_idx].skip_execution = KernelData::SkipKernelExecution(prim_params); + + if (!kd.internalBufferSizes.empty()) { + // Pre-quantizing kernel was generated. Update the kernel and intermediate buffers or disable it. + if (execute_type == KernelType::DEFAULT) { + kd.kernels[0].skip_execution = true; + } else { + kd.kernels[0].skip_execution = false; + size_t input_f = get_input_bf_size(prim_params).second; + size_t input_size = input_f * dispatchData.tile_m * dispatchData.gws[2]; + + if (kd.internalBufferSizes[0] < input_size) { + kd.internalBufferSizes.clear(); + kd.internalBufferSizes.push_back(input_size); // quantized input is char type + kd.internalBufferSizes.push_back(input_size / quantize_grp_size * 2); // de_quan_scale is half type + } + + kd.kernels[0].params.workGroups.global = {std::max((input_size / quantize_grp_size), (size_t)1), 1, 1}; + kd.kernels[0].params.workGroups.local = {16, 1, 1}; + } + } }; } } @@ -572,34 +659,51 @@ KernelsData FullyConnected_bf_tiled::GetTunedKernelsDataByIndex(const Params &pa weights_layout = WeightsLayout::os_iyx_osv64; } - auto kernels_data = GetCommonKernelsData(params, - fc_params.inputs[0].GetLayout(), - weights_layout, - tparams.exec_options, - autoTuneIndex); + KernelsData kernels_data; + if (should_dynamic_quantize(fc_params)) { + // Use seperate 2 kernels for dynamic quantizing : quantizing_kernel + fc_kernel + // 1st kernel : Dynamic quantizing by quantize_grp_size + // 2nd kernel : fully connected kernel with KernelType::DEFAULT. Quantized inputs and scale values could be used. + // 3rd kernel : (optional) fully connected shape_agnostic kernel with KernelType::SLM. Quantized inputs and scale values would be used. + kernels_data = GetMultiKernelsData(params, + fc_params.inputs[0].GetLayout(), + weights_layout, + tparams.exec_options, + autoTuneIndex); + OPENVINO_ASSERT(!kernels_data.empty() && !kernels_data[0].kernels.empty(), "[GPU] Error to create multi kernel for dynamic quantizing."); - // In case of dynamic params try to configure additional optimized SLM kernel for large batches - if (params.is_shape_agnostic) { - auto tparams = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex); - auto can_select_slm_kernel = tparams.kernel_type == KernelType::SLM; + if (params.is_shape_agnostic) + GetUpdateDispatchDataFunc(kernels_data[0]); + } else { + kernels_data = GetCommonKernelsData(params, + fc_params.inputs[0].GetLayout(), + weights_layout, + tparams.exec_options, + autoTuneIndex, + 0); - if (!can_select_slm_kernel) - return kernels_data; + if (params.is_shape_agnostic) { + auto tparams = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex); + auto can_select_slm_kernel = tparams.kernel_type == KernelType::SLM; - auto slm_kernel = GetCommonKernelsData(params, - fc_params.inputs[0].GetLayout(), - weights_layout, - tparams.exec_options, - autoTuneIndex, - 1); + if (!can_select_slm_kernel) + return kernels_data; - if (slm_kernel.empty() || slm_kernel[0].kernels.empty()) - return kernels_data; + auto slm_kernel = GetCommonKernelsData(params, + fc_params.inputs[0].GetLayout(), + weights_layout, + tparams.exec_options, + autoTuneIndex, + 1); - kernels_data[0].kernels.push_back(slm_kernel[0].kernels.back()); + if (slm_kernel.empty() || slm_kernel[0].kernels.empty()) + return kernels_data; - // Update default update_dispatch_data_func function - GetUpdateDispatchDataFunc(kernels_data[0]); + kernels_data[0].kernels.push_back(slm_kernel[0].kernels.back()); + + // Update default update_dispatch_data_func function + GetUpdateDispatchDataFunc(kernels_data[0]); + } } return kernels_data; @@ -630,4 +734,157 @@ KernelsData FullyConnected_bf_tiled::GetKernelsData(const Params& params) const return res; } + + +KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, + DataLayout dl, + WeightsLayout wl, + const std::string exeMode, + int autoTuneIndex) const { + if (!Validate(params)) { + return KernelsData(); + } + + const auto& fc_params = static_cast(params); + + bool bProperInput = fc_params.inputs[0].GetLayout() == dl; + if (!bProperInput && !fc_params.inputs[0].PitchesDifferFromLogicalDims()) { + bProperInput = (dl == DataLayout::fb && fc_params.inputs[0].GetLayout() == DataLayout::fyxb) || + (dl == DataLayout::bf && fc_params.inputs[0].GetLayout() == DataLayout::bfyx); + } + + KernelData kd = KernelData::Default(params, 2); + fully_connected_params& new_params = *static_cast(kd.params.get()); + + if (!bProperInput) { + new_params.inputs[0] = new_params.inputs[0].TransformIgnorePadding(dl); + kd.reorderInput = true; + } + + bool succeed = UpdateWeightsParams(new_params, + wl, + kd.weightsReorderParams, + GetSupportedKey()); + if (!succeed) { + return {}; + } + + int inputs_count = 1; + if (new_params.compressed) { + inputs_count++; + if (new_params.has_decompression_zp && !new_params.scalar_zp) + inputs_count++; + } + + // Generate dispatch data for KernelType::DEFAULT + int kernel_number = 0; + const DispatchData dispatchData = SetDefault(new_params, autoTuneIndex, kernel_number); + + // Dynamic-quantize kernel + { + auto& quan_kernel = kd.kernels[0]; + DispatchData dyn_quan_dispatch = dispatchData; + dyn_quan_dispatch.gws = {std::max((fc_params.inputs[0].PhysicalSize() / quantize_grp_size), (size_t)1), 1, 1}; + dyn_quan_dispatch.lws = {16, 1, 1}; + quan_kernel.params.workGroups.global = dyn_quan_dispatch.gws; + quan_kernel.params.workGroups.local = dyn_quan_dispatch.lws; + quan_kernel.skip_execution = false; + + auto quan_entry_point = GetEntryPoint(kernelName, fc_params.layerID, params, kernel_number); + auto quan_cldnn_jit = GetJitConstants(new_params, dyn_quan_dispatch); + quan_cldnn_jit.AddConstant(MakeJitConstant("FC_KERNEL_DYNAMIC_QUANTIZE", 1)); + auto quan_jit = CreateJit(kernelName, quan_cldnn_jit, quan_entry_point); + + + FillCLKernelData(quan_kernel, + dyn_quan_dispatch, + params.engineInfo, + kernelName, + quan_jit, + quan_entry_point, + exeMode, // No exec mode + false, + false, + 1, // Only INPUT_0 is used for quantizing + 0, // No fused ops + 0, // No output + fc_params.is_shape_agnostic); + + quan_kernel.params.arguments.clear(); // Clear original output argument + quan_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INPUT, 0}); + quan_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0}); + quan_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1}); + kd.internalBufferSizes.push_back(fc_params.inputs[0].PhysicalSize()); + kd.internalBufferSizes.push_back(fc_params.inputs[0].PhysicalSize() / quantize_grp_size * 2); + kernel_number++; + } + kd.internalBufferDataType = Datatype::F16; + + // FC kernel for dynamic quantized input with KernelType::DEFAULT + { + auto entry_point = GetEntryPoint(kernelName, fc_params.layerID, params, kernel_number); + auto cldnn_jit = GetJitConstants(new_params, dispatchData); + auto jit = CreateJit(kernelName, cldnn_jit, entry_point); + + auto& fc_kernel = kd.kernels[1]; + fc_kernel.params.workGroups.global = dispatchData.gws; + fc_kernel.params.workGroups.local = dispatchData.lws; + fc_kernel.skip_execution = false; + + FillCLKernelData(fc_kernel, + dispatchData, + params.engineInfo, + kernelName, + jit, + entry_point, + exeMode, + true, + !fc_params.bias.empty(), + inputs_count, + GetFusedPrimitiveInputsCount(params), + 1, + fc_params.is_shape_agnostic); + + fc_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0}); + fc_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1}); + kernel_number++; + } + + const DispatchData slm_Data = SetDefault(new_params, autoTuneIndex, kernel_number); + auto slm_params = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex); + auto can_select_slm_kernel = slm_params.kernel_type == KernelType::SLM; + // FC kernel for dynamic quantized input with KernelType::SLM + if (params.is_shape_agnostic && can_select_slm_kernel) { + kd.kernels.resize(kernel_number + 1); + + auto entry_point = GetEntryPoint(kernelName, fc_params.layerID, params, kernel_number); + auto cldnn_jit = GetJitConstants(new_params, slm_Data); + auto jit = CreateJit(kernelName, cldnn_jit, entry_point); + + auto& sa_kernel = kd.kernels[2]; + sa_kernel.params.workGroups.global = slm_Data.gws; + sa_kernel.params.workGroups.local = slm_Data.lws; + sa_kernel.skip_execution = false; + + FillCLKernelData(sa_kernel, + slm_Data, + params.engineInfo, + kernelName, + jit, + entry_point, + slm_params.exec_options, + true, + !fc_params.bias.empty(), + inputs_count, + GetFusedPrimitiveInputsCount(params), + 1, + fc_params.is_shape_agnostic); + + sa_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0}); + sa_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1}); + } + + kd.autoTuneIndex = autoTuneIndex; + return {kd}; +} } // namespace kernel_selector diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.h index 62ad0ba2d07..d40a985a85a 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.h @@ -31,6 +31,12 @@ public: ParamsKey GetSupportedKey() const override; DeviceFeaturesKey get_required_device_features_key(const Params& params) const override; + KernelsData GetMultiKernelsData(const Params ¶ms, + DataLayout dl, + WeightsLayout wl, + const std::string exeMode = EXE_MODE_DEFAULT, + int autoTuneIndex = -1) const; + struct tune_params { tune_params(unsigned tile_b, unsigned tile_ofm, diff --git a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_params.h b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_params.h index 845b3207a75..01428a5fbed 100644 --- a/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_params.h +++ b/src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_params.h @@ -15,6 +15,7 @@ struct fully_connected_params : public weight_bias_params { fully_connected_params() : weight_bias_params(KernelType::FULLY_CONNECTED) {} QuantizationType quantization = QuantizationType::NONE; + size_t dynamic_quantization_group_size = 0; ParamsKey GetParamsKey() const override { ParamsKey k = weight_bias_params::GetParamsKey(); diff --git a/src/plugins/intel_gpu/src/plugin/compiled_model.cpp b/src/plugins/intel_gpu/src/plugin/compiled_model.cpp index 2d49b01cd1b..b4d3e658e41 100644 --- a/src/plugins/intel_gpu/src/plugin/compiled_model.cpp +++ b/src/plugins/intel_gpu/src/plugin/compiled_model.cpp @@ -257,7 +257,8 @@ ov::Any CompiledModel::get_property(const std::string& name) const { ov::PropertyName{ov::hint::num_requests.name(), PropertyMutability::RO}, ov::PropertyName{ov::hint::inference_precision.name(), PropertyMutability::RO}, ov::PropertyName{ov::device::id.name(), PropertyMutability::RO}, - ov::PropertyName{ov::execution_devices.name(), PropertyMutability::RO} + ov::PropertyName{ov::execution_devices.name(), PropertyMutability::RO}, + ov::PropertyName{ov::hint::dynamic_quantization_group_size.name(), PropertyMutability::RO} }; } else if (name == ov::model_name) { return decltype(ov::model_name)::value_type {m_model_name}; diff --git a/src/plugins/intel_gpu/src/plugin/plugin.cpp b/src/plugins/intel_gpu/src/plugin/plugin.cpp index d49ea91b44d..b6db0233ee9 100644 --- a/src/plugins/intel_gpu/src/plugin/plugin.cpp +++ b/src/plugins/intel_gpu/src/plugin/plugin.cpp @@ -555,6 +555,7 @@ std::vector Plugin::get_supported_properties() const { ov::PropertyName{ov::hint::inference_precision.name(), PropertyMutability::RW}, ov::PropertyName{ov::hint::enable_cpu_pinning.name(), PropertyMutability::RW}, ov::PropertyName{ov::device::id.name(), PropertyMutability::RW}, + ov::PropertyName{ov::hint::dynamic_quantization_group_size.name(), PropertyMutability::RW} }; return supported_properties; diff --git a/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp b/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp index 3708a339b35..bac76dfec64 100644 --- a/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp +++ b/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp @@ -181,6 +181,7 @@ static void print_help_messages() { message_list.emplace_back("OV_GPU_DisableRuntimeSkipReorder", "Disable runtime skip reorder."); message_list.emplace_back("OV_GPU_DisablePrimitiveFusing", "Disable primitive fusing"); message_list.emplace_back("OV_GPU_DisableFakeAlignment", "Disable fake alignment"); + message_list.emplace_back("OV_GPU_EnableDynamicQuantize", "Enable Dynamic quantization for fully connected primitive"); message_list.emplace_back("OV_GPU_DumpIteration", "Dump n-th execution of network, separated by space."); message_list.emplace_back("OV_GPU_MemPreallocationOptions", "Controls buffer pre-allocation feature. Expects 4 values separated by space in " "the following order: number of iterations for pre-allocation(int), max size of single iteration in bytes(int), " @@ -245,7 +246,8 @@ debug_configuration::debug_configuration() , disable_build_time_weight_reorder_for_dynamic_nodes(0) , disable_runtime_skip_reorder(0) , disable_primitive_fusing(0) - , disable_fake_alignment(0) { + , disable_fake_alignment(0) + , enable_dynamic_quantize(0) { #ifdef GPU_DEBUG_CONFIG get_gpu_debug_env_var("Help", help); get_common_debug_env_var("Verbose", verbose); @@ -296,6 +298,7 @@ debug_configuration::debug_configuration() get_gpu_debug_env_var("DisableRuntimeSkipReorder", disable_runtime_skip_reorder); get_gpu_debug_env_var("DisablePrimitiveFusing", disable_primitive_fusing); get_gpu_debug_env_var("DisableFakeAlignment", disable_fake_alignment); + get_gpu_debug_env_var("EnableDynamicQuantize", enable_dynamic_quantize); std::string dump_iteration_str; get_gpu_debug_env_var("DumpIteration", dump_iteration_str); std::string mem_preallocation_params_str; diff --git a/src/plugins/intel_gpu/src/runtime/execution_config.cpp b/src/plugins/intel_gpu/src/runtime/execution_config.cpp index b0edfe39c90..b7bb9947717 100644 --- a/src/plugins/intel_gpu/src/runtime/execution_config.cpp +++ b/src/plugins/intel_gpu/src/runtime/execution_config.cpp @@ -56,6 +56,7 @@ void ExecutionConfig::set_default() { std::make_tuple(ov::internal::exclusive_async_requests, false), std::make_tuple(ov::internal::query_model_ratio, 1.0f), std::make_tuple(ov::cache_mode, ov::CacheMode::OPTIMIZE_SPEED), + std::make_tuple(ov::hint::dynamic_quantization_group_size, 0), // Legacy API properties std::make_tuple(ov::intel_gpu::nv12_two_inputs, false), diff --git a/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp b/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp index 34ef987a31f..cbe4eb25fef 100644 --- a/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp +++ b/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp @@ -1255,6 +1255,116 @@ public: } } + void test_compressed_int4_scale_dyn_quan(bool is_caching_test, bool is_dynamic, int batch = 1) { + tests::random_generator rg(GET_SUITE_NAME); + auto& engine = get_test_engine(); + + if (engine.get_device_info().dev_type == device_type::discrete_gpu) + GTEST_SKIP(); + + long int batch_num = batch; + long int ifm_num = 1024; + long int ofm_num = 4096; + long int scales_group_size = 32; + + bool is_3d = true; + + auto input_ps = is_3d ? ov::PartialShape{ batch_num, 1, ifm_num } : ov::PartialShape{ batch_num, ifm_num}; + auto dyn_input_ps = is_3d ? ov::PartialShape{ -1, 1, ifm_num } : ov::PartialShape{ -1, ifm_num}; + auto input_mem = engine.allocate_memory({ input_ps, data_types::f16, format::bfyx }); + + auto weights_mem = engine.allocate_memory({ {ofm_num, ifm_num}, data_types::u4, format::bfyx }); + auto scale_mem = engine.allocate_memory({ {ofm_num, ifm_num / scales_group_size}, data_types::f16, format::bfyx }); + + auto input_data = rg.generate_random_1d(batch_num * ifm_num, -16.0f, 16.0f); + set_values(input_mem, input_data); + + auto weigths_data = rg.generate_random_1d(ofm_num * ifm_num / 2, 0, 10); + set_values(weights_mem, weigths_data); + + auto scale_data = rg.generate_random_1d(ofm_num * ifm_num / scales_group_size, -4.0f, 4.0f); + set_values(scale_mem, scale_data); + + auto in_layout = is_dynamic ? layout{ dyn_input_ps, data_types::f16, format::bfyx } + : layout{ input_ps, data_types::f16, format::bfyx }; + + auto fc_prim = fully_connected("fc_prim", input_info("input"), "weights", "", "scale", "", data_types::f16, padding(), is_3d ? 3 : 2, 2); + fc_prim.decompression_zero_point_scalar = 0; + + // Implemented dynamic quantize kernel + auto get_ref_results = [&]() { + topology topology( + input_layout("input", in_layout), + data("weights", weights_mem), + data("scale", scale_mem), + fc_prim + ); + + auto config = get_test_default_config(engine); + config.set_property(ov::intel_gpu::allow_new_shape_infer(true)); + config.set_property(ov::intel_gpu::optimize_data(true)); + + network network(engine, topology, config); + network.set_input_data("input", input_mem); + + auto outputs = network.execute(); + OPENVINO_ASSERT(outputs.size() == 1); + OPENVINO_ASSERT(outputs.begin()->first == "fc_prim"); + + auto output_layout = outputs.begin()->second.get_layout(); + auto output_mem = outputs.begin()->second.get_memory(); + + return engine.reinterpret_buffer(*output_mem, output_layout); + }; + + topology topology( + input_layout("input", in_layout), + data("weights", weights_mem), + data("scale", scale_mem), + fc_prim + ); + + auto config = get_test_default_config(engine); + config.set_property(ov::intel_gpu::allow_new_shape_infer(true)); + config.set_property(ov::intel_gpu::optimize_data(true)); + config.set_property(ov::hint::dynamic_quantization_group_size(32)); + + network::ptr network = get_network(engine, topology, config, get_test_stream_ptr(), is_caching_test); + + if (is_dynamic && !engine.get_device_info().supports_immad) { + auto inst = network->get_primitive("fc_prim"); + auto impl = inst->get_impl(); + ASSERT_TRUE(impl != NULL); + ASSERT_EQ(impl->get_kernels().size(), size_t((is_dynamic ? 3 : 2))); // shape-agnostic kernels + } + + network->set_input_data("input", input_mem); + + auto outputs = network->execute(); + ASSERT_EQ(outputs.size(), size_t(1)); + ASSERT_EQ(outputs.begin()->first, "fc_prim"); + + auto output_mem = outputs.begin()->second.get_memory(); + cldnn::mem_lock output_ptr (output_mem, get_test_stream()); + + auto ref_output_mem = get_ref_results(); + cldnn::mem_lock output_ptr_ref (ref_output_mem, get_test_stream()); + + size_t count = 0; + float max_diff = 0.f; + float avg = 0.f; + for (size_t i = 0; i < output_ptr_ref.size(); ++i) { + auto abs_diff = std::abs(output_ptr_ref[i] - output_ptr[i]); + if (max_diff < abs_diff) + max_diff = abs_diff; + avg = abs_diff; + count++; + OPENVINO_ASSERT(abs_diff < 256); + } + GPU_DEBUG_LOG << "---> count: " << count << ", max_diff:" << max_diff << ", avg_diff: " << (avg/count) << std::endl; + } + + void test_compressed_int4_scale(bool is_caching_test, bool is_dynamic, long int batch_num, long int scales_group_size = 128) { tests::random_generator rg(GET_SUITE_NAME); auto& engine = get_test_engine(); @@ -3158,6 +3268,40 @@ TEST_F(fully_connected_gpu_tests, compressed_int4_scale_b1g128) { this->test_compressed_int4_scale(false, false, 1, 128); } +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_single_batch) { + this->test_compressed_int4_scale_dyn_quan(false, false); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan) { + this->test_compressed_int4_scale_dyn_quan(false, false, 512); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_unaligned) { + this->test_compressed_int4_scale_dyn_quan(false, false, 511); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_dynamic_single_batch) { + this->test_compressed_int4_scale_dyn_quan(false, true, 1); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_dynamic) { + this->test_compressed_int4_scale_dyn_quan(false, true, 512); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_dynamic_unaligned) { + this->test_compressed_int4_scale_dyn_quan(false, true, 511); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_cache) { + this->test_compressed_int4_scale_dyn_quan(true, false, 512); +} + +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_cache_dynamic) { + this->test_compressed_int4_scale_dyn_quan(true, true, 512); +} + + + TEST_F(fully_connected_gpu_tests, compressed_scale_bias) { this->test_compressed_scale_bias(false); }