From 70d57436fae0f20e42a8a5c74b822d5b8bae8d08 Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Tue, 7 May 2024 18:58:55 +0900 Subject: [PATCH 01/16] [GPU] Implement dynamic quantize FC + Added fc_bf_tiled_dyn_quan to bf_tiled + Applied dynamic quantize kernel for fp16 fc_gpu_bf_tiled + Defined DYNAMIC_QUANTIZE + Added test-case Signed-off-by: Min, Byung-il --- .../fully_connected_gpu_bf_tiled.cl | 657 +++++++++++++++++- .../include/batch_headers/int4_utils.cl | 71 ++ .../fully_connected_kernel_bf_tiled.cpp | 16 + .../test_cases/fully_connected_gpu_test.cpp | 112 +++ 4 files changed, 833 insertions(+), 23 deletions(-) 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..e3663c05cf2 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 @@ -50,8 +50,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_MIXED_INT4 UNPACK_MIXED_INT4x2_OSV32_ISV2 #else #define UNPACK_INT4 UNPACK_INT4x2 +#define UNPACK_MIXED_INT4 UNPACK_MIXED_INT4x2 #endif // Macros for vectorized types. #define INPUT_VEC_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, TILE_IFM) @@ -79,6 +81,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 @@ -686,6 +689,565 @@ inline void FUNC(fc_bf_tiled_kernel_default)( // ===================================================================================================================================== } + +// Dyc Quantize +#define INPUT_LOAD_SIZE 4 +#define DQ_TYPE char +#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, +#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 +#if USE_SLM + , __local int* wei_local_mem +#endif +#if BIAS_TERM + , const __global BIAS_TYPE* biases +#endif +#if HAS_FUSED_OPS_DECLS + , FUSED_OPS_DECLS +#endif +) { +#if USE_SLM + uint gid = (uint)get_group_id(0); + uint local_id = (uint)get_local_id(2); +#else + uint gid = (uint)get_group_id(0); +#endif + + 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); + +#if USE_SLM + uint out_f = gid * (TILE_OFM * SIMD); + uint out_b = LWS_BATCHES * TILE_B * (uint)get_group_id(2) + local_id * TILE_B; +#else + FILTER_VEC_TYPE wei = 0; + uint out_f = (feature_mega_block * DISPATCH_FSV + feature_mini_block) * (TILE_OFM * SIMD); + uint out_b = ((batch_mega_block * DISPATCH_BSV + batch_mini_block) * TILE_B); +#endif + +#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 + +#if COMPRESSED_WEIGHTS_INT4 + uint weights_offset = out_f * (INPUT_ELEMENTS_COUNT / 2); +#else + uint weights_offset = out_f * INPUT_ELEMENTS_COUNT; +#endif + + ACCUMULATOR_VEC_TYPE acc[TILE_B] = { }; + + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 + // Dyn Quan + MAKE_VECTOR_TYPE(INPUT0_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 max[HALF_TILE_B] = { }; + #else + INPUT_VEC_TYPE in_0[TILE_B] = { }; + #endif + +#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 + +#if REALIGN_FP16_OFFSET + // For fp16 we need to ensure that all block reads are aligned to 4 byte (2 words) boundary. + // To do this solve first input feature separately. + { + INPUT0_TYPE tmp_input = input[input_offset + get_sub_group_local_id() % TILE_B * TILE_IN_B_PITCH]; + ACCUMULATOR_VEC_TYPE tmp_wei = TO_ACCUMULATOR_VEC_TYPE(BLOCK_READN(FILTER_TYPE, TILE_OFM, weights, weights_offset)); + #if COMPRESSED_WEIGHTS + tmp_wei = (tmp_wei - d_zp) * d_scale; + #endif + unroll_for(uint bi = 0; bi < TILE_B; ++bi) { + acc[bi] = _sub_group_shuffle(tmp_input, bi) * tmp_wei; + } + + weights_offset += TILE_OFM * SIMD; + input_offset += 1; + } +#endif + // ===================================================================================================================================== + // Main computation loop + uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD); + uint idx_sglid = (sglid * TILE_K) % 32; // same index for sglid 0~7 : to tile_k direction + uint batch_sglid = (sglid * TILE_K) / 32; // 0 to 1 : to batch direction + + __attribute__((opencl_unroll_hint(1))) + for (uint ni = 0; ni < iterations; ++ni) { + // Load input. + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 + // Packing : Get 4(B)x4(K) integer vector (packing to 4x1 vector) + uint in_offset = input_offset + (idx_sglid + batch_sglid * TILE_IN_B_PITCH); + for (uint bi = 0; bi < HALF_TILE_B; ++bi) { + tiled_input_0[bi] = vload4(0, &input[in_offset]); + + // Next batch + in_offset += (TILE_IN_B_PITCH * 2); + } + + input_offset += TILE_IFM * SIMD; + #else + #define LOAD_IN_0(bi) do { \ + in_0[bi] = INPUT_BLOCK_READ(input, input_offset); \ + input_offset += TILE_IN_B_PITCH; \ + } while (false) + + CONST_LOOP(TILE_B, LOAD_IN_0); + #undef LOAD_IN_0 + input_offset += TILE_IFM * SIMD - TILE_IN_B_PITCH * TILE_B; + #endif + + #if USE_SLM && DYNAMIC_QUANTIZE + MAKE_VECTOR_TYPE(int, TILE_OFM) acc_tmp[TILE_B] = { }; + #else + ACCUMULATOR_VEC_TYPE acc_tmp[TILE_B] = { }; + #endif + + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 + #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 + + #if DYNAMIC_QUANTIZE + // Quantizing for loaded input using max value + MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) de_quantize_scale = 1; + MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) dq_max_input; + MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) quan = 128; + unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { + max[bi] = fmax(fmax(fabs(tiled_input_0[bi][0]), fabs(tiled_input_0[bi][1])), fmax(fabs(tiled_input_0[bi][2]), fabs(tiled_input_0[bi][3]))); + dq_max_input[bi] = sub_group_reduce_max(max[bi]); + } + de_quantize_scale = dq_max_input / quan; + // Packing 4 of converted inputs to integer type + unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { + packed_in_0[bi] = as_int(CAT(convert_, MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE))(tiled_input_0[bi] / de_quantize_scale[bi])); + } + #endif + + // __local SLM_FILTER_VEC* char_slm_weight = (__local SLM_FILTER_VEC*)wei_local_mem; + __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) { + // uchar4 wei_packed = as_uchar4(_sub_group_block_read_uc4((const __global uchar *)(weights) + (weights_idx))); + 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_MIXED_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 w_idx = fi * 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); + #endif // USE_SLM && COMPRESSED_WEIGHTS_INT4 + + unroll_for(uint ki = 0; ki < (TILE_IFM * SIMD) / TILE_K; ++ki) { + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 + #if (TILE_K != 1) && (TILE_K != 2) && (TILE_K != 4) + #error "FC bf_tiled kernel: unsupported TILE_K size for SLM kernel" + #endif + #elif COMPRESSED_WEIGHTS_INT4 + FILTER_PACKED_VEC_TYPE wei_packed = FILTER_BLOCK_READ(weights, weights_offset); + wei = UNPACK_INT4(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&wei_packed)); + #else + wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); + #endif + + #if COMPRESSED_WEIGHTS && !USE_SLM + ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); + unroll_for(uint kii = 0; kii < TILE_K; ++kii) { + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + const uint w_idx = kii * TILE_OFM + fi; + const uint offset_ofm = out_f + fi*SIMD + sglid; + // Valid only if DECOMPRESSION_SCALE_POST_OP is enabled + ACCUMULATOR_TYPE ds = ACCUMULATOR_VAL_ONE; + + #if DECOMPRESSION_ZP_TERM + #if DECOMPRESSION_ZP_SCALAR + ACCUMULATOR_TYPE dzp = DECOMPRESSION_ZP_VALUE; + #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 + const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + + ((kii + ki*TILE_K + ni*TILE_IFM*SIMD) / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; + ACCUMULATOR_TYPE dzp = decompression_zp[zp_offset]; + #else + ACCUMULATOR_TYPE dzp = d_zps[fi % DECOMPRESSION_ZP_LENGTH]; + #endif + #else + ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; + #endif + w[w_idx] = (w[w_idx] - dzp) * ds; + } + } + #endif + + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 + // Error if TILE_OFM != 2 + #if DYNAMIC_QUANTIZE + // Compute input * weight : packed char4 type + char4 input_val = AS_DQ_TYPE_4(_sub_group_shuffle(packed_in_0[0], ki)); + 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) { + acc_tmp[bi][0] = imad_SW(acc_tmp[bi][0], input_val, first_weight); + acc_tmp[bi][1] = imad_SW(acc_tmp[bi][1], input_val, second_weight); + input_val = as_char4(_sub_group_shuffle(packed_in_0[(bi+1) / 2], ((bi+1) % 2) * 8 + ki)); + } + #else + 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) { + half4 in_val = as_half4(_sub_group_shuffle(((int2*)(&tiled_input_0[bi/2]))[0], (bi % 2) * 8 + ki)); + unroll_for (uint kii = 0; kii < TILE_K; ++kii) { + ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[0] += in_val[kii] * convert_half(first_weight[kii]); + ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[1] += in_val[kii] * convert_half(second_weight[kii]); + } + } + #endif + #else + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + unroll_for (uint kii = 0; kii < TILE_K; ++kii) { + const uint total_k = ki * TILE_K + kii; + INPUT0_TYPE in_val = _sub_group_shuffle(((INPUT0_TYPE*)(&in_0[bi]))[total_k / SIMD], total_k % SIMD); + unroll_for (uint fi = 0; fi < TILE_OFM; ++fi) { + ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[kii * TILE_OFM + fi]; + } + } + } + + #endif + + 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 + + #if USE_SLM && DYNAMIC_QUANTIZE + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi / 2]; + acc_tmp[bi][fi] = 0; + #else + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; + acc_tmp[bi][fi] = 0; + #endif + } + } + #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 + + #if USE_SLM && DYNAMIC_QUANTIZE + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi / 2]; + #else + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; + #endif + } + } + #endif + } // Done main compute loop : ni + + // ===================================================================================================================================== + // Leftovers +#if MAIN_LOOP_ELEMENTS_COUNT % (TILE_IFM * SIMD) != 0 + // Handle leftovers in normal case without alignment correction. + #define LEFTOVER_IFM (MAIN_LOOP_ELEMENTS_COUNT % (TILE_IFM * SIMD)) + { + #define LOAD_IN_0(bi) do { \ + in_0[bi] = INPUT_BLOCK_READ(input, input_offset); \ + input_offset += TILE_IN_B_PITCH; \ + } while (false) + + CONST_LOOP(TILE_B, LOAD_IN_0); + #undef LOAD_IN_0 + input_offset += TILE_IFM * SIMD - TILE_IN_B_PITCH * TILE_B; + unroll_for(uint ki = 0; ki < CEIL_DIV(LEFTOVER_IFM, TILE_K); ++ki) { + #if USE_SLM + FILTER_VEC_TYPE wei = 0; + #endif + + #if COMPRESSED_WEIGHTS_INT4 + FILTER_PACKED_VEC_TYPE wei_packed = FILTER_BLOCK_READ(weights, weights_offset); + wei = UNPACK_INT4(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&wei_packed)); + #else + wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); + #endif + + #if COMPRESSED_WEIGHTS + ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); + unroll_for(uint kii = 0; kii < TILE_K; ++kii) { + unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { + const uint w_idx = kii * TILE_OFM + fi; + uint offset_ofm = out_f + fi*SIMD + get_sub_group_local_id(); + #if DECOMPRESSION_SCALE_GROUPS_NUM > 1 + const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH + + ((kii + ki*TILE_K + iterations*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; + ACCUMULATOR_TYPE ds = decompression_scale[scale_offset]; + #else + ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; + #endif + + #if DECOMPRESSION_ZP_TERM + #if DECOMPRESSION_ZP_SCALAR + ACCUMULATOR_TYPE dzp = DECOMPRESSION_ZP_VALUE; + #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 + const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + + ((kii + ki*TILE_K + iterations*TILE_IFM*SIMD) / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; + ACCUMULATOR_TYPE dzp = decompression_zp[zp_offset]; + #else + ACCUMULATOR_TYPE dzp = d_zps[fi % DECOMPRESSION_ZP_LENGTH]; + #endif + #else + ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; + #endif + w[w_idx] = (w[w_idx] - dzp) * ds; + } + } + #endif + weights_offset += TILE_K_OFM_PACKED * SIMD; + + unroll_for (uint kii = 0; kii < TILE_K; ++kii) { + unroll_for (uint fi = 0; fi < TILE_OFM; ++fi) { + unroll_for (uint bi = 0; bi < TILE_B; ++bi) { + const uint total_k = ki * TILE_K + kii; + if (total_k < LEFTOVER_IFM) { + INPUT0_TYPE in_val = _sub_group_shuffle(((INPUT0_TYPE*)(&in_0[bi]))[total_k / SIMD], total_k % SIMD); + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[kii * TILE_OFM + fi]; + } + } + } + } + } + } + #undef LEFTOVER_IFM +#endif // MAIN_LOOP_ELEMENTS_COUNT % (TILE_IFM * SIMD) != 0 + + // ===================================================================================================================================== + // 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; + } + } + // ===================================================================================================================================== +} + + REQD_SUB_GROUP_SIZE(SIMD) KERNEL(fc)( OPTIONAL_SHAPE_INFO_ARG @@ -706,7 +1268,8 @@ KERNEL(fc)( #endif ) { #if USE_SLM - __local ACCUMULATOR_TYPE wei_local_mem[TILE_IFM * SIMD * TILE_OFM * SIMD]; + __local int dq_wei_local_mem[SIMD * TILE_OFM * SIMD]; + __local ACCUMULATOR_TYPE wei_local_mem[TILE_IFM * SIMD * TILE_OFM * SIMD]; #endif #if IS_DYNAMIC && COMPRESSED_WEIGHTS_INT4 const int batch_size = BATCH_SIZE; @@ -843,6 +1406,76 @@ KERNEL(fc)( , FUSED_OPS_ARGS #endif ); + } else { + if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE/* && !FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2*/) { + FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( + 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 + , dq_wei_local_mem + #endif + #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 + ); + } + } +#else + if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE/* && !FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2*/) { + FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( + 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 + , dq_wei_local_mem + #endif + #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 @@ -866,28 +1499,6 @@ KERNEL(fc)( #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 } 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..e830ab0828b 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,69 @@ 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_mixed_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_mixed_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_mixed_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_mixed_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); +} + + inline float2 unpack_to_float(uint4x2_t v) __attribute__((overloadable)) { return convert_float2(cvt_uint4x2_to_uint8x2(v)); } @@ -120,3 +189,5 @@ inline half8 unpack_to_half_osv32_isv2(int4x8_t v) __attribute__((overloadable)) #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_MIXED_INT4x2(target_type, value) CAT(unpack_mixed_to_, target_type)(value) +#define UNPACK_MIXED_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_mixed_to_, target_type), _osv32_isv2)(value) \ No newline at end of file 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 250e38e694b..4ea046b512f 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 @@ -444,10 +444,26 @@ 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 + const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; + if ((scale_group_size % simd == 0) && + params.inputs[0].GetDType() == Datatype::F16 && params.outputs[0].GetDType() == Datatype::F16 && + (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && + params.inputs[0].Y().v > 16 && dispatchData.tile_m > 1 && dispatchData.tile_n == 2 && + params.decompression_zero_point.Feature().v == 1) { + jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 1)); + jit.AddConstant(MakeJitConstant("DECOMPRESSION_SCALE_POST_OP", 1)); + } 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)); 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 1be8616c656..a39c35b49db 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,113 @@ public: } } + void test_compressed_int4_scale_my(bool is_caching_test, bool is_dynamic) { + 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 = 512; + 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 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{ {-1, ifm_num}, 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; + + auto get_implemented_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::intel_gpu::force_implementations(ov::intel_gpu::ImplForcingMap{{"fc_prim", { format::bfyx, "fully_connected_gpu_bfyx_ref"}}})); + + network::ptr network = get_network(engine, topology, config, get_test_stream_ptr(), is_caching_test); + + if (is_dynamic) { + auto inst = network->get_primitive("fc_prim"); + auto impl = inst->get_impl(); + ASSERT_EQ(impl->get_kernels().size(), size_t(2)); // Two 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 impl_output_mem = get_implemented_results(); + cldnn::mem_lock output_ptr_impl (impl_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_impl.size(); ++i) { + auto abs_diff = std::abs(output_ptr_impl[i] - output_ptr[i]); + if (max_diff < abs_diff) + max_diff = abs_diff; + avg = abs_diff; + count++; + OPENVINO_ASSERT(abs_diff < 256); + } + std::cout << "---> 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(); @@ -3073,6 +3180,11 @@ TEST_F(fully_connected_gpu_tests, compressed_scale_zp_bias_cached) { this->test_compressed_scale_zp_bias(true); } +// Testing +TEST_F(fully_connected_gpu_tests, compressed_int4_scale_my) { + this->test_compressed_int4_scale_my(false, false); +} + TEST_F(fully_connected_gpu_tests, compressed_int4_scale) { this->test_compressed_int4_scale(false, false, 256); } From df9b1bdb31e572c48af37fc8922e9d10b9ee1abf Mon Sep 17 00:00:00 2001 From: "Min, Byungil" Date: Thu, 9 May 2024 03:20:26 +0900 Subject: [PATCH 02/16] Minor bugfix Signed-off-by: Min, Byungil --- .../fully_connected_gpu_bf_tiled.cl | 17 +++++++------ .../include/batch_headers/int4_utils.cl | 24 ++++++++++++++++--- 2 files changed, 29 insertions(+), 12 deletions(-) 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 e3663c05cf2..866eeb4f03b 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 @@ -50,7 +50,7 @@ // 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_MIXED_INT4 UNPACK_MIXED_INT4x2_OSV32_ISV2 +#define UNPACK_MIXED_INT4 UNPACK_INT4x2_OSV32_ISV2 #else #define UNPACK_INT4 UNPACK_INT4x2 #define UNPACK_MIXED_INT4 UNPACK_MIXED_INT4x2 @@ -737,6 +737,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( uint gid = (uint)get_group_id(0); #endif + uint bgid = (uint)get_group_id(2); uint sglid = (uint)get_sub_group_local_id(); // Dispatch as bs_fs_bsv_fsv, where bsv = DISPATCH_BSV and fsv = DISPATCH_FSV. @@ -751,7 +752,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #if USE_SLM uint out_f = gid * (TILE_OFM * SIMD); - uint out_b = LWS_BATCHES * TILE_B * (uint)get_group_id(2) + local_id * TILE_B; + uint out_b = LWS_BATCHES * TILE_B * bgid + local_id * TILE_B; #else FILTER_VEC_TYPE wei = 0; uint out_f = (feature_mega_block * DISPATCH_FSV + feature_mini_block) * (TILE_OFM * SIMD); @@ -889,6 +890,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( max[bi] = fmax(fmax(fabs(tiled_input_0[bi][0]), fabs(tiled_input_0[bi][1])), fmax(fabs(tiled_input_0[bi][2]), fabs(tiled_input_0[bi][3]))); dq_max_input[bi] = sub_group_reduce_max(max[bi]); } + de_quantize_scale = dq_max_input / quan; // Packing 4 of converted inputs to integer type unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { @@ -916,12 +918,11 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( 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 w_idx = fi * 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]; + dzp[W_IDX] = decompression_zp[zp_offset]; } } #else @@ -980,7 +981,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); unroll_for(uint kii = 0; kii < TILE_K; ++kii) { unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { - const uint w_idx = kii * TILE_OFM + fi; const uint offset_ofm = out_f + fi*SIMD + sglid; // Valid only if DECOMPRESSION_SCALE_POST_OP is enabled ACCUMULATOR_TYPE ds = ACCUMULATOR_VAL_ONE; @@ -998,7 +998,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #else ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; #endif - w[w_idx] = (w[w_idx] - dzp) * ds; + w[W_IDX] = (w[W_IDX] - dzp) * ds; } } #endif @@ -1034,7 +1034,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( const uint total_k = ki * TILE_K + kii; INPUT0_TYPE in_val = _sub_group_shuffle(((INPUT0_TYPE*)(&in_0[bi]))[total_k / SIMD], total_k % SIMD); unroll_for (uint fi = 0; fi < TILE_OFM; ++fi) { - ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[kii * TILE_OFM + fi]; + ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[W_IDX]; } } } @@ -1121,7 +1121,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); unroll_for(uint kii = 0; kii < TILE_K; ++kii) { unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { - const uint w_idx = kii * TILE_OFM + fi; uint offset_ofm = out_f + fi*SIMD + get_sub_group_local_id(); #if DECOMPRESSION_SCALE_GROUPS_NUM > 1 const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH + @@ -1144,7 +1143,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #else ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; #endif - w[w_idx] = (w[w_idx] - dzp) * ds; + w[W_IDX] = (w[W_IDX] - dzp) * ds; } } #endif diff --git a/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/int4_utils.cl b/src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/int4_utils.cl index e830ab0828b..9403258a16b 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 @@ -85,7 +85,7 @@ inline char8 unpack_mixed_to_char_osv32_isv2(uint4x8_t v) __attribute__((overloa 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); + return (char8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); } inline uchar8 unpack_mixed_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) { @@ -93,10 +93,9 @@ inline uchar8 unpack_mixed_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overl 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); + 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)); } @@ -185,8 +184,27 @@ 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_MIXED_INT4x2(target_type, value) CAT(unpack_mixed_to_, target_type)(value) From 3e8f67be3b010568dcc028d6490b74c1caf2c8fe Mon Sep 17 00:00:00 2001 From: "Min, Byungil" Date: Thu, 9 May 2024 11:46:57 +0900 Subject: [PATCH 03/16] Fix Accuracy issue + double reduce_max Signed-off-by: Min, Byungil --- .../fully_connected_gpu_bf_tiled.cl | 44 ++++++++++++------- .../fully_connected_kernel_bf_tiled.cpp | 3 +- 2 files changed, 30 insertions(+), 17 deletions(-) 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 866eeb4f03b..027a3fa5762 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 @@ -737,7 +737,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( uint gid = (uint)get_group_id(0); #endif - uint bgid = (uint)get_group_id(2); uint sglid = (uint)get_sub_group_local_id(); // Dispatch as bs_fs_bsv_fsv, where bsv = DISPATCH_BSV and fsv = DISPATCH_FSV. @@ -752,7 +751,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #if USE_SLM uint out_f = gid * (TILE_OFM * SIMD); - uint out_b = LWS_BATCHES * TILE_B * bgid + local_id * TILE_B; + uint out_b = LWS_BATCHES * TILE_B * (uint)get_group_id(2) + local_id * TILE_B; #else FILTER_VEC_TYPE wei = 0; uint out_f = (feature_mega_block * DISPATCH_FSV + feature_mini_block) * (TILE_OFM * SIMD); @@ -779,7 +778,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // Dyn Quan MAKE_VECTOR_TYPE(INPUT0_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 max[HALF_TILE_B] = { }; #else INPUT_VEC_TYPE in_0[TILE_B] = { }; #endif @@ -842,7 +840,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( __attribute__((opencl_unroll_hint(1))) for (uint ni = 0; ni < iterations; ++ni) { - // Load input. #if USE_SLM && COMPRESSED_WEIGHTS_INT4 // Packing : Get 4(B)x4(K) integer vector (packing to 4x1 vector) uint in_offset = input_offset + (idx_sglid + batch_sglid * TILE_IN_B_PITCH); @@ -883,18 +880,23 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #if DYNAMIC_QUANTIZE // Quantizing for loaded input using max value - MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) de_quantize_scale = 1; - MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) dq_max_input; + INPUT0_TYPE max[2][HALF_TILE_B] = { 0 }; + MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) de_quantize_scale[2] = { }; + MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) dq_max_input[2] = { }; MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) quan = 128; unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { - max[bi] = fmax(fmax(fabs(tiled_input_0[bi][0]), fabs(tiled_input_0[bi][1])), fmax(fabs(tiled_input_0[bi][2]), fabs(tiled_input_0[bi][3]))); - dq_max_input[bi] = sub_group_reduce_max(max[bi]); + max[batch_sglid][bi] = fmax(fmax(fabs(tiled_input_0[bi][0]), fabs(tiled_input_0[bi][1])), fmax(fabs(tiled_input_0[bi][2]), fabs(tiled_input_0[bi][3]))); } + unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { + dq_max_input[0][bi] = sub_group_reduce_max(max[0][bi]); + dq_max_input[1][bi] = sub_group_reduce_max(max[1][bi]); + } + de_quantize_scale[0] = dq_max_input[0] / quan; + de_quantize_scale[1] = dq_max_input[1] / quan; - de_quantize_scale = dq_max_input / quan; // Packing 4 of converted inputs to integer type unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { - packed_in_0[bi] = as_int(CAT(convert_, MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE))(tiled_input_0[bi] / de_quantize_scale[bi])); + packed_in_0[bi] = as_int(CAT(convert_, MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE))(tiled_input_0[bi] / de_quantize_scale[batch_sglid][bi])); } #endif @@ -1043,7 +1045,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( weights_offset += TILE_K_OFM_PACKED * SIMD; - #if DECOMPRESSION_SCALE_POST_OP && (TILE_IFM * SIMD > DECOMPRESSION_SCALE_GROUP_SIZE) + #if (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; @@ -1057,7 +1059,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #if USE_SLM && DYNAMIC_QUANTIZE - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi / 2]; + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi % 2][bi / 2]; acc_tmp[bi][fi] = 0; #else ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; @@ -1068,7 +1070,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif } // Whole tile_k elements of each iteration : ki - #if DECOMPRESSION_SCALE_POST_OP && (TILE_IFM * SIMD <= DECOMPRESSION_SCALE_GROUP_SIZE) + #if (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) { @@ -1082,7 +1084,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #if USE_SLM && DYNAMIC_QUANTIZE - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi / 2]; + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi % 2][bi / 2]; #else ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; #endif @@ -1155,7 +1157,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( const uint total_k = ki * TILE_K + kii; if (total_k < LEFTOVER_IFM) { INPUT0_TYPE in_val = _sub_group_shuffle(((INPUT0_TYPE*)(&in_0[bi]))[total_k / SIMD], total_k % SIMD); - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[kii * TILE_OFM + fi]; + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[W_IDX]; } } } @@ -1407,6 +1409,12 @@ KERNEL(fc)( ); } else { if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE/* && !FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2*/) { + // if (get_sub_group_local_id() == 0 && get_group_id(0) == 0 && get_group_id(2) == 0 && + // get_local_id(0) == 0 && get_local_id(2) == 0) { + // printf(">>>> DYNAMIC : ELEMENTS_COUNT(%d) batch(%d) TILE_B(%d) OSV32_ISV2(%d) TILE_OFM(%d) => group0(%d) local0(%d) / group2(%d) local2(%d)\n", + // (int)MAIN_LOOP_ELEMENTS_COUNT, (int)INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM, (int)TILE_B, (int)FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2, (int)TILE_OFM, + // (int)get_global_size(0), (int)get_local_size(0), (int)get_global_size(2), (int)get_local_size(2)); + // } FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, @@ -1454,6 +1462,12 @@ KERNEL(fc)( } #else if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE/* && !FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2*/) { + // if (get_sub_group_local_id() == 0 && get_group_id(0) == 0 && get_group_id(2) == 0 && + // get_local_id(0) == 0 && get_local_id(2) == 0) { + // printf(">>>> STATIC : MAIN_LOOP_ELEMENTS_COUNT(%d) batch(%d) => group0(%d) local0(%d) / group2(%d) local2(%d)\n", + // (int)MAIN_LOOP_ELEMENTS_COUNT, (int)INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM, + // (int)get_global_size(0), (int)get_local_size(0), (int)get_global_size(2), (int)get_local_size(2)); + // } FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, 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 4ea046b512f..8e3922f200c 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 @@ -453,10 +453,9 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para if ((scale_group_size % simd == 0) && params.inputs[0].GetDType() == Datatype::F16 && params.outputs[0].GetDType() == Datatype::F16 && (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && - params.inputs[0].Y().v > 16 && dispatchData.tile_m > 1 && dispatchData.tile_n == 2 && + params.inputs[0].Y().v > 16 && dispatchData.tile_n == 2 && params.decompression_zero_point.Feature().v == 1) { jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 1)); - jit.AddConstant(MakeJitConstant("DECOMPRESSION_SCALE_POST_OP", 1)); } else { jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 0)); } From 311ee2d87fb576a557df39cce69c853d94305285 Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Fri, 24 May 2024 20:37:26 +0900 Subject: [PATCH 04/16] Implement 2 kernel approach Signed-off-by: Min, Byung-il --- .../fully_connected_gpu_bf_tiled.cl | 172 ++++++---- .../fully_connected_kernel_bf_tiled.cpp | 300 +++++++++++++++--- .../fully_connected_kernel_bf_tiled.h | 7 + 3 files changed, 372 insertions(+), 107 deletions(-) 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 027a3fa5762..f8f0dc845b9 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,47 @@ // 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 gid = get_group_id(0); + const uint local_id = get_local_id(0); + uint offset = gid * 32*16 + local_id * 32; + 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[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[offset + i * 4]); + } + + de_quan_scale[gid * 16 + local_id] = quan_scale; + // if (gid % 8 == 0 && gid <= 32 && local_id < 4) { + // printf(" -- global_id(%d) local_id(%d) offset(%d): %.2f, %.2f, %.2f, %.2f => Quan : %.2f, %.2f, %.2f, %.2f\n", (int)gid, (int)local_id, (int)offset, + // (float)input_0[0][0], (float)input_0[0][1], (float)input_0[0][2], (float)input_0[0][3], + // (float)quantized_value[0][0], (float)quantized_value[0][1], (float)quantized_value[0][2], (float)quantized_value[0][3]); + // // printf(" -- max_value (%.2f), quan_scale (%.2f)\n", max_value, de_quan_scale[gid * 16 + local_id]); + // printf(" -- gid(%d) local_id(%d) scale_offset(%d) : quan_scale (%.2f,%.2f) max_value(%.2f)\n", + // (int)gid, (int)local_id, (int)gid * 16 + local_id, (float)quan_scale, (float)de_quan_scale[gid * 16 + local_id], (float)max_value); + // } +} +#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}" @@ -648,28 +689,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 = @@ -689,10 +708,7 @@ inline void FUNC(fc_bf_tiled_kernel_default)( // ===================================================================================================================================== } - // Dyc Quantize -#define INPUT_LOAD_SIZE 4 -#define DQ_TYPE char #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) @@ -712,6 +728,10 @@ inline void FUNC(fc_bf_tiled_kernel_default)( inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_ARG const __global INPUT0_TYPE* input, +#if DYNAMIC_QUANTIZE + __global char* quantized_input, + __global INPUT0_TYPE* scale, +#endif #if DECOMPRESSION_SCALE_TERM const __global DECOMPRESSION_SCALE_TYPE* decompression_scale, #endif @@ -774,10 +794,11 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( ACCUMULATOR_VEC_TYPE acc[TILE_B] = { }; - #if USE_SLM && COMPRESSED_WEIGHTS_INT4 + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 && DYNAMIC_QUANTIZE // Dyn Quan - MAKE_VECTOR_TYPE(INPUT0_TYPE, INPUT_LOAD_SIZE) tiled_input_0[HALF_TILE_B] = { }; // Load 4 linear inputs for packing + 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 + MAKE_VECTOR_TYPE(INPUT0_TYPE, TILE_B) de_quantize_scale; #else INPUT_VEC_TYPE in_0[TILE_B] = { }; #endif @@ -835,22 +856,33 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // ===================================================================================================================================== // Main computation loop uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD); + // Each sub-group loads 2 Batch uint idx_sglid = (sglid * TILE_K) % 32; // same index for sglid 0~7 : to tile_k direction uint batch_sglid = (sglid * TILE_K) / 32; // 0 to 1 : to batch direction __attribute__((opencl_unroll_hint(1))) for (uint ni = 0; ni < iterations; ++ni) { - #if USE_SLM && COMPRESSED_WEIGHTS_INT4 - // Packing : Get 4(B)x4(K) integer vector (packing to 4x1 vector) + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 & DYNAMIC_QUANTIZE uint in_offset = input_offset + (idx_sglid + batch_sglid * TILE_IN_B_PITCH); + uint scale_offset = input_offset / 32; for (uint bi = 0; bi < HALF_TILE_B; ++bi) { - tiled_input_0[bi] = vload4(0, &input[in_offset]); + // 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/32]; + + // 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/32 * 2); } input_offset += TILE_IFM * SIMD; + + // Packing + MAKE_VECTOR_TYPE(int, TILE_B) acc_tmp[TILE_OFM] = { }; #else #define LOAD_IN_0(bi) do { \ in_0[bi] = INPUT_BLOCK_READ(input, input_offset); \ @@ -860,11 +892,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( CONST_LOOP(TILE_B, LOAD_IN_0); #undef LOAD_IN_0 input_offset += TILE_IFM * SIMD - TILE_IN_B_PITCH * TILE_B; - #endif - #if USE_SLM && DYNAMIC_QUANTIZE - MAKE_VECTOR_TYPE(int, TILE_OFM) acc_tmp[TILE_B] = { }; - #else ACCUMULATOR_VEC_TYPE acc_tmp[TILE_B] = { }; #endif @@ -878,7 +906,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( barrier(CLK_LOCAL_MEM_FENCE); #endif - #if DYNAMIC_QUANTIZE + #if 0 && DYNAMIC_QUANTIZE // Quantizing for loaded input using max value INPUT0_TYPE max[2][HALF_TILE_B] = { 0 }; MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) de_quantize_scale[2] = { }; @@ -1005,31 +1033,30 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( } #endif - #if USE_SLM && COMPRESSED_WEIGHTS_INT4 + #if USE_SLM && COMPRESSED_WEIGHTS_INT4 && DYNAMIC_QUANTIZE // Error if TILE_OFM != 2 - #if DYNAMIC_QUANTIZE - // Compute input * weight : packed char4 type - char4 input_val = AS_DQ_TYPE_4(_sub_group_shuffle(packed_in_0[0], ki)); - 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) { - acc_tmp[bi][0] = imad_SW(acc_tmp[bi][0], input_val, first_weight); - acc_tmp[bi][1] = imad_SW(acc_tmp[bi][1], input_val, second_weight); - input_val = as_char4(_sub_group_shuffle(packed_in_0[(bi+1) / 2], ((bi+1) % 2) * 8 + ki)); - } - #else - 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) { - half4 in_val = as_half4(_sub_group_shuffle(((int2*)(&tiled_input_0[bi/2]))[0], (bi % 2) * 8 + ki)); - unroll_for (uint kii = 0; kii < TILE_K; ++kii) { - ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[0] += in_val[kii] * convert_half(first_weight[kii]); - ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[1] += in_val[kii] * convert_half(second_weight[kii]); - } - } - #endif + // Compute input * weight : packed char4 type + // char4 input_val = AS_DQ_TYPE_4(_sub_group_shuffle(packed_in_0[0], ki)); + 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); + // input_val = as_char4(_sub_group_shuffle(packed_in_0[(bi+1) / 2], ((bi+1) % 2) * 8 + ki)); + } + // !DYNAMIC_QUANTIZE + // 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) { + // half4 in_val = as_half4(_sub_group_shuffle(((int2*)(&tiled_input_0[bi/2]))[0], (bi % 2) * 8 + ki)); + // unroll_for (uint kii = 0; kii < TILE_K; ++kii) { + // ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[0] += in_val[kii] * convert_half(first_weight[kii]); + // ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[1] += in_val[kii] * convert_half(second_weight[kii]); + // } + // } #else unroll_for (uint bi = 0; bi < TILE_B; ++bi) { unroll_for (uint kii = 0; kii < TILE_K; ++kii) { @@ -1045,7 +1072,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( weights_offset += TILE_K_OFM_PACKED * SIMD; - #if (TILE_IFM * SIMD > DECOMPRESSION_SCALE_GROUP_SIZE) + #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; @@ -1059,8 +1086,8 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #if USE_SLM && DYNAMIC_QUANTIZE - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi % 2][bi / 2]; - acc_tmp[bi][fi] = 0; + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; + acc_tmp[fi][bi] = 0; #else ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; acc_tmp[bi][fi] = 0; @@ -1070,7 +1097,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif } // Whole tile_k elements of each iteration : ki - #if (TILE_IFM * SIMD <= DECOMPRESSION_SCALE_GROUP_SIZE) + #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) { @@ -1084,7 +1111,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #if USE_SLM && DYNAMIC_QUANTIZE - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[bi]))[fi]) * ds * de_quantize_scale[bi % 2][bi / 2]; + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; #else ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; #endif @@ -1267,6 +1294,10 @@ 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 int dq_wei_local_mem[SIMD * TILE_OFM * SIMD]; @@ -1408,7 +1439,7 @@ KERNEL(fc)( #endif ); } else { - if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE/* && !FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2*/) { + if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE) { // if (get_sub_group_local_id() == 0 && get_group_id(0) == 0 && get_group_id(2) == 0 && // get_local_id(0) == 0 && get_local_id(2) == 0) { // printf(">>>> DYNAMIC : ELEMENTS_COUNT(%d) batch(%d) TILE_B(%d) OSV32_ISV2(%d) TILE_OFM(%d) => group0(%d) local0(%d) / group2(%d) local2(%d)\n", @@ -1418,6 +1449,10 @@ KERNEL(fc)( FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, + #if DYNAMIC_QUANTIZE + quantized_input, + de_quan_scale, + #endif #if DECOMPRESSION_SCALE_TERM decompression_scale, #endif @@ -1461,7 +1496,7 @@ KERNEL(fc)( } } #else - if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE/* && !FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2*/) { + if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE) { // if (get_sub_group_local_id() == 0 && get_group_id(0) == 0 && get_group_id(2) == 0 && // get_local_id(0) == 0 && get_local_id(2) == 0) { // printf(">>>> STATIC : MAIN_LOOP_ELEMENTS_COUNT(%d) batch(%d) => group0(%d) local0(%d) / group2(%d) local2(%d)\n", @@ -1471,6 +1506,10 @@ KERNEL(fc)( FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, + #if DYNAMIC_QUANTIZE + quantized_input, + de_quan_scale, + #endif #if DECOMPRESSION_SCALE_TERM decompression_scale, #endif @@ -1514,6 +1553,7 @@ KERNEL(fc)( } #endif } +#endif // !FC_KERNEL_DYNAMIC_QUANTIZE #undef INPUT_VEC_TYPE #undef ACCUMULATOR_VEC_TYPE 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 8e3922f200c..7b32a209e22 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,7 +3,7 @@ // #include "fully_connected_kernel_bf_tiled.h" - +#include "kernel_selector_utils.h" #include #include #include "common_types.h" @@ -12,6 +12,20 @@ static constexpr size_t simd = 16; namespace kernel_selector { +static const size_t group_size = 32; + +// DYNAMIC_QUANTIZE +static bool is_dynamic_quantize(const fully_connected_params& params) { + const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; + if ((scale_group_size % simd == 0) && (params.is_shape_agnostic || params.inputs[0].Batch().v > 8) && + params.inputs[0].GetDType() == Datatype::F16 && params.outputs[0].GetDType() == Datatype::F16 && + (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && + params.inputs[0].Y().v > 16 && params.decompression_zero_point.Feature().v == 1) + 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) @@ -258,10 +272,10 @@ 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) { - 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)); - } + // if (params.is_shape_agnostic) { + // 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)); + // } selector.Case(tune_params(8, 2, 2, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM)) .Case(tune_params(8, 2, 1, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM)); } @@ -449,17 +463,15 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para } // Validated perf gain, Dynamic quantize force enable SCALE_POST_OP for char type multiplication - const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; - if ((scale_group_size % simd == 0) && - params.inputs[0].GetDType() == Datatype::F16 && params.outputs[0].GetDType() == Datatype::F16 && - (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && - params.inputs[0].Y().v > 16 && dispatchData.tile_n == 2 && - params.decompression_zero_point.Feature().v == 1) { + if (is_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)); } else { jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 0)); } + jit.AddConstant(MakeJitConstant("DQ_TYPE", "char")); + jit.AddConstant(MakeJitConstant("SIMD", simd)); jit.AddConstant(MakeJitConstant("TILE_B", dispatchData.tile_m)); jit.AddConstant(MakeJitConstant("HALF_TILE_B", dispatchData.tile_m/2)); @@ -537,29 +549,55 @@ 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 = prim_params.outputs[0].Batch().v; if (prim_params.outputs[0].GetLayout() == DataLayout::bfyx) output_batch *= prim_params.outputs[0].Feature().v; - // Choose one of the two shape agnostic kernels: - // - kd.kernels[0] for batches <= 240 (default version) - // - kd.kernels[1] for batches >= 256 (slm version) + // Get index of the added shape-agnostic kernel + int kernel_num = kd.kernels.size() - 1; + + // 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; + const auto execute_kernel_idx = (output_batch + default_alignment > 256) ? kernel_num : kernel_num - 1; + const auto skip_kernel_idx = (execute_kernel_idx == kernel_num) ? (kernel_num - 1) : kernel_num; + // 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_kernel_idx == kernel_num ? "SLM" : "Default") << " shape-agnostic kernel version " << "will be used for batch size = " << output_batch << "\n"; + // printf(" -- previous kernel idx(%d) in kernels.size(%lu) for sa : gws(%lu, %lu, %lu)\n", execute_kernel_idx, kd.kernels.size(), + // kd.kernels[execute_kernel_idx].params.workGroups.global[0], kd.kernels[execute_kernel_idx].params.workGroups.global[1], kd.kernels[execute_kernel_idx].params.workGroups.global[2]); + auto dispatchData = SetDefault(prim_params, -1, execute_kernel_idx); 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); + + // printf(" -- update kernel idx(%d) for sa : gws(%lu, %lu, %lu) lws(%lu, %lu, %lu) physical(%lu)\n", execute_kernel_idx, + // dispatchData.gws[0], dispatchData.gws[1], dispatchData.gws[2], dispatchData.lws[0], dispatchData.lws[1], dispatchData.lws[2], prim_params.inputs[0].PhysicalSize()); + // printf(" -- relevant values : Physical(%lu) Y(%lu) TILE_B(%u) gws(%lu) lws(%lu) \n", prim_params.inputs[0].PhysicalSize(), prim_params.inputs[0].Y().v, dispatchData.tile_m, dispatchData.gws[2], dispatchData.lws[2]); + + if (!kd.internalBufferSizes.empty()) { + size_t input_size = prim_params.inputs[0].Y().v * dispatchData.tile_m * dispatchData.gws[2]; + // printf(" -- Require update intermediate buffer\n"); + // printf(" -- previous kernel(%d) for sa : input size(%lu) gws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, input_size, + // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], (int)kd.internalBufferSizes.size()); + + kd.kernels[0].params.workGroups.global = {std::max((input_size / (16 * 2)), (size_t)1), 1, 1}; + kd.kernels[0].params.workGroups.local = {16, 1, 1}; + kd.internalBufferSizes.clear(); + kd.internalBufferSizes.push_back(prim_params.inputs[0].PhysicalSize() * 2); + kd.internalBufferSizes.push_back(prim_params.inputs[0].PhysicalSize() / 32 * 2); + + // printf(" -- update kernel(%d) for sa : gws(%lu, %lu, %lu) lws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, + // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], + // kd.kernels[0].params.workGroups.local[0], kd.kernels[0].params.workGroups.local[1], kd.kernels[0].params.workGroups.local[2], (int)kd.internalBufferSizes.size()); + } }; } } @@ -586,34 +624,57 @@ 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; + // printf(">> GetTunedKernelsDataByIndex \n"); + if (is_dynamic_quantize(fc_params)) { + // Use seperate 2 kernels for dynamic quantizing : quantizing_kernel + fc_kernel + // First kernel : Dynamic quantizing with 32 group size + // Second kernel : fully connected with char type size using quantized inputs and scale values + kernels_data = GetMultiKernelsData(params, + fc_params.inputs[0].GetLayout(), + weights_layout, + tparams.exec_options, + autoTuneIndex, + 0); + 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; + // GPU_DEBUG_TRACE_DETAI << " Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; + // std::cout << " >> Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; + 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; + // GPU_DEBUG_TRACE_DETAI << " No Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; + // std::cout << " >> No! Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; - auto slm_kernel = GetCommonKernelsData(params, - fc_params.inputs[0].GetLayout(), - weights_layout, - tparams.exec_options, - autoTuneIndex, - 1); + if (params.is_shape_agnostic) { + auto tparams = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex); + auto can_select_slm_kernel = tparams.kernel_type == KernelType::SLM; - if (slm_kernel.empty() || slm_kernel[0].kernels.empty()) - return kernels_data; + if (!can_select_slm_kernel) + return kernels_data; - kernels_data[0].kernels.push_back(slm_kernel[0].kernels.back()); + auto slm_kernel = GetCommonKernelsData(params, + fc_params.inputs[0].GetLayout(), + weights_layout, + tparams.exec_options, + autoTuneIndex, + 1); - // Update default update_dispatch_data_func function - GetUpdateDispatchDataFunc(kernels_data[0]); + if (slm_kernel.empty() || slm_kernel[0].kernels.empty()) + return kernels_data; + + 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; @@ -644,4 +705,161 @@ 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, + int kernel_number) 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++; + } + + // Quantize kernel + const DispatchData dispatchData = SetDefault(new_params, autoTuneIndex, kernel_number); + { + auto& quan_kernel = kd.kernels[0]; + DispatchData dyn_quan_dispatch = dispatchData; + dyn_quan_dispatch.gws = {std::max((fc_params.inputs[0].PhysicalSize() / (16 * 2)), (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; + + // printf(" >> GetMultiKernelsData : Input0.Batch(%lu) Input0.Feature(%lu) Input0.Y(%lu) Input0.pysical(%lu)\n", + // fc_params.inputs[0].Batch().v, fc_params.inputs[0].Feature().v, fc_params.inputs[0].Y().v, fc_params.inputs[0].PhysicalSize()); + // printf(" -- Quantizing kernel : gws(%ld, %ld, %ld), lws(%ld, %ld, %ld)\n", dyn_quan_dispatch.gws[0], dyn_quan_dispatch.gws[1], dyn_quan_dispatch.gws[2], + // dyn_quan_dispatch.lws[0], dyn_quan_dispatch.lws[1], dyn_quan_dispatch.lws[2]); + + 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() * 2); + kd.internalBufferSizes.push_back(fc_params.inputs[0].PhysicalSize() / 32 * 2); + kernel_number++; + } + kd.internalBufferDataType = Datatype::F16; + + // FC kernel for dynamic quantized input + { + 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; + if (params.is_shape_agnostic && can_select_slm_kernel) { + // std::cout << " -- is_shape_agnostic quantizing : previous kernels size " << kd.kernels.size() << "\n"; + 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, // input : intermediate buffers (quantized input + scale value) + GetFusedPrimitiveInputsCount(params), + 1, // Output + 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}); + } + + // TODO Pass estimated time only through DispatchData + 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..7440b99ae99 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,13 @@ 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, + int kernel_number = 0) const; + struct tune_params { tune_params(unsigned tile_b, unsigned tile_ofm, From 83cb824b673342d4d003b92281834893f33ab33d Mon Sep 17 00:00:00 2001 From: "Min, Byungil" Date: Sun, 26 May 2024 13:07:24 +0900 Subject: [PATCH 05/16] Bugfix in dynamic quantizing Signed-off-by: Min, Byungil --- .../fully_connected_gpu_bf_tiled.cl | 9 +++- .../fully_connected_kernel_bf_tiled.cpp | 49 +++++++++++-------- 2 files changed, 36 insertions(+), 22 deletions(-) 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 f8f0dc845b9..5bee8f32368 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 @@ -757,6 +757,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( uint gid = (uint)get_group_id(0); #endif + uint sglid = (uint)get_sub_group_local_id(); // Dispatch as bs_fs_bsv_fsv, where bsv = DISPATCH_BSV and fsv = DISPATCH_FSV. @@ -769,11 +770,13 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( 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; + #if USE_SLM uint out_f = gid * (TILE_OFM * SIMD); uint out_b = LWS_BATCHES * TILE_B * (uint)get_group_id(2) + local_id * TILE_B; #else - FILTER_VEC_TYPE wei = 0; + // FILTER_VEC_TYPE wei = 0; uint out_f = (feature_mega_block * DISPATCH_FSV + feature_mini_block) * (TILE_OFM * SIMD); uint out_b = ((batch_mega_block * DISPATCH_BSV + batch_mini_block) * TILE_B); #endif @@ -798,7 +801,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // Dyn Quan 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 - MAKE_VECTOR_TYPE(INPUT0_TYPE, TILE_B) de_quantize_scale; + INPUT0_TYPE de_quantize_scale[TILE_B]; #else INPUT_VEC_TYPE in_0[TILE_B] = { }; #endif @@ -1440,6 +1443,8 @@ KERNEL(fc)( ); } else { if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE) { + // fc_bf_tiled_kernel_dyn_quan kernel is for dynamic quantizing. DECOMPRESSION_SCALE_POST_OP is required. + // It shows better performance with weight SLM and 4bit weight. // if (get_sub_group_local_id() == 0 && get_group_id(0) == 0 && get_group_id(2) == 0 && // get_local_id(0) == 0 && get_local_id(2) == 0) { // printf(">>>> DYNAMIC : ELEMENTS_COUNT(%d) batch(%d) TILE_B(%d) OSV32_ISV2(%d) TILE_OFM(%d) => group0(%d) local0(%d) / group2(%d) local2(%d)\n", 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 7b32a209e22..37ad3591e4c 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 @@ -17,7 +17,7 @@ static const size_t group_size = 32; // DYNAMIC_QUANTIZE static bool is_dynamic_quantize(const fully_connected_params& params) { const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; - if ((scale_group_size % simd == 0) && (params.is_shape_agnostic || params.inputs[0].Batch().v > 8) && + if ((scale_group_size % simd == 0) && (params.is_shape_agnostic || params.inputs[0].Batch().v > 256) && params.inputs[0].GetDType() == Datatype::F16 && params.outputs[0].GetDType() == Datatype::F16 && (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && params.inputs[0].Y().v > 16 && params.decompression_zero_point.Feature().v == 1) @@ -272,10 +272,10 @@ 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) { - // 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)); - // } + if (params.is_shape_agnostic && !is_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)); + } selector.Case(tune_params(8, 2, 2, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM)) .Case(tune_params(8, 2, 1, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM)); } @@ -561,8 +561,9 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { // - 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 execute_kernel_idx = (output_batch + default_alignment > 256) ? kernel_num : kernel_num - 1; - const auto skip_kernel_idx = (execute_kernel_idx == kernel_num) ? (kernel_num - 1) : kernel_num; + const auto execute_type = (output_batch + default_alignment > 256) ? KernelType::SLM : KernelType::DEFAULT; + const auto execute_kernel_idx = (execute_type == KernelType::SLM) ? kernel_num : kernel_num - 1; + const auto skip_kernel_idx = (execute_type == KernelType::SLM) ? (kernel_num - 1) : kernel_num; // Check default or SLM version FC, and disable remain version kd.kernels[skip_kernel_idx].skip_execution = true; @@ -573,7 +574,7 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { // printf(" -- previous kernel idx(%d) in kernels.size(%lu) for sa : gws(%lu, %lu, %lu)\n", execute_kernel_idx, kd.kernels.size(), // kd.kernels[execute_kernel_idx].params.workGroups.global[0], kd.kernels[execute_kernel_idx].params.workGroups.global[1], kd.kernels[execute_kernel_idx].params.workGroups.global[2]); - auto dispatchData = SetDefault(prim_params, -1, execute_kernel_idx); + auto dispatchData = SetDefault(prim_params, -1, (int)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); @@ -583,20 +584,28 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { // printf(" -- relevant values : Physical(%lu) Y(%lu) TILE_B(%u) gws(%lu) lws(%lu) \n", prim_params.inputs[0].PhysicalSize(), prim_params.inputs[0].Y().v, dispatchData.tile_m, dispatchData.gws[2], dispatchData.lws[2]); if (!kd.internalBufferSizes.empty()) { - size_t input_size = prim_params.inputs[0].Y().v * dispatchData.tile_m * dispatchData.gws[2]; - // printf(" -- Require update intermediate buffer\n"); - // printf(" -- previous kernel(%d) for sa : input size(%lu) gws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, input_size, - // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], (int)kd.internalBufferSizes.size()); + // 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_size = prim_params.inputs[0].Y().v * dispatchData.tile_m * dispatchData.gws[2]; + // printf(" -- Require update intermediate buffer\n"); + // printf(" -- previous kernel(%d) for sa : input size(%lu) gws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, input_size, + // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], (int)kd.internalBufferSizes.size()); - kd.kernels[0].params.workGroups.global = {std::max((input_size / (16 * 2)), (size_t)1), 1, 1}; - kd.kernels[0].params.workGroups.local = {16, 1, 1}; - kd.internalBufferSizes.clear(); - kd.internalBufferSizes.push_back(prim_params.inputs[0].PhysicalSize() * 2); - kd.internalBufferSizes.push_back(prim_params.inputs[0].PhysicalSize() / 32 * 2); + if (kd.kernels[0].params.workGroups.global[0] < (input_size / (16 * 2))) { + kd.internalBufferSizes.clear(); + kd.internalBufferSizes.push_back(input_size * 2); + kd.internalBufferSizes.push_back(input_size / 32 * 2); + } - // printf(" -- update kernel(%d) for sa : gws(%lu, %lu, %lu) lws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, - // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], - // kd.kernels[0].params.workGroups.local[0], kd.kernels[0].params.workGroups.local[1], kd.kernels[0].params.workGroups.local[2], (int)kd.internalBufferSizes.size()); + kd.kernels[0].params.workGroups.global = {std::max((input_size / (16 * 2)), (size_t)1), 1, 1}; + kd.kernels[0].params.workGroups.local = {16, 1, 1}; + // printf(" -- update kernel(%d) for sa : gws(%lu, %lu, %lu) lws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, + // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], + // kd.kernels[0].params.workGroups.local[0], kd.kernels[0].params.workGroups.local[1], kd.kernels[0].params.workGroups.local[2], (int)kd.internalBufferSizes.size()); + } } }; } From 840b1e8defa533cc7505ba7262bfb492fadca6d8 Mon Sep 17 00:00:00 2001 From: "Min, Byungil" Date: Tue, 28 May 2024 10:44:18 +0900 Subject: [PATCH 06/16] Use dyn quan for float output Signed-off-by: Min, Byungil --- .../kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 37ad3591e4c..72735bf66a0 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 @@ -18,7 +18,7 @@ static const size_t group_size = 32; static bool is_dynamic_quantize(const fully_connected_params& params) { const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; if ((scale_group_size % simd == 0) && (params.is_shape_agnostic || params.inputs[0].Batch().v > 256) && - params.inputs[0].GetDType() == Datatype::F16 && params.outputs[0].GetDType() == Datatype::F16 && + params.inputs[0].GetDType() == Datatype::F16 && (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && params.inputs[0].Y().v > 16 && params.decompression_zero_point.Feature().v == 1) return true; From 916358be780b84ddb3def2e4652b52681e9435f6 Mon Sep 17 00:00:00 2001 From: "Min, Byungil" Date: Tue, 28 May 2024 21:09:40 +0900 Subject: [PATCH 07/16] Generalize dynamic quantize scope Signed-off-by: Min, Byungil --- .../fully_connected_kernel_bf_tiled.cpp | 65 ++++++++++++++----- 1 file changed, 48 insertions(+), 17 deletions(-) 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 72735bf66a0..11db0cbe049 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 @@ -9,20 +9,53 @@ #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 const size_t group_size = 32; +static std::pair get_input_threads(const fully_connected_params& params) { + size_t ifm_threads = params.inputs[0].Feature().v; + size_t batch_threads = params.inputs[0].Batch().v; + if (params.outputs[0].GetLayout() == DataLayout::bfyx) { + ifm_threads = params.inputs[0].Y().v; + batch_threads = params.inputs[0].Batch().v * params.inputs[0].Feature().v; + } + + return {batch_threads, ifm_threads}; +} + +static std::pair get_output_threads(const fully_connected_params& params, uint32_t tile_ofm) { + size_t feature_threads = CeilDiv(params.outputs[0].Feature().v, 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, tile_ofm * simd); + batch_threads = params.outputs[0].Batch().v * params.outputs[0].Feature().v; + } + + return {batch_threads, feature_threads}; +} // DYNAMIC_QUANTIZE static bool is_dynamic_quantize(const fully_connected_params& params) { + auto threads = get_input_threads(params); + auto batch_threads = threads.first; + auto ifm_threads = threads.second; + const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v; - if ((scale_group_size % simd == 0) && (params.is_shape_agnostic || params.inputs[0].Batch().v > 256) && + if ((scale_group_size % simd == 0) && (ifm_threads % quantize_grp_size == 0) && + (params.is_shape_agnostic || (params.inputs[0].Batch().v > 1 && batch_threads > min_slm_size)) && params.inputs[0].GetDType() == Datatype::F16 && (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) && - params.inputs[0].Y().v > 16 && params.decompression_zero_point.Feature().v == 1) + (params.decompression_zero_point.Feature().v == 1)) return true; + GPU_DEBUG_TRACE_DETAIL << " Dynamic quantizing for FC : scale_group_size " << scale_group_size << ", Shape_Agnostic(" << (int)params.is_shape_agnostic << + ") : Input (" << kernel_selector::toString(params.inputs[0].GetDType()) << + ") B: " << params.inputs[0].Batch().v << ", F: " << params.inputs[0].Feature().v << ", Y: " << params.inputs[0].Y().v << + ", Weight " << kernel_selector::toString(params.weights.GetDType()) << ", zp_f " << params.decompression_zero_point.Feature().v << + ", format : " << kernel_selector::toString(params.outputs[0].GetLayout()) << std ::endl; + return false; } @@ -196,7 +229,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; @@ -357,12 +390,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; - } + auto threads = get_output_threads(params, tparams.tile_ofm); + auto batch_threads = threads.first; + auto feature_threads = threads.second; batch_threads = CeilDiv(batch_threads, tparams.tile_b); @@ -560,8 +590,8 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { // - 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 execute_type = (output_batch + default_alignment > 256) ? KernelType::SLM : KernelType::DEFAULT; + // 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) ? kernel_num : kernel_num - 1; const auto skip_kernel_idx = (execute_type == KernelType::SLM) ? (kernel_num - 1) : kernel_num; @@ -589,18 +619,19 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { kd.kernels[0].skip_execution = true; } else { kd.kernels[0].skip_execution = false; - size_t input_size = prim_params.inputs[0].Y().v * dispatchData.tile_m * dispatchData.gws[2]; + size_t ifm_threads = get_input_threads(prim_params).second; + size_t input_size = ifm_threads * dispatchData.tile_m * dispatchData.gws[2]; // printf(" -- Require update intermediate buffer\n"); // printf(" -- previous kernel(%d) for sa : input size(%lu) gws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, input_size, // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], (int)kd.internalBufferSizes.size()); - if (kd.kernels[0].params.workGroups.global[0] < (input_size / (16 * 2))) { + if (kd.kernels[0].params.workGroups.global[0] < (input_size / quantize_grp_size)) { kd.internalBufferSizes.clear(); kd.internalBufferSizes.push_back(input_size * 2); - kd.internalBufferSizes.push_back(input_size / 32 * 2); + kd.internalBufferSizes.push_back(input_size / quantize_grp_size * 2); } - kd.kernels[0].params.workGroups.global = {std::max((input_size / (16 * 2)), (size_t)1), 1, 1}; + 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}; // printf(" -- update kernel(%d) for sa : gws(%lu, %lu, %lu) lws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], @@ -762,7 +793,7 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, { auto& quan_kernel = kd.kernels[0]; DispatchData dyn_quan_dispatch = dispatchData; - dyn_quan_dispatch.gws = {std::max((fc_params.inputs[0].PhysicalSize() / (16 * 2)), (size_t)1), 1, 1}; + 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; From 98f0904c82020fb801f672912b64dc0b5e5e495a Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Sun, 2 Jun 2024 15:01:58 +0900 Subject: [PATCH 08/16] [GPU] Productize code + Applied code review + Remove debugging code + Fixed quantizing kernel Signed-off-by: Min, Byung-il --- .../fully_connected_gpu_bf_tiled.cl | 78 ++--------- .../fully_connected_kernel_bf_tiled.cpp | 123 +++++++----------- .../fully_connected_kernel_bf_tiled.h | 3 +- 3 files changed, 58 insertions(+), 146 deletions(-) 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 5bee8f32368..37ef63d5a22 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 @@ -24,15 +24,17 @@ KERNEL(quantize_input)( const __global INPUT0_TYPE* input, __global char* quantized_input, __global INPUT0_TYPE* de_quan_scale) { - const uint gid = get_group_id(0); - const uint local_id = get_local_id(0); - uint offset = gid * 32*16 + local_id * 32; + const uint offset = get_global_id(0); + if (offset != (get_group_id(0) * 16 + get_local_id(0))) + printf("!!!!!!!!!!!!!!!!!!!!!!\n") + + 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[offset + i * 4]); + 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]))); } @@ -43,18 +45,10 @@ KERNEL(quantize_input)( 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[offset + i * 4]); + vstore4(quantized_value[i], 0, &quantized_input[input_offset + i * 4]); } - de_quan_scale[gid * 16 + local_id] = quan_scale; - // if (gid % 8 == 0 && gid <= 32 && local_id < 4) { - // printf(" -- global_id(%d) local_id(%d) offset(%d): %.2f, %.2f, %.2f, %.2f => Quan : %.2f, %.2f, %.2f, %.2f\n", (int)gid, (int)local_id, (int)offset, - // (float)input_0[0][0], (float)input_0[0][1], (float)input_0[0][2], (float)input_0[0][3], - // (float)quantized_value[0][0], (float)quantized_value[0][1], (float)quantized_value[0][2], (float)quantized_value[0][3]); - // // printf(" -- max_value (%.2f), quan_scale (%.2f)\n", max_value, de_quan_scale[gid * 16 + local_id]); - // printf(" -- gid(%d) local_id(%d) scale_offset(%d) : quan_scale (%.2f,%.2f) max_value(%.2f)\n", - // (int)gid, (int)local_id, (int)gid * 16 + local_id, (float)quan_scale, (float)de_quan_scale[gid * 16 + local_id], (float)max_value); - // } + de_quan_scale[offset] = quan_scale; } #else // !FC_KERNEL_DYNAMIC_QUANTIZE @@ -798,7 +792,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( ACCUMULATOR_VEC_TYPE acc[TILE_B] = { }; #if USE_SLM && COMPRESSED_WEIGHTS_INT4 && DYNAMIC_QUANTIZE - // Dyn Quan + // 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]; @@ -856,6 +850,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( input_offset += 1; } #endif + // ===================================================================================================================================== // Main computation loop uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD); @@ -909,28 +904,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( barrier(CLK_LOCAL_MEM_FENCE); #endif - #if 0 && DYNAMIC_QUANTIZE - // Quantizing for loaded input using max value - INPUT0_TYPE max[2][HALF_TILE_B] = { 0 }; - MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) de_quantize_scale[2] = { }; - MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) dq_max_input[2] = { }; - MAKE_VECTOR_TYPE(INPUT0_TYPE, HALF_TILE_B) quan = 128; - unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { - max[batch_sglid][bi] = fmax(fmax(fabs(tiled_input_0[bi][0]), fabs(tiled_input_0[bi][1])), fmax(fabs(tiled_input_0[bi][2]), fabs(tiled_input_0[bi][3]))); - } - unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { - dq_max_input[0][bi] = sub_group_reduce_max(max[0][bi]); - dq_max_input[1][bi] = sub_group_reduce_max(max[1][bi]); - } - de_quantize_scale[0] = dq_max_input[0] / quan; - de_quantize_scale[1] = dq_max_input[1] / quan; - - // Packing 4 of converted inputs to integer type - unroll_for (uint bi = 0; bi < HALF_TILE_B; ++bi) { - packed_in_0[bi] = as_int(CAT(convert_, MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE))(tiled_input_0[bi] / de_quantize_scale[batch_sglid][bi])); - } - #endif - // __local SLM_FILTER_VEC* char_slm_weight = (__local SLM_FILTER_VEC*)wei_local_mem; __local int* char_slm_weight = (__local int*)wei_local_mem; @@ -939,7 +912,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // 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) { - // uchar4 wei_packed = as_uchar4(_sub_group_block_read_uc4((const __global uchar *)(weights) + (weights_idx))); 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_MIXED_INT4(DQ_TYPE, *((uint4x8_t *)&wei_packed)); @@ -1037,9 +1009,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif #if USE_SLM && COMPRESSED_WEIGHTS_INT4 && DYNAMIC_QUANTIZE - // Error if TILE_OFM != 2 // Compute input * weight : packed char4 type - // char4 input_val = AS_DQ_TYPE_4(_sub_group_shuffle(packed_in_0[0], ki)); 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; @@ -1047,19 +1017,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( 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); - // input_val = as_char4(_sub_group_shuffle(packed_in_0[(bi+1) / 2], ((bi+1) % 2) * 8 + ki)); } - // !DYNAMIC_QUANTIZE - // 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) { - // half4 in_val = as_half4(_sub_group_shuffle(((int2*)(&tiled_input_0[bi/2]))[0], (bi % 2) * 8 + ki)); - // unroll_for (uint kii = 0; kii < TILE_K; ++kii) { - // ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[0] += in_val[kii] * convert_half(first_weight[kii]); - // ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[1] += in_val[kii] * convert_half(second_weight[kii]); - // } - // } #else unroll_for (uint bi = 0; bi < TILE_B; ++bi) { unroll_for (uint kii = 0; kii < TILE_K; ++kii) { @@ -1121,7 +1079,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( } } #endif - } // Done main compute loop : ni + } // Main compute loop : ni // ===================================================================================================================================== // Leftovers @@ -1443,14 +1401,6 @@ KERNEL(fc)( ); } else { if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE) { - // fc_bf_tiled_kernel_dyn_quan kernel is for dynamic quantizing. DECOMPRESSION_SCALE_POST_OP is required. - // It shows better performance with weight SLM and 4bit weight. - // if (get_sub_group_local_id() == 0 && get_group_id(0) == 0 && get_group_id(2) == 0 && - // get_local_id(0) == 0 && get_local_id(2) == 0) { - // printf(">>>> DYNAMIC : ELEMENTS_COUNT(%d) batch(%d) TILE_B(%d) OSV32_ISV2(%d) TILE_OFM(%d) => group0(%d) local0(%d) / group2(%d) local2(%d)\n", - // (int)MAIN_LOOP_ELEMENTS_COUNT, (int)INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM, (int)TILE_B, (int)FILTER_LAYOUT_OS_IS_YX_OSV32_ISV2, (int)TILE_OFM, - // (int)get_global_size(0), (int)get_local_size(0), (int)get_global_size(2), (int)get_local_size(2)); - // } FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, @@ -1502,12 +1452,6 @@ KERNEL(fc)( } #else if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE) { - // if (get_sub_group_local_id() == 0 && get_group_id(0) == 0 && get_group_id(2) == 0 && - // get_local_id(0) == 0 && get_local_id(2) == 0) { - // printf(">>>> STATIC : MAIN_LOOP_ELEMENTS_COUNT(%d) batch(%d) => group0(%d) local0(%d) / group2(%d) local2(%d)\n", - // (int)MAIN_LOOP_ELEMENTS_COUNT, (int)INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM, - // (int)get_global_size(0), (int)get_local_size(0), (int)get_global_size(2), (int)get_local_size(2)); - // } FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, 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 11db0cbe049..943f37e443a 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 @@ -14,37 +14,41 @@ static constexpr size_t min_slm_size = 256; namespace kernel_selector { -static std::pair get_input_threads(const fully_connected_params& params) { - size_t ifm_threads = params.inputs[0].Feature().v; - size_t batch_threads = params.inputs[0].Batch().v; +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) { - ifm_threads = params.inputs[0].Y().v; - batch_threads = params.inputs[0].Batch().v * params.inputs[0].Feature().v; + input_f = params.inputs[0].Y().v; + input_batch = params.inputs[0].Batch().v * params.inputs[0].Feature().v; } - return {batch_threads, ifm_threads}; + return {input_batch, input_f}; } -static std::pair get_output_threads(const fully_connected_params& params, uint32_t tile_ofm) { - size_t feature_threads = CeilDiv(params.outputs[0].Feature().v, tile_ofm * simd); - size_t batch_threads = params.outputs[0].Batch().v; +static std::pair get_output_aligned_bf_size(const fully_connected_params& params, bool is_aligned, uint32_t align_b = 1, uint32_t align_f = 1) { + size_t output_f = (is_aligned == 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) { - feature_threads = CeilDiv(params.outputs[0].Y().v, tile_ofm * simd); - batch_threads = params.outputs[0].Batch().v * params.outputs[0].Feature().v; + output_f = (is_aligned == 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; } - return {batch_threads, feature_threads}; + output_b = (is_aligned == true) ? CeilDiv(output_b, align_b) : output_b; + + return {output_b, output_f}; } // DYNAMIC_QUANTIZE static bool is_dynamic_quantize(const fully_connected_params& params) { - auto threads = get_input_threads(params); - auto batch_threads = threads.first; - auto ifm_threads = threads.second; + 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) && (ifm_threads % quantize_grp_size == 0) && - (params.is_shape_agnostic || (params.inputs[0].Batch().v > 1 && batch_threads > min_slm_size)) && + 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)) @@ -201,12 +205,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) && @@ -275,14 +276,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); @@ -390,12 +387,10 @@ FullyConnected_bf_tiled::SetDefault(const fully_connected_params& params, int au auto tparams = GetAutoTuneParams(params, kernel_type, autoTuneIndex); - auto threads = get_output_threads(params, tparams.tile_ofm); + 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; - batch_threads = CeilDiv(batch_threads, tparams.tile_b); - 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) const bool can_use_slm = tparams.kernel_type == KernelType::SLM; @@ -423,9 +418,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) @@ -579,9 +572,7 @@ 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); - 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; + size_t output_batch = get_output_aligned_bf_size(prim_params, false).first; // Get index of the added shape-agnostic kernel int kernel_num = kd.kernels.size() - 1; @@ -601,18 +592,11 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { GPU_DEBUG_TRACE_DETAIL << "FC bf tiled: " << (execute_kernel_idx == kernel_num ? "SLM" : "Default") << " shape-agnostic kernel version " << "will be used for batch size = " << output_batch << "\n"; - // printf(" -- previous kernel idx(%d) in kernels.size(%lu) for sa : gws(%lu, %lu, %lu)\n", execute_kernel_idx, kd.kernels.size(), - // kd.kernels[execute_kernel_idx].params.workGroups.global[0], kd.kernels[execute_kernel_idx].params.workGroups.global[1], kd.kernels[execute_kernel_idx].params.workGroups.global[2]); - auto dispatchData = SetDefault(prim_params, -1, (int)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); - // printf(" -- update kernel idx(%d) for sa : gws(%lu, %lu, %lu) lws(%lu, %lu, %lu) physical(%lu)\n", execute_kernel_idx, - // dispatchData.gws[0], dispatchData.gws[1], dispatchData.gws[2], dispatchData.lws[0], dispatchData.lws[1], dispatchData.lws[2], prim_params.inputs[0].PhysicalSize()); - // printf(" -- relevant values : Physical(%lu) Y(%lu) TILE_B(%u) gws(%lu) lws(%lu) \n", prim_params.inputs[0].PhysicalSize(), prim_params.inputs[0].Y().v, dispatchData.tile_m, dispatchData.gws[2], dispatchData.lws[2]); - if (!kd.internalBufferSizes.empty()) { // Pre-quantizing kernel was generated. Update the kernel and intermediate buffers or disable it. if (execute_type == KernelType::DEFAULT) { @@ -621,21 +605,15 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { kd.kernels[0].skip_execution = false; size_t ifm_threads = get_input_threads(prim_params).second; size_t input_size = ifm_threads * dispatchData.tile_m * dispatchData.gws[2]; - // printf(" -- Require update intermediate buffer\n"); - // printf(" -- previous kernel(%d) for sa : input size(%lu) gws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, input_size, - // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], (int)kd.internalBufferSizes.size()); if (kd.kernels[0].params.workGroups.global[0] < (input_size / quantize_grp_size)) { kd.internalBufferSizes.clear(); - kd.internalBufferSizes.push_back(input_size * 2); - kd.internalBufferSizes.push_back(input_size / quantize_grp_size * 2); + 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}; - // printf(" -- update kernel(%d) for sa : gws(%lu, %lu, %lu) lws(%lu, %lu, %lu) kd.internalBufferSizes.size(%d)\n", 0, - // kd.kernels[0].params.workGroups.global[0], kd.kernels[0].params.workGroups.global[1], kd.kernels[0].params.workGroups.global[2], - // kd.kernels[0].params.workGroups.local[0], kd.kernels[0].params.workGroups.local[1], kd.kernels[0].params.workGroups.local[2], (int)kd.internalBufferSizes.size()); } } }; @@ -665,21 +643,18 @@ KernelsData FullyConnected_bf_tiled::GetTunedKernelsDataByIndex(const Params &pa } KernelsData kernels_data; - // printf(">> GetTunedKernelsDataByIndex \n"); if (is_dynamic_quantize(fc_params)) { // Use seperate 2 kernels for dynamic quantizing : quantizing_kernel + fc_kernel - // First kernel : Dynamic quantizing with 32 group size - // Second kernel : fully connected with char type size using quantized inputs and scale values + // 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 : fully connected 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, - 0); + autoTuneIndex); OPENVINO_ASSERT(!kernels_data.empty() && !kernels_data[0].kernels.empty(), "[GPU] Error to create multi kernel for dynamic quantizing."); - // GPU_DEBUG_TRACE_DETAI << " Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; - // std::cout << " >> Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; if (params.is_shape_agnostic) GetUpdateDispatchDataFunc(kernels_data[0]); } else { @@ -690,9 +665,6 @@ KernelsData FullyConnected_bf_tiled::GetTunedKernelsDataByIndex(const Params &pa autoTuneIndex, 0); - // GPU_DEBUG_TRACE_DETAI << " No Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; - // std::cout << " >> No! Dynamic quantizing : Use kernels size " << kernels_data[0].kernels.size() << "\n"; - if (params.is_shape_agnostic) { auto tparams = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex); auto can_select_slm_kernel = tparams.kernel_type == KernelType::SLM; @@ -751,8 +723,7 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, DataLayout dl, WeightsLayout wl, const std::string exeMode, - int autoTuneIndex, - int kernel_number) const { + int autoTuneIndex) const { if (!Validate(params)) { return KernelsData(); } @@ -788,8 +759,11 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, inputs_count++; } - // Quantize kernel + // 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; @@ -799,14 +773,10 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, quan_kernel.params.workGroups.local = dyn_quan_dispatch.lws; quan_kernel.skip_execution = false; - // printf(" >> GetMultiKernelsData : Input0.Batch(%lu) Input0.Feature(%lu) Input0.Y(%lu) Input0.pysical(%lu)\n", - // fc_params.inputs[0].Batch().v, fc_params.inputs[0].Feature().v, fc_params.inputs[0].Y().v, fc_params.inputs[0].PhysicalSize()); - // printf(" -- Quantizing kernel : gws(%ld, %ld, %ld), lws(%ld, %ld, %ld)\n", dyn_quan_dispatch.gws[0], dyn_quan_dispatch.gws[1], dyn_quan_dispatch.gws[2], - // dyn_quan_dispatch.lws[0], dyn_quan_dispatch.lws[1], dyn_quan_dispatch.lws[2]); - 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)); + quan_cldnn_jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", quantize_grp_size)); auto quan_jit = CreateJit(kernelName, quan_cldnn_jit, quan_entry_point); @@ -834,7 +804,7 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, } kd.internalBufferDataType = Datatype::F16; - // FC kernel for dynamic quantized input + // 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); @@ -867,8 +837,8 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, 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) { - // std::cout << " -- is_shape_agnostic quantizing : previous kernels size " << kd.kernels.size() << "\n"; kd.kernels.resize(kernel_number + 1); auto entry_point = GetEntryPoint(kernelName, fc_params.layerID, params, kernel_number); @@ -889,16 +859,15 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, slm_params.exec_options, true, !fc_params.bias.empty(), - inputs_count, // input : intermediate buffers (quantized input + scale value) + inputs_count, GetFusedPrimitiveInputsCount(params), - 1, // Output + 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}); } - // TODO Pass estimated time only through DispatchData kd.autoTuneIndex = autoTuneIndex; return {kd}; } 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 7440b99ae99..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 @@ -35,8 +35,7 @@ public: DataLayout dl, WeightsLayout wl, const std::string exeMode = EXE_MODE_DEFAULT, - int autoTuneIndex = -1, - int kernel_number = 0) const; + int autoTuneIndex = -1) const; struct tune_params { tune_params(unsigned tile_b, From 04c51a979f280078dbe2a49f84f1ab1211e5f430 Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Sun, 2 Jun 2024 15:34:17 +0900 Subject: [PATCH 09/16] Bugfix Signed-off-by: Min, Byung-il --- .../fully_connected/fully_connected_kernel_bf_tiled.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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 943f37e443a..7005ef3ff02 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 @@ -603,8 +603,8 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { kd.kernels[0].skip_execution = true; } else { kd.kernels[0].skip_execution = false; - size_t ifm_threads = get_input_threads(prim_params).second; - size_t input_size = ifm_threads * dispatchData.tile_m * dispatchData.gws[2]; + 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.kernels[0].params.workGroups.global[0] < (input_size / quantize_grp_size)) { kd.internalBufferSizes.clear(); From 606ea65ad8b0e29d38250fd3776d08ac6d5fee02 Mon Sep 17 00:00:00 2001 From: "Min, Byungil" Date: Sun, 2 Jun 2024 03:45:35 +0900 Subject: [PATCH 10/16] [GPU] Add config dynamic_quantization_group_size Signed-off-by: Min, Byungil --- .../intel_gpu/src/graph/impls/ocl/fully_connected.cpp | 2 ++ .../cl_kernels/fully_connected_gpu_bf_tiled.cl | 2 -- .../fully_connected_kernel_bf_tiled.cpp | 11 ++++------- .../kernels/fully_connected/fully_connected_params.h | 1 + src/plugins/intel_gpu/src/plugin/compiled_model.cpp | 3 ++- src/plugins/intel_gpu/src/plugin/plugin.cpp | 1 + .../intel_gpu/src/runtime/execution_config.cpp | 1 + 7 files changed, 11 insertions(+), 10 deletions(-) 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 37ef63d5a22..02571263369 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 @@ -25,8 +25,6 @@ KERNEL(quantize_input)( __global char* quantized_input, __global INPUT0_TYPE* de_quan_scale) { const uint offset = get_global_id(0); - if (offset != (get_group_id(0) * 16 + get_local_id(0))) - printf("!!!!!!!!!!!!!!!!!!!!!!\n") uint input_offset = offset * QUANTIZE_GROUP_SIZE; half4 input_0[8]; 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 7005ef3ff02..2c16d57476d 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 @@ -42,6 +42,9 @@ static std::pair get_output_aligned_bf_size(const fully_connecte // DYNAMIC_QUANTIZE static bool is_dynamic_quantize(const fully_connected_params& params) { + if (params.dynamic_quantization_group_size < 32) + return false; + auto threads = get_input_bf_size(params); auto input_b = threads.first; auto input_f = threads.second; @@ -54,12 +57,6 @@ static bool is_dynamic_quantize(const fully_connected_params& params) { (params.decompression_zero_point.Feature().v == 1)) return true; - GPU_DEBUG_TRACE_DETAIL << " Dynamic quantizing for FC : scale_group_size " << scale_group_size << ", Shape_Agnostic(" << (int)params.is_shape_agnostic << - ") : Input (" << kernel_selector::toString(params.inputs[0].GetDType()) << - ") B: " << params.inputs[0].Batch().v << ", F: " << params.inputs[0].Feature().v << ", Y: " << params.inputs[0].Y().v << - ", Weight " << kernel_selector::toString(params.weights.GetDType()) << ", zp_f " << params.decompression_zero_point.Feature().v << - ", format : " << kernel_selector::toString(params.outputs[0].GetLayout()) << std ::endl; - return false; } @@ -589,7 +586,7 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { // 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 == kernel_num ? "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, (int)execute_type); 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..0f4dfcfddd7 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; + int 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 dd94f5347b6..5364e89fe0b 100644 --- a/src/plugins/intel_gpu/src/plugin/plugin.cpp +++ b/src/plugins/intel_gpu/src/plugin/plugin.cpp @@ -554,6 +554,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/execution_config.cpp b/src/plugins/intel_gpu/src/runtime/execution_config.cpp index 66b8d3e70ca..e27055b6823 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, 32), // Legacy API properties std::make_tuple(ov::intel_gpu::nv12_two_inputs, false), From 85585a2de1846f3d0c9d11b49e99271f044a266a Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Tue, 4 Jun 2024 10:25:51 +0900 Subject: [PATCH 11/16] Minor fix + Removed hard-coded values. + Minor bugfix Signed-off-by: Min, Byung-il --- .../cl_kernels/fully_connected_gpu_bf_tiled.cl | 6 +++--- .../cl_kernels/include/batch_headers/int4_utils.cl | 2 +- .../fully_connected_kernel_bf_tiled.cpp | 11 +++++++---- 3 files changed, 11 insertions(+), 8 deletions(-) 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 02571263369..6ae910b15f8 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 @@ -860,19 +860,19 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( for (uint ni = 0; ni < iterations; ++ni) { #if USE_SLM && COMPRESSED_WEIGHTS_INT4 & DYNAMIC_QUANTIZE uint in_offset = input_offset + (idx_sglid + batch_sglid * TILE_IN_B_PITCH); - uint scale_offset = input_offset / 32; + 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/32]; + 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/32 * 2); + scale_offset += (TILE_IN_B_PITCH/QUANTIZE_GROUP_SIZE * 2); } input_offset += TILE_IFM * SIMD; 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 9403258a16b..b373eac6c88 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 @@ -208,4 +208,4 @@ inline uchar8 unpack_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadabl #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_MIXED_INT4x2(target_type, value) CAT(unpack_mixed_to_, target_type)(value) -#define UNPACK_MIXED_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_mixed_to_, target_type), _osv32_isv2)(value) \ No newline at end of file +#define UNPACK_MIXED_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_mixed_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 2c16d57476d..4ed23f5fe27 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 @@ -42,7 +42,10 @@ static std::pair get_output_aligned_bf_size(const fully_connecte // DYNAMIC_QUANTIZE static bool is_dynamic_quantize(const fully_connected_params& params) { - if (params.dynamic_quantization_group_size < 32) + if (params.inputs[0].GetFirstElementOffset() != 0) + return false; + + if (params.dynamic_quantization_group_size < quantize_grp_size) return false; auto threads = get_input_bf_size(params); @@ -491,6 +494,7 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para } jit.AddConstant(MakeJitConstant("DQ_TYPE", "char")); + jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", quantize_grp_size)); jit.AddConstant(MakeJitConstant("SIMD", simd)); jit.AddConstant(MakeJitConstant("TILE_B", dispatchData.tile_m)); @@ -773,7 +777,6 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, 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)); - quan_cldnn_jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", quantize_grp_size)); auto quan_jit = CreateJit(kernelName, quan_cldnn_jit, quan_entry_point); @@ -795,8 +798,8 @@ KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params ¶ms, 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() * 2); - kd.internalBufferSizes.push_back(fc_params.inputs[0].PhysicalSize() / 32 * 2); + 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; From a3c765c49ccb977ffc84d7ec222a868b201b9acc Mon Sep 17 00:00:00 2001 From: "Min, Byungil" Date: Mon, 3 Jun 2024 20:49:51 +0900 Subject: [PATCH 12/16] [GPU] Update fc_gpu_bf_tiled kernel + Modified bf_tiled_dyn_quan + Removed redundancy + Set default value of dynamic_quantization_group_size 0 Signed-off-by: Min, Byungil --- .../fully_connected_gpu_bf_tiled.cl | 345 ++++++------------ .../fully_connected/fully_connected_params.h | 2 +- .../src/runtime/execution_config.cpp | 2 +- 3 files changed, 116 insertions(+), 233 deletions(-) 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 6ae910b15f8..2aa0c395479 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 @@ -717,13 +717,12 @@ inline void FUNC(fc_bf_tiled_kernel_default)( #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) +#if USE_SLM && DYNAMIC_QUANTIZE inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_ARG const __global INPUT0_TYPE* input, -#if DYNAMIC_QUANTIZE __global char* quantized_input, __global INPUT0_TYPE* scale, -#endif #if DECOMPRESSION_SCALE_TERM const __global DECOMPRESSION_SCALE_TYPE* decompression_scale, #endif @@ -732,9 +731,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif __global OUTPUT_TYPE* output, const __global FILTER_TYPE* weights -#if USE_SLM , __local int* wei_local_mem -#endif #if BIAS_TERM , const __global BIAS_TYPE* biases #endif @@ -742,14 +739,8 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( , FUSED_OPS_DECLS #endif ) { -#if USE_SLM uint gid = (uint)get_group_id(0); uint local_id = (uint)get_local_id(2); -#else - uint gid = (uint)get_group_id(0); -#endif - - uint sglid = (uint)get_sub_group_local_id(); // Dispatch as bs_fs_bsv_fsv, where bsv = DISPATCH_BSV and fsv = DISPATCH_FSV. @@ -764,14 +755,8 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( FILTER_VEC_TYPE wei = 0; -#if USE_SLM uint out_f = gid * (TILE_OFM * SIMD); uint out_b = LWS_BATCHES * TILE_B * (uint)get_group_id(2) + local_id * TILE_B; -#else - // FILTER_VEC_TYPE wei = 0; - uint out_f = (feature_mega_block * DISPATCH_FSV + feature_mini_block) * (TILE_OFM * SIMD); - uint out_b = ((batch_mega_block * DISPATCH_BSV + batch_mini_block) * TILE_B); -#endif #if OUTPUT_3D uint out_b0 = out_b / OUTPUT_FEATURE_NUM; @@ -781,22 +766,14 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( uint input_offset = out_b * TILE_IN_B_PITCH + INPUT0_OFFSET; #endif -#if COMPRESSED_WEIGHTS_INT4 uint weights_offset = out_f * (INPUT_ELEMENTS_COUNT / 2); -#else - uint weights_offset = out_f * INPUT_ELEMENTS_COUNT; -#endif ACCUMULATOR_VEC_TYPE acc[TILE_B] = { }; - #if USE_SLM && COMPRESSED_WEIGHTS_INT4 && DYNAMIC_QUANTIZE - // 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]; - #else - INPUT_VEC_TYPE in_0[TILE_B] = { }; - #endif + // 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 @@ -831,24 +808,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( ACCUMULATOR_TYPE* d_zps = (ACCUMULATOR_TYPE*)(&d_zp); #endif -#if REALIGN_FP16_OFFSET - // For fp16 we need to ensure that all block reads are aligned to 4 byte (2 words) boundary. - // To do this solve first input feature separately. - { - INPUT0_TYPE tmp_input = input[input_offset + get_sub_group_local_id() % TILE_B * TILE_IN_B_PITCH]; - ACCUMULATOR_VEC_TYPE tmp_wei = TO_ACCUMULATOR_VEC_TYPE(BLOCK_READN(FILTER_TYPE, TILE_OFM, weights, weights_offset)); - #if COMPRESSED_WEIGHTS - tmp_wei = (tmp_wei - d_zp) * d_scale; - #endif - unroll_for(uint bi = 0; bi < TILE_B; ++bi) { - acc[bi] = _sub_group_shuffle(tmp_input, bi) * tmp_wei; - } - - weights_offset += TILE_OFM * SIMD; - input_offset += 1; - } -#endif - // ===================================================================================================================================== // Main computation loop uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD); @@ -858,176 +817,114 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( __attribute__((opencl_unroll_hint(1))) for (uint ni = 0; ni < iterations; ++ni) { - #if USE_SLM && COMPRESSED_WEIGHTS_INT4 & DYNAMIC_QUANTIZE - 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)]; + 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]); + // 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); - } + // Next batch + in_offset += (TILE_IN_B_PITCH * 2); + scale_offset += (TILE_IN_B_PITCH/QUANTIZE_GROUP_SIZE * 2); + } - input_offset += TILE_IFM * SIMD; + input_offset += TILE_IFM * SIMD; - // Packing - MAKE_VECTOR_TYPE(int, TILE_B) acc_tmp[TILE_OFM] = { }; - #else - #define LOAD_IN_0(bi) do { \ - in_0[bi] = INPUT_BLOCK_READ(input, input_offset); \ - input_offset += TILE_IN_B_PITCH; \ - } while (false) + // Packing + MAKE_VECTOR_TYPE(int, TILE_B) acc_tmp[TILE_OFM] = { }; - CONST_LOOP(TILE_B, LOAD_IN_0); - #undef LOAD_IN_0 - input_offset += TILE_IFM * SIMD - TILE_IN_B_PITCH * TILE_B; - - ACCUMULATOR_VEC_TYPE acc_tmp[TILE_B] = { }; + #if TILE_OFM != 2 + #error "FC bf_tiled kernel: can't use SLM optimization with TILE_OFM != 2" #endif - #if USE_SLM && COMPRESSED_WEIGHTS_INT4 - #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 SLM_FILTER_VEC* char_slm_weight = (__local SLM_FILTER_VEC*)wei_local_mem; - __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_MIXED_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; - + // 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 // USE_SLM && COMPRESSED_WEIGHTS_INT4 + #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_MIXED_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 USE_SLM && COMPRESSED_WEIGHTS_INT4 - #if (TILE_K != 1) && (TILE_K != 2) && (TILE_K != 4) - #error "FC bf_tiled kernel: unsupported TILE_K size for SLM kernel" - #endif - #elif COMPRESSED_WEIGHTS_INT4 - FILTER_PACKED_VEC_TYPE wei_packed = FILTER_BLOCK_READ(weights, weights_offset); - wei = UNPACK_INT4(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&wei_packed)); - #else - wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); + #if TILE_K != 4 + #error "FC bf_tiled kernel: unsupported TILE_K size for SLM kernel" #endif - #if COMPRESSED_WEIGHTS && !USE_SLM - ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); - unroll_for(uint kii = 0; kii < TILE_K; ++kii) { - unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { - const uint offset_ofm = out_f + fi*SIMD + sglid; - // Valid only if DECOMPRESSION_SCALE_POST_OP is enabled - ACCUMULATOR_TYPE ds = ACCUMULATOR_VAL_ONE; - - #if DECOMPRESSION_ZP_TERM - #if DECOMPRESSION_ZP_SCALAR - ACCUMULATOR_TYPE dzp = DECOMPRESSION_ZP_VALUE; - #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 - const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + - ((kii + ki*TILE_K + ni*TILE_IFM*SIMD) / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; - ACCUMULATOR_TYPE dzp = decompression_zp[zp_offset]; - #else - ACCUMULATOR_TYPE dzp = d_zps[fi % DECOMPRESSION_ZP_LENGTH]; - #endif - #else - ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; - #endif - w[W_IDX] = (w[W_IDX] - dzp) * ds; - } - } - #endif - - #if USE_SLM && COMPRESSED_WEIGHTS_INT4 && DYNAMIC_QUANTIZE - // 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); - } - #else - unroll_for (uint bi = 0; bi < TILE_B; ++bi) { - unroll_for (uint kii = 0; kii < TILE_K; ++kii) { - const uint total_k = ki * TILE_K + kii; - INPUT0_TYPE in_val = _sub_group_shuffle(((INPUT0_TYPE*)(&in_0[bi]))[total_k / SIMD], total_k % SIMD); - unroll_for (uint fi = 0; fi < TILE_OFM; ++fi) { - ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[W_IDX]; - } - } - } - - #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; @@ -1044,13 +941,8 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; #endif - #if USE_SLM && DYNAMIC_QUANTIZE - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; - acc_tmp[fi][bi] = 0; - #else - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; - acc_tmp[bi][fi] = 0; - #endif + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; + acc_tmp[fi][bi] = 0; } } #endif @@ -1069,11 +961,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; #endif - #if USE_SLM && DYNAMIC_QUANTIZE - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; - #else - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += ((ACCUMULATOR_TYPE*)(&acc_tmp[bi]))[fi] * ds; - #endif + ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi]; } } #endif @@ -1233,7 +1121,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( } // ===================================================================================================================================== } - +#endif REQD_SUB_GROUP_SIZE(SIMD) KERNEL(fc)( @@ -1259,8 +1147,11 @@ KERNEL(fc)( #endif ) { #if USE_SLM + #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; @@ -1398,14 +1289,12 @@ KERNEL(fc)( #endif ); } else { - if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE) { + #if USE_SLM && DYNAMIC_QUANTIZE FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, - #if DYNAMIC_QUANTIZE quantized_input, de_quan_scale, - #endif #if DECOMPRESSION_SCALE_TERM decompression_scale, #endif @@ -1414,9 +1303,7 @@ KERNEL(fc)( #endif output, weights - #if USE_SLM , dq_wei_local_mem - #endif #if BIAS_TERM , biases #endif @@ -1424,7 +1311,7 @@ KERNEL(fc)( , FUSED_OPS_ARGS #endif ); - } else { + #else FUNC_CALL(fc_bf_tiled_kernel_default)( OPTIONAL_SHAPE_INFO_TENSOR input, @@ -1446,17 +1333,15 @@ KERNEL(fc)( , FUSED_OPS_ARGS #endif ); - } + #endif } #else - if ((INPUT0_FEATURE_NUM*INPUT0_BATCH_NUM > 256) && USE_SLM && DYNAMIC_QUANTIZE) { + #if USE_SLM && DYNAMIC_QUANTIZE FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_TENSOR input, - #if DYNAMIC_QUANTIZE quantized_input, de_quan_scale, - #endif #if DECOMPRESSION_SCALE_TERM decompression_scale, #endif @@ -1465,9 +1350,7 @@ KERNEL(fc)( #endif output, weights - #if USE_SLM , dq_wei_local_mem - #endif #if BIAS_TERM , biases #endif @@ -1475,7 +1358,7 @@ KERNEL(fc)( , FUSED_OPS_ARGS #endif ); - } else { + #else FUNC_CALL(fc_bf_tiled_kernel_default)( OPTIONAL_SHAPE_INFO_TENSOR input, @@ -1497,7 +1380,7 @@ KERNEL(fc)( , FUSED_OPS_ARGS #endif ); - } + #endif #endif } #endif // !FC_KERNEL_DYNAMIC_QUANTIZE 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 0f4dfcfddd7..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,7 +15,7 @@ struct fully_connected_params : public weight_bias_params { fully_connected_params() : weight_bias_params(KernelType::FULLY_CONNECTED) {} QuantizationType quantization = QuantizationType::NONE; - int dynamic_quantization_group_size = 0; + 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/runtime/execution_config.cpp b/src/plugins/intel_gpu/src/runtime/execution_config.cpp index 0f1ae84aa6f..b7bb9947717 100644 --- a/src/plugins/intel_gpu/src/runtime/execution_config.cpp +++ b/src/plugins/intel_gpu/src/runtime/execution_config.cpp @@ -56,7 +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, 32), + std::make_tuple(ov::hint::dynamic_quantization_group_size, 0), // Legacy API properties std::make_tuple(ov::intel_gpu::nv12_two_inputs, false), From a49d8acbbf527da95bbe2b601be3e7b3ab8b2d9d Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Wed, 5 Jun 2024 11:44:04 +0900 Subject: [PATCH 13/16] Pass CPPLINT Signed-off-by: Min, Byung-il --- .../cl_kernels/fully_connected_gpu_bf_tiled.cl | 6 +++--- .../fully_connected_kernel_bf_tiled.cpp | 13 ++++++++----- 2 files changed, 11 insertions(+), 8 deletions(-) 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 2aa0c395479..ea502f5a409 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 @@ -810,10 +810,10 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // ===================================================================================================================================== // Main computation loop - uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD); + const uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD); // Each sub-group loads 2 Batch - uint idx_sglid = (sglid * TILE_K) % 32; // same index for sglid 0~7 : to tile_k direction - uint batch_sglid = (sglid * TILE_K) / 32; // 0 to 1 : to batch direction + 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) { 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 4ed23f5fe27..cc672ce33b7 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 @@ -576,7 +576,9 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { size_t output_batch = get_output_aligned_bf_size(prim_params, false).first; // Get index of the added shape-agnostic kernel - int kernel_num = kd.kernels.size() - 1; + int kernel_offset = 0; + if (kd.kernels.size() == 3) + kernel_offset = 1; // quantize kernel exists // Choose one of the two shape agnostic kernels: N == added kernel number // - kd.kernels[N-1] for batches <= 240 (default version) @@ -584,8 +586,9 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { const auto default_alignment = 16; // 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) ? kernel_num : kernel_num - 1; - const auto skip_kernel_idx = (execute_type == KernelType::SLM) ? (kernel_num - 1) : kernel_num; + 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; @@ -593,7 +596,7 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { 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, (int)execute_type); + 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); @@ -648,7 +651,7 @@ KernelsData FullyConnected_bf_tiled::GetTunedKernelsDataByIndex(const Params &pa // 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 : fully connected kernel with KernelType::SLM. Quantized inputs and scale values would 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, From 725624b27ef631daa6b6334e5942b37f8d6f05f0 Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Thu, 6 Jun 2024 16:28:45 +0900 Subject: [PATCH 14/16] [GPU] Add test-cases for dynamic quantization Signed-off-by: Min, Byung-il --- .../fully_connected_kernel_bf_tiled.cpp | 6 ++- .../test_cases/fully_connected_gpu_test.cpp | 48 ++++++++++++------- 2 files changed, 37 insertions(+), 17 deletions(-) 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 cc672ce33b7..2782e82f384 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 @@ -57,8 +57,12 @@ static bool is_dynamic_quantize(const fully_connected_params& params) { (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)) + (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; } 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 8048c4aa535..1774f9fcdc5 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,14 +1255,14 @@ public: } } - void test_compressed_int4_scale_my(bool is_caching_test, bool is_dynamic) { + 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 = 512; + long int batch_num = batch; long int ifm_num = 1024; long int ofm_num = 4096; long int scales_group_size = 32; @@ -1270,6 +1270,7 @@ public: 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 }); @@ -1284,13 +1285,14 @@ public: 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{ {-1, ifm_num}, data_types::f16, format::bfyx } + 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; - auto get_implemented_results = [&]() { + // Implemented dynamic quantize kernel + auto get_ref_results = [&]() { topology topology( input_layout("input", in_layout), data("weights", weights_mem), @@ -1325,14 +1327,15 @@ public: 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::intel_gpu::force_implementations(ov::intel_gpu::ImplForcingMap{{"fc_prim", { format::bfyx, "fully_connected_gpu_bfyx_ref"}}})); + 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) { + if (is_dynamic && !engine.get_device_info().supports_immad) { auto inst = network->get_primitive("fc_prim"); auto impl = inst->get_impl(); - ASSERT_EQ(impl->get_kernels().size(), size_t(2)); // Two shape-agnostic kernels + 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); @@ -1344,14 +1347,14 @@ public: auto output_mem = outputs.begin()->second.get_memory(); cldnn::mem_lock output_ptr (output_mem, get_test_stream()); - auto impl_output_mem = get_implemented_results(); - cldnn::mem_lock output_ptr_impl (impl_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_impl.size(); ++i) { - auto abs_diff = std::abs(output_ptr_impl[i] - output_ptr[i]); + 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; @@ -3116,11 +3119,6 @@ TEST_F(fully_connected_gpu_tests, compressed_scale_zp_bias_cached) { this->test_compressed_scale_zp_bias(true); } -// Testing -TEST_F(fully_connected_gpu_tests, compressed_int4_scale_my) { - this->test_compressed_int4_scale_my(false, false); -} - TEST_F(fully_connected_gpu_tests, compressed_int4_scale) { this->test_compressed_int4_scale(false, false, 256); } @@ -3165,6 +3163,24 @@ 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_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_scale_bias) { this->test_compressed_scale_bias(false); } From 88460a900689a487b9035336f87b51a09e387576 Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Tue, 11 Jun 2024 18:17:47 +0900 Subject: [PATCH 15/16] [GPU] Apply comments and add test-cases, debug_config Signed-off-by: Min, Byung-il --- .../intel_gpu/runtime/debug_configuration.hpp | 1 + .../fully_connected_gpu_bf_tiled.cl | 82 +------------------ .../include/batch_headers/int4_utils.cl | 12 +-- .../fully_connected_kernel_bf_tiled.cpp | 31 ++++--- .../src/runtime/debug_configuration.cpp | 5 +- .../test_cases/fully_connected_gpu_test.cpp | 20 ++++- 6 files changed, 51 insertions(+), 100 deletions(-) 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 3ec28d1e32a..92c7da1fa85 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 @@ -139,6 +139,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/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 ea502f5a409..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 @@ -83,10 +83,10 @@ KERNEL(quantize_input)( // 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_MIXED_INT4 UNPACK_INT4x2_OSV32_ISV2 +#define UNPACK_TRANSPOSED_INT4 UNPACK_INT4x2_OSV32_ISV2 #else #define UNPACK_INT4 UNPACK_INT4x2 -#define UNPACK_MIXED_INT4 UNPACK_MIXED_INT4x2 +#define UNPACK_TRANSPOSED_INT4 UNPACK_TRANSPOSED_INT4x2 #endif // Macros for vectorized types. #define INPUT_VEC_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, TILE_IFM) @@ -701,6 +701,7 @@ 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) @@ -717,7 +718,6 @@ inline void FUNC(fc_bf_tiled_kernel_default)( #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) -#if USE_SLM && DYNAMIC_QUANTIZE inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( OPTIONAL_SHAPE_INFO_ARG const __global INPUT0_TYPE* input, @@ -855,7 +855,7 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( // 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_MIXED_INT4(DQ_TYPE, *((uint4x8_t *)&wei_packed)); + 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 @@ -967,80 +967,6 @@ inline void FUNC(fc_bf_tiled_kernel_dyn_quan)( #endif } // Main compute loop : ni - // ===================================================================================================================================== - // Leftovers -#if MAIN_LOOP_ELEMENTS_COUNT % (TILE_IFM * SIMD) != 0 - // Handle leftovers in normal case without alignment correction. - #define LEFTOVER_IFM (MAIN_LOOP_ELEMENTS_COUNT % (TILE_IFM * SIMD)) - { - #define LOAD_IN_0(bi) do { \ - in_0[bi] = INPUT_BLOCK_READ(input, input_offset); \ - input_offset += TILE_IN_B_PITCH; \ - } while (false) - - CONST_LOOP(TILE_B, LOAD_IN_0); - #undef LOAD_IN_0 - input_offset += TILE_IFM * SIMD - TILE_IN_B_PITCH * TILE_B; - unroll_for(uint ki = 0; ki < CEIL_DIV(LEFTOVER_IFM, TILE_K); ++ki) { - #if USE_SLM - FILTER_VEC_TYPE wei = 0; - #endif - - #if COMPRESSED_WEIGHTS_INT4 - FILTER_PACKED_VEC_TYPE wei_packed = FILTER_BLOCK_READ(weights, weights_offset); - wei = UNPACK_INT4(ACCUMULATOR_TYPE, *((INT4_PACKED_TYPE*)&wei_packed)); - #else - wei = TO_FILTER_VEC_TYPE(FILTER_BLOCK_READ(weights, weights_offset)); - #endif - - #if COMPRESSED_WEIGHTS - ACCUMULATOR_TYPE* w = (ACCUMULATOR_TYPE*)(&wei); - unroll_for(uint kii = 0; kii < TILE_K; ++kii) { - unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) { - uint offset_ofm = out_f + fi*SIMD + get_sub_group_local_id(); - #if DECOMPRESSION_SCALE_GROUPS_NUM > 1 - const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH + - ((kii + ki*TILE_K + iterations*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH; - ACCUMULATOR_TYPE ds = decompression_scale[scale_offset]; - #else - ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH]; - #endif - - #if DECOMPRESSION_ZP_TERM - #if DECOMPRESSION_ZP_SCALAR - ACCUMULATOR_TYPE dzp = DECOMPRESSION_ZP_VALUE; - #elif DECOMPRESSION_ZP_GROUPS_NUM > 1 - const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH + - ((kii + ki*TILE_K + iterations*TILE_IFM*SIMD) / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH; - ACCUMULATOR_TYPE dzp = decompression_zp[zp_offset]; - #else - ACCUMULATOR_TYPE dzp = d_zps[fi % DECOMPRESSION_ZP_LENGTH]; - #endif - #else - ACCUMULATOR_TYPE dzp = ACCUMULATOR_VAL_ZERO; - #endif - w[W_IDX] = (w[W_IDX] - dzp) * ds; - } - } - #endif - weights_offset += TILE_K_OFM_PACKED * SIMD; - - unroll_for (uint kii = 0; kii < TILE_K; ++kii) { - unroll_for (uint fi = 0; fi < TILE_OFM; ++fi) { - unroll_for (uint bi = 0; bi < TILE_B; ++bi) { - const uint total_k = ki * TILE_K + kii; - if (total_k < LEFTOVER_IFM) { - INPUT0_TYPE in_val = _sub_group_shuffle(((INPUT0_TYPE*)(&in_0[bi]))[total_k / SIMD], total_k % SIMD); - ((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += in_val * ((ACCUMULATOR_TYPE*)(&wei))[W_IDX]; - } - } - } - } - } - } - #undef LEFTOVER_IFM -#endif // MAIN_LOOP_ELEMENTS_COUNT % (TILE_IFM * SIMD) != 0 - // ===================================================================================================================================== // Post-processing: bias, activation, fused-ops ACTIVATION_VEC_TYPE activated[TILE_B] = { }; 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 b373eac6c88..caf656e6fe3 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 @@ -64,7 +64,7 @@ inline char8 unpack_to_char(uint4x8_t v) __attribute__((overloadable)) { return (char8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1); } -inline char8 unpack_mixed_to_char(uint4x8_t v) __attribute__((overloadable)) { +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); @@ -72,7 +72,7 @@ inline char8 unpack_mixed_to_char(uint4x8_t v) __attribute__((overloadable)) { return (char8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); } -inline uchar8 unpack_mixed_to_uchar(uint4x8_t v) __attribute__((overloadable)) { +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); @@ -80,7 +80,7 @@ inline uchar8 unpack_mixed_to_uchar(uint4x8_t v) __attribute__((overloadable)) { return (uchar8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); } -inline char8 unpack_mixed_to_char_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) { +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); @@ -88,7 +88,7 @@ inline char8 unpack_mixed_to_char_osv32_isv2(uint4x8_t v) __attribute__((overloa return (char8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1); } -inline uchar8 unpack_mixed_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) { +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); @@ -207,5 +207,5 @@ inline uchar8 unpack_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadabl #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_MIXED_INT4x2(target_type, value) CAT(unpack_mixed_to_, target_type)(value) -#define UNPACK_MIXED_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_mixed_to_, target_type), _osv32_isv2)(value) +#define UNPACK_MIXED_INT4x2(target_type, value) CAT(unpack_transposed_to_, target_type)(value) +#define UNPACK_MIXED_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 2782e82f384..f63820ffb99 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 @@ -26,26 +26,32 @@ static std::pair get_input_bf_size(const fully_connected_params& return {input_batch, input_f}; } -static std::pair get_output_aligned_bf_size(const fully_connected_params& params, bool is_aligned, uint32_t align_b = 1, uint32_t align_f = 1) { - size_t output_f = (is_aligned == true) ? CeilDiv(params.outputs[0].Feature().v, align_f) : params.outputs[0].Feature().v; +static std::pair get_output_aligned_bf_size(const fully_connected_params& params, bool needs_align, uint32_t align_b = 1, uint32_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 = (is_aligned == true) ? CeilDiv(params.outputs[0].Y().v, align_f) : params.outputs[0].Y().v; + 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 = (is_aligned == true) ? CeilDiv(output_b, align_b) : output_b; + output_b = (needs_align == true) ? CeilDiv(output_b, align_b) : output_b; return {output_b, output_f}; } // DYNAMIC_QUANTIZE -static bool is_dynamic_quantize(const fully_connected_params& params) { +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 (params.dynamic_quantization_group_size < quantize_grp_size) + if (dynamic_quantization_group_size < quantize_grp_size) return false; auto threads = get_input_bf_size(params); @@ -306,7 +312,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 && !is_dynamic_quantize(params)) { + 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)); } @@ -490,16 +496,15 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para } // Validated perf gain, Dynamic quantize force enable SCALE_POST_OP for char type multiplication - if (is_dynamic_quantize(params) && dispatchData.tile_m > 1 && dispatchData.tile_n == 2) { + 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("DQ_TYPE", "char")); - jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", quantize_grp_size)); - jit.AddConstant(MakeJitConstant("SIMD", simd)); jit.AddConstant(MakeJitConstant("TILE_B", dispatchData.tile_m)); jit.AddConstant(MakeJitConstant("HALF_TILE_B", dispatchData.tile_m/2)); @@ -614,7 +619,7 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const { 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.kernels[0].params.workGroups.global[0] < (input_size / quantize_grp_size)) { + 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 @@ -651,7 +656,7 @@ KernelsData FullyConnected_bf_tiled::GetTunedKernelsDataByIndex(const Params &pa } KernelsData kernels_data; - if (is_dynamic_quantize(fc_params)) { + 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. diff --git a/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp b/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp index fc9c1601470..62691906847 100644 --- a/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp +++ b/src/plugins/intel_gpu/src/runtime/debug_configuration.cpp @@ -180,6 +180,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), " @@ -243,7 +244,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); @@ -293,6 +295,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/tests/unit/test_cases/fully_connected_gpu_test.cpp b/src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp index e44aca7dd28..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 @@ -1361,7 +1361,7 @@ public: count++; OPENVINO_ASSERT(abs_diff < 256); } - std::cout << "---> count: " << count << ", max_diff:" << max_diff << ", avg_diff: " << (avg/count) << std::endl; + GPU_DEBUG_LOG << "---> count: " << count << ", max_diff:" << max_diff << ", avg_diff: " << (avg/count) << std::endl; } @@ -3276,15 +3276,31 @@ 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); From 6a4134ab61dbaa735a9dd329ec7cf6b668d02fb9 Mon Sep 17 00:00:00 2001 From: "Min, Byung-il" Date: Wed, 12 Jun 2024 08:23:56 +0900 Subject: [PATCH 16/16] [GPU] resovle CI issues Signed-off-by: Min, Byung-il --- .../cl_kernels/include/batch_headers/int4_utils.cl | 4 ++-- .../fully_connected/fully_connected_kernel_bf_tiled.cpp | 5 ++++- 2 files changed, 6 insertions(+), 3 deletions(-) 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 caf656e6fe3..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 @@ -207,5 +207,5 @@ inline uchar8 unpack_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadabl #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_MIXED_INT4x2(target_type, value) CAT(unpack_transposed_to_, target_type)(value) -#define UNPACK_MIXED_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_transposed_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 f63820ffb99..7930f4220f0 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 @@ -26,7 +26,10 @@ static std::pair get_input_bf_size(const fully_connected_params& 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, uint32_t align_f = 1) { +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