natas_hw
/
openvino
0
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge 104e150dbf into 43de73cbd6

This commit is contained in:
Min, Byungil 2024-06-12 14:40:31 +09:00 committed by GitHub
parent 43de73cbd6 104e150dbf
commit 140ec7acad
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
12 changed files with 1034 additions and 108 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
src/plugins/intel_gpu/include/intel_gpu/runtime/debug_configuration.hpp
View File

@ -140,6 +140,7 @@ public:
int disable_runtime_skip_reorder; // Disable runtime skip reorder
int disable_primitive_fusing; // Disable primitive fusing
int disable_fake_alignment; // Disable fake alignment
int enable_dynamic_quantize; // Enable Dynamic quantization for fully connected primitive
std::set<int64_t> dump_iteration; // Dump n-th execution of network.
std::vector<std::string> load_layers_raw_dump; // List of layers to load dumped raw binary and filenames
static const debug_configuration *get_instance();

2
src/plugins/intel_gpu/src/graph/impls/ocl/fully_connected.cpp
View File

@ -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;
}

510
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/fully_connected_gpu_bf_tiled.cl
View File

@ -17,6 +17,39 @@
// DISPATCH_FSV - output coordinates for each sub-group are calculated from linearized coordinates
// DISPATCH_BSV as if they laid in bs_fs_bsv_fsv format, these macros describe fsv and bsv factors;
#define INPUT_LOAD_SIZE 4
#if FC_KERNEL_DYNAMIC_QUANTIZE
KERNEL(quantize_input)(
const __global INPUT0_TYPE* input,
__global char* quantized_input,
__global INPUT0_TYPE* de_quan_scale) {
const uint offset = get_global_id(0);
uint input_offset = offset * QUANTIZE_GROUP_SIZE;
half4 input_0[8];
char4 quantized_value[8];
half max[8];
unroll_for (uint i = 0 ; i < 8 ; ++i) {
input_0[i] = vload4(0, &input[input_offset + i * 4]);
max[i] = fmax(fmax(fabs(input_0[i][0]), fabs(input_0[i][1])), fmax(fabs(input_0[i][2]), fabs(input_0[i][3])));
}
half max_value = fmax(fmax(fmax(max[0], max[1]), fmax(max[2], max[3])),
fmax(fmax(max[4], max[5]), fmax(max[6], max[7])));
half quan_scale = max_value / 128;
unroll_for (uint i = 0 ; i < 8 ; ++i) {
quantized_value[i] = CAT(convert_, MAKE_VECTOR_TYPE(char, INPUT_LOAD_SIZE))(input_0[i] / (half4)quan_scale);
vstore4(quantized_value[i], 0, &quantized_input[input_offset + i * 4]);
}
de_quan_scale[offset] = quan_scale;
}
#else // !FC_KERNEL_DYNAMIC_QUANTIZE
// Verify JIT parameters.
#if SIMD != 8 && SIMD != 16
# error "fully_connected_gpu_bf_tiled.cl - SIMD must be one of {8, 16}"
@ -50,8 +83,10 @@
// Data stored in memory : f0k0k1|f16k0k1|f0k2k3|f16k2k3
// => unpack as f0k0k1|f0k2k3|f16k0k1|f16k2k3 so that the weight access order is preserved
#define UNPACK_INT4 UNPACK_INT4x2_OSV32_ISV2
#define UNPACK_TRANSPOSED_INT4 UNPACK_INT4x2_OSV32_ISV2
#else
#define UNPACK_INT4 UNPACK_INT4x2
#define UNPACK_TRANSPOSED_INT4 UNPACK_TRANSPOSED_INT4x2
#endif
// Macros for vectorized types.
#define INPUT_VEC_TYPE MAKE_VECTOR_TYPE(INPUT0_TYPE, TILE_IFM)
@ -79,6 +114,7 @@
// Check alignment restrictions for using block writes on output.
#define USE_BLOCK_WRITE ((OUTPUT_TYPE_SIZE * TILE_OUT_B_PITCH) % 16 == 0 && (OUTPUT_TYPE_SIZE * OUTPUT_OFFSET) % 16 == 0)
#if !REALIGN_FP16_OFFSET
# if OUTPUT_3D
# define MAIN_LOOP_ELEMENTS_COUNT INPUT0_SIZE_Y
@ -645,28 +681,6 @@ inline void FUNC(fc_bf_tiled_kernel_default)(
#undef WRITE_OUTPUT
} else {
output_offset += sglid;
// TODO: Investigate why below code doesn't compile and check how it affects performance.
//#define WRITE_OUTPUT_FEATURE(fi) do { \
// const bool should_write = \
// TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0 || \
// out_f + (fi) * SIMD + sglid < TILE_OUT_F_NUM; \
// if (should_write) { \
// output[output_offset] = result[out_bi][fi]; \
// } \
// output_offset += SIMD; \
// } while (false)
//
//#define WRITE_OUTPUT(bi) do { \
// const uint out_bi = bi; \
// CONST_LOOP(TILE_OFM, WRITE_OUTPUT_FEATURE); \
// output_offset += TILE_OUT_B_PITCH - TILE_OFM * SIMD; \
// } while (false)
//
//CONST_LOOP(TILE_B, WRITE_OUTPUT);
//#undef WRITE_OUTPUT
//#undef WRITE_OUTPUT_FEATURE
for (uint bi = 0; bi < TILE_B; ++bi) {
for (uint fi = 0; fi < TILE_OFM; ++fi) {
const bool should_write =
@ -686,6 +700,355 @@ inline void FUNC(fc_bf_tiled_kernel_default)(
// =====================================================================================================================================
}
// Dyc Quantize
#if USE_SLM && DYNAMIC_QUANTIZE
#define PACKED_DQ_TYPE int
#define DQ_VEC_TYPE MAKE_VECTOR_TYPE(DQ_TYPE, TILE_IFM)
#define DQ_SLM_FILTER_VEC MAKE_VECTOR_TYPE(DQ_TYPE, 4)
#define DQ_SLM_FILTER_PACKED_VEC MAKE_VECTOR_TYPE(FILTER_TYPE, FILTER_LOAD_BLOCK_SIZE)
#define DQ_SLM_FILTER_UNPACKED_VEC MAKE_VECTOR_TYPE(DQ_TYPE, FILTER_ELEMENTS_PER_LOAD)
#define DQ_FILTER_VEC_TYPE MAKE_VECTOR_TYPE(DQ_TYPE, TILE_K_OFM)
#define TO_DQ_TYPE(x) CAT(CAT(convert_, DQ_TYPE),_sat)(x)
#define TO_DQ_VEC_TYPE(x) CAT(convert_, DQ_VEC_TYPE)(x)
#define TO_DQ_SLM_FILTER_UNPACKED_VEC(x) CAT(convert_, DQ_SLM_FILTER_UNPACKED_VEC)(x)
#define TO_DQ_FILTER_VEC_TYPE(x) CAT(convert_, DQ_FILTER_VEC_TYPE)(x)
#define AS_TYPE_N_(type, n, x) as_##type##n(x)
#define AS_TYPE_N(type, n, x) AS_TYPE_N_(type, n, x)
#define AS_DQ_TYPE_4(x) AS_TYPE_N(DQ_TYPE, INPUT_LOAD_SIZE, x)
inline void FUNC(fc_bf_tiled_kernel_dyn_quan)(
OPTIONAL_SHAPE_INFO_ARG
const __global INPUT0_TYPE* input,
__global char* quantized_input,
__global INPUT0_TYPE* scale,
#if DECOMPRESSION_SCALE_TERM
const __global DECOMPRESSION_SCALE_TYPE* decompression_scale,
#endif
#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR
const __global DECOMPRESSION_ZP_TYPE* decompression_zp,
#endif
__global OUTPUT_TYPE* output,
const __global FILTER_TYPE* weights
, __local int* wei_local_mem
#if BIAS_TERM
, const __global BIAS_TYPE* biases
#endif
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_DECLS
#endif
) {
uint gid = (uint)get_group_id(0);
uint local_id = (uint)get_local_id(2);
uint sglid = (uint)get_sub_group_local_id();
// Dispatch as bs_fs_bsv_fsv, where bsv = DISPATCH_BSV and fsv = DISPATCH_FSV.
// This allows more fine grained control over dispatch order than using work-groups and
// avoids requirement of threads being available for whole work-group.
// It could hovewer have some drawbacks like not providing physical locality or not using
// full dispatch pipeline.
uint feature_mini_block = gid % DISPATCH_FSV;
uint batch_mini_block = gid / DISPATCH_FSV % DISPATCH_BSV;
uint feature_mega_block = gid / (DISPATCH_FSV * DISPATCH_BSV) % (CEIL_DIV(TILE_OUT_F_NUM, TILE_OFM * SIMD) / DISPATCH_FSV);
uint batch_mega_block = gid / (DISPATCH_FSV * DISPATCH_BSV * CEIL_DIV(TILE_OUT_F_NUM, TILE_OFM * SIMD) / DISPATCH_FSV);
FILTER_VEC_TYPE wei = 0;
uint out_f = gid * (TILE_OFM * SIMD);
uint out_b = LWS_BATCHES * TILE_B * (uint)get_group_id(2) + local_id * TILE_B;
#if OUTPUT_3D
uint out_b0 = out_b / OUTPUT_FEATURE_NUM;
uint out_b1 = out_b % OUTPUT_FEATURE_NUM;
uint input_offset = out_b0 * INPUT0_BATCH_PITCH + out_b1 * INPUT0_FEATURE_PITCH + INPUT0_OFFSET;
#else
uint input_offset = out_b * TILE_IN_B_PITCH + INPUT0_OFFSET;
#endif
uint weights_offset = out_f * (INPUT_ELEMENTS_COUNT / 2);
ACCUMULATOR_VEC_TYPE acc[TILE_B] = { };
// Dynamic Quantize
MAKE_VECTOR_TYPE(DQ_TYPE, INPUT_LOAD_SIZE) tiled_input_0[HALF_TILE_B] = { }; // Load 4 linear inputs for packing
PACKED_DQ_TYPE packed_in_0[HALF_TILE_B] = { }; // Packing char4 inputs to 1 integer
INPUT0_TYPE de_quantize_scale[TILE_B];
#if COMPRESSED_WEIGHTS && DECOMPRESSION_SCALE_GROUPS_NUM == 1
#if DECOMPRESSION_SCALE_LENGTH > 1 && DECOMPRESSION_SCALE_LENGTH % (TILE_OFM * SIMD) == 0
ACCUMULATOR_VEC_TYPE d_scale = TO_ACCUMULATOR_VEC_TYPE(BLOCK_READN(DECOMPRESSION_SCALE_TYPE, TILE_OFM, decompression_scale, out_f));
#elif DECOMPRESSION_SCALE_LENGTH > 1 && DECOMPRESSION_SCALE_LENGTH % (TILE_OFM * SIMD) != 0
ACCUMULATOR_VEC_TYPE d_scale = 0;
unroll_for(uint of = 0; of < TILE_OFM; ++of) {
uint offset = out_f + of*SIMD + get_sub_group_local_id();
if (offset < DECOMPRESSION_SCALE_LENGTH)
((ACCUMULATOR_TYPE*)(&d_scale))[of] = decompression_scale[offset];
}
#else
ACCUMULATOR_VEC_TYPE d_scale = decompression_scale[0];
#endif
ACCUMULATOR_TYPE* d_scales = (ACCUMULATOR_TYPE*)(&d_scale);
#endif
#if COMPRESSED_WEIGHTS && DECOMPRESSION_ZP_TERM && DECOMPRESSION_ZP_GROUPS_NUM == 1 && !DECOMPRESSION_ZP_SCALAR
#if DECOMPRESSION_ZP_LENGTH > 1 && DECOMPRESSION_ZP_LENGTH % (TILE_OFM * SIMD) == 0
ACCUMULATOR_VEC_TYPE d_zp = TO_ACCUMULATOR_VEC_TYPE(BLOCK_READN(DECOMPRESSION_ZP_TYPE, TILE_OFM, decompression_zp, out_f));
#elif DECOMPRESSION_ZP_LENGTH > 1 && DECOMPRESSION_ZP_LENGTH % (TILE_OFM * SIMD) != 0
ACCUMULATOR_VEC_TYPE d_zp = 0;
unroll_for(uint of = 0; of < TILE_OFM; ++of) {
uint offset = out_f + of*SIMD + get_sub_group_local_id();
if (offset < DECOMPRESSION_ZP_LENGTH)
((ACCUMULATOR_TYPE*)(&d_zp))[of] = decompression_zp[offset];
}
#else
ACCUMULATOR_VEC_TYPE d_zp = decompression_zp[0];
#endif
ACCUMULATOR_TYPE* d_zps = (ACCUMULATOR_TYPE*)(&d_zp);
#endif
// =====================================================================================================================================
// Main computation loop
const uint iterations = MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD);
// Each sub-group loads 2 Batch
uint idx_sglid = (sglid * TILE_K) % QUANTIZE_GROUP_SIZE; // same index for sglid 0~7 : to tile_k direction
uint batch_sglid = (sglid * TILE_K) / QUANTIZE_GROUP_SIZE; // 0 to 1 : to batch direction
__attribute__((opencl_unroll_hint(1)))
for (uint ni = 0; ni < iterations; ++ni) {
uint in_offset = input_offset + (idx_sglid + batch_sglid * TILE_IN_B_PITCH);
uint scale_offset = input_offset / QUANTIZE_GROUP_SIZE;
for (uint bi = 0; bi < HALF_TILE_B; ++bi) {
// Load quantizing info from pre-quantizing kernel
tiled_input_0[bi] = vload4(0, &quantized_input[in_offset]);
de_quantize_scale[bi * 2] = scale[scale_offset];
de_quantize_scale[bi * 2 + 1] = scale[scale_offset+ (TILE_IN_B_PITCH/QUANTIZE_GROUP_SIZE)];
// Packing : Get 4(B)x4(K) integer vector (packing to 4x1 vector)
packed_in_0[bi] = as_int(tiled_input_0[bi]);
// Next batch
in_offset += (TILE_IN_B_PITCH * 2);
scale_offset += (TILE_IN_B_PITCH/QUANTIZE_GROUP_SIZE * 2);
}
input_offset += TILE_IFM * SIMD;
// Packing
MAKE_VECTOR_TYPE(int, TILE_B) acc_tmp[TILE_OFM] = { };
#if TILE_OFM != 2
#error "FC bf_tiled kernel: can't use SLM optimization with TILE_OFM != 2"
#endif
// Skip first barrier synchronization if there is only single outer loop iteration.
#if MAIN_LOOP_ELEMENTS_COUNT / (TILE_IFM * SIMD) > 1
barrier(CLK_LOCAL_MEM_FENCE);
#endif
__local int* char_slm_weight = (__local int*)wei_local_mem;
uint weights_idx = weights_offset + local_id * SIMD * FILTER_LOAD_ITERS * FILTER_LOAD_BLOCK_SIZE;
uint wei_local_idx = local_id * SIMD * FILTER_LOAD_ITERS * (FILTER_LOAD_BLOCK_SIZE/2) + sglid * 2;
// DECOMPRESSION_SCALE_POST_OP SHOULD be enabled for dynamic quantize FC : scale is ACCUMULATOR_VAL_ONE
unroll_for(uint load_iter = 0; load_iter < FILTER_LOAD_ITERS; ++load_iter) {
SLM_FILTER_PACKED_VEC wei_packed = BLOCK_READN(FILTER_TYPE, FILTER_LOAD_BLOCK_SIZE, weights, weights_idx);
DQ_SLM_FILTER_UNPACKED_VEC dq_wei_unpacked = UNPACK_TRANSPOSED_INT4(DQ_TYPE, *((uint4x8_t *)&wei_packed));
// Calculate zero-point and scale only for DECOMPRESSION_SCALE_POST_OP enabled
#if DECOMPRESSION_ZP_TERM
#if DECOMPRESSION_ZP_SCALAR
DQ_SLM_FILTER_UNPACKED_VEC dzp = (DQ_SLM_FILTER_UNPACKED_VEC)(DECOMPRESSION_ZP_VALUE);
#elif DECOMPRESSION_ZP_GROUPS_NUM > 1
DQ_SLM_FILTER_UNPACKED_VEC dzp;
unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) {
unroll_for(uint kii = 0; kii < FILTER_LOAD_BLOCK_SIZE; ++kii) {
const uint offset_ofm = out_f + fi*SIMD + sglid;
const uint offset_ifm = ni * TILE_IFM * SIMD + local_id * FILTER_LOAD_ITERS * FILTER_LOAD_BLOCK_SIZE + load_iter * FILTER_LOAD_BLOCK_SIZE + kii;
const uint zp_offset = (offset_ofm % DECOMPRESSION_ZP_BATCH_NUM) * DECOMPRESSION_ZP_BATCH_PITCH +
(offset_ifm / DECOMPRESSION_ZP_GROUP_SIZE) * DECOMPRESSION_ZP_FEATURE_PITCH;
dzp[W_IDX] = decompression_zp[zp_offset];
}
}
#else
DQ_SLM_FILTER_UNPACKED_VEC dzp = (DQ_SLM_FILTER_UNPACKED_VEC)(d_zps[0]);
#endif
#else
DQ_SLM_FILTER_UNPACKED_VEC dzp = (DQ_SLM_FILTER_UNPACKED_VEC)(ACCUMULATOR_VAL_ZERO);
#endif
// Calculate weight : w = (w - dzp) * ds
dq_wei_unpacked -= dzp;
#if FILTER_LOAD_BLOCK_SIZE == 2
DQ_SLM_FILTER_VEC wei_1 = {dq_wei_unpacked.s01, dq_wei_unpacked.s23};
char_slm_weight[wei_local_idx] = as_int(wei_1);
#elif FILTER_LOAD_BLOCK_SIZE == 4
DQ_SLM_FILTER_VEC wei_1 = {dq_wei_unpacked.s01, dq_wei_unpacked.s23};
char_slm_weight[wei_local_idx] = as_int(wei_1);
DQ_SLM_FILTER_VEC wei_2 = {dq_wei_unpacked.s45, dq_wei_unpacked.s67};
char_slm_weight[wei_local_idx+1] = as_int(wei_2);
#elif FILTER_LOAD_BLOCK_SIZE == 8
DQ_SLM_FILTER_VEC wei_1 = {dq_wei_unpacked.s01, dq_wei_unpacked.s23};
char_slm_weight[wei_local_idx] = as_int(wei_1);
DQ_SLM_FILTER_VEC wei_2 = {dq_wei_unpacked.s45, dq_wei_unpacked.s67};
char_slm_weight[wei_local_idx+1] = as_int(wei_2);
DQ_SLM_FILTER_VEC wei_3 = {dq_wei_unpacked.s89, dq_wei_unpacked.sab};
char_slm_weight[wei_local_idx+2] = as_int(wei_3);
DQ_SLM_FILTER_VEC wei_4 = {dq_wei_unpacked.scd, dq_wei_unpacked.sef};
char_slm_weight[wei_local_idx+3] = as_int(wei_4);
#else
#error "FC bf_tiled kernel: unsupported FILTER_LOAD_BLOCK_SIZE for SLM kernel"
#endif
wei_local_idx += SIMD * (FILTER_LOAD_BLOCK_SIZE/2);
weights_idx += SIMD * FILTER_LOAD_BLOCK_SIZE;
}
wei_local_idx = sglid * 2;
barrier(CLK_LOCAL_MEM_FENCE);
unroll_for(uint ki = 0; ki < (TILE_IFM * SIMD) / TILE_K; ++ki) {
#if TILE_K != 4
#error "FC bf_tiled kernel: unsupported TILE_K size for SLM kernel"
#endif
// Compute input * weight : packed char4 type
char8 weight = vload8(0, (__local char *)(&char_slm_weight[wei_local_idx + 16*2*ki]));
char4 first_weight = weight.s0123;
char4 second_weight = weight.s4567;
unroll_for (uint bi = 0; bi < TILE_B; ++bi) {
char4 input_val = as_char4(_sub_group_shuffle(packed_in_0[bi / 2], (bi % 2) * 8 + ki));
acc_tmp[0][bi] = imad_SW(acc_tmp[0][bi], input_val, first_weight);
acc_tmp[1][bi] = imad_SW(acc_tmp[1][bi], input_val, second_weight);
}
weights_offset += TILE_K_OFM_PACKED * SIMD;
#if DECOMPRESSION_SCALE_POST_OP && (TILE_IFM * SIMD > DECOMPRESSION_SCALE_GROUP_SIZE)
unroll_for (uint bi = 0; bi < TILE_B; ++bi) {
unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) {
const uint offset_ofm = out_f + fi*SIMD + sglid;
#if DECOMPRESSION_SCALE_GROUPS_NUM > 1
const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH +
((ni*TILE_IFM*SIMD + ki*TILE_K) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH;
ACCUMULATOR_TYPE ds = decompression_scale[scale_offset];
#else
ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH];
#endif
((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi];
acc_tmp[fi][bi] = 0;
}
}
#endif
} // Whole tile_k elements of each iteration : ki
#if DECOMPRESSION_SCALE_POST_OP && (TILE_IFM * SIMD <= DECOMPRESSION_SCALE_GROUP_SIZE)
const uint ni_offset = ((ni*TILE_IFM*SIMD) / DECOMPRESSION_SCALE_GROUP_SIZE)*DECOMPRESSION_SCALE_FEATURE_PITCH;
unroll_for (uint bi = 0; bi < TILE_B; ++bi) {
unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) {
const uint offset_ofm = out_f + fi*SIMD + sglid;
#if DECOMPRESSION_SCALE_GROUPS_NUM > 1
const uint scale_offset = (offset_ofm % DECOMPRESSION_SCALE_BATCH_NUM) * DECOMPRESSION_SCALE_BATCH_PITCH + ni_offset;
ACCUMULATOR_TYPE ds = decompression_scale[scale_offset];
#else
ACCUMULATOR_TYPE ds = d_scales[fi % DECOMPRESSION_SCALE_LENGTH];
#endif
((ACCUMULATOR_TYPE*)(&acc[bi]))[fi] += convert_half(((int *)(&acc_tmp[fi]))[bi]) * ds * de_quantize_scale[bi];
}
}
#endif
} // Main compute loop : ni
// =====================================================================================================================================
// Post-processing: bias, activation, fused-ops
ACTIVATION_VEC_TYPE activated[TILE_B] = { };
for (uint bi = 0; bi < TILE_B; ++bi) {
activated[bi] = TO_ACTIVATION_VEC_TYPE(acc[bi]);
}
#if BIAS_TERM
#if TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0
BIAS_VEC_TYPE bias = BIAS_BLOCK_READ(biases, out_f);
#else
BIAS_VEC_TYPE bias = 0;
unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) {
((BIAS_TYPE*)(&bias))[fi] = biases[out_f + sglid + fi * SIMD];
}
#endif
unroll_for (uint bi = 0; bi < TILE_B; ++bi) {
activated[bi] += TO_ACTIVATION_VEC_TYPE(bias);
}
#endif
OUTPUT_VEC_TYPE result[TILE_B] = { };
#if HAS_FUSED_OPS
unroll_for (uint bi = 0; bi < TILE_B; ++bi) {
#if TILE_OFM > 1
unroll_for(uint fi = 0; fi < TILE_OFM; ++fi) {
FUSED_OPS_VEC;
result[bi][fi] = FUSED_OPS_RESULT_VEC;
}
#else
FUSED_OPS_SCALAR;
result[bi] = FUSED_OPS_RESULT_SCALAR;
#endif // TILE_OFM > 1
}
#else
unroll_for (uint bi = 0; bi < TILE_B; ++bi) {
result[bi] = TO_OUTPUT_VEC_TYPE(ACTIVATION_TYPED(activated[bi], ACTIVATION_PARAMS_TYPED));
}
#endif
// =====================================================================================================================================
// Write results
uint output_offset = out_f * TILE_OUT_F_PITCH + out_b * TILE_OUT_B_PITCH + OUTPUT_OFFSET;
if (USE_BLOCK_WRITE && (TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0 || out_f + (TILE_OFM * SIMD) <= TILE_OUT_F_NUM)) {
#if IS_DYNAMIC
#define WRITE_OUTPUT(bi) do { \
if (bi + out_b < BATCH_SIZE) \
OUTPUT_BLOCK_WRITE(output, output_offset, result[bi]); \
output_offset += TILE_OUT_B_PITCH; \
} while (false)
#else
#define WRITE_OUTPUT(bi) do { \
OUTPUT_BLOCK_WRITE(output, output_offset, result[bi]); \
output_offset += TILE_OUT_B_PITCH; \
} while (false)
#endif
CONST_LOOP(TILE_B, WRITE_OUTPUT);
#undef WRITE_OUTPUT
} else {
output_offset += sglid;
for (uint bi = 0; bi < TILE_B; ++bi) {
for (uint fi = 0; fi < TILE_OFM; ++fi) {
const bool should_write =
#if IS_DYNAMIC
bi + out_b < BATCH_SIZE &&
#endif
(TILE_OUT_F_NUM % (TILE_OFM * SIMD) == 0 ||
out_f + fi * SIMD + sglid < TILE_OUT_F_NUM);
if (should_write) {
output[output_offset] = ((OUTPUT_TYPE*)(&result[bi]))[fi];
}
output_offset += SIMD;
}
output_offset += TILE_OUT_B_PITCH - TILE_OFM * SIMD;
}
}
// =====================================================================================================================================
}
#endif
REQD_SUB_GROUP_SIZE(SIMD)
KERNEL(fc)(
OPTIONAL_SHAPE_INFO_ARG
@ -704,9 +1067,17 @@ KERNEL(fc)(
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_DECLS
#endif
#if DYNAMIC_QUANTIZE
, __global char* quantized_input
, __global INPUT0_TYPE* de_quan_scale
#endif
) {
#if USE_SLM
__local ACCUMULATOR_TYPE wei_local_mem[TILE_IFM * SIMD * TILE_OFM * SIMD];
#if DYNAMIC_QUANTIZE
__local int dq_wei_local_mem[SIMD * TILE_OFM * SIMD];
#else
__local ACCUMULATOR_TYPE wei_local_mem[TILE_IFM * SIMD * TILE_OFM * SIMD];
#endif
#endif
#if IS_DYNAMIC && COMPRESSED_WEIGHTS_INT4
const int batch_size = BATCH_SIZE;
@ -844,6 +1215,76 @@ KERNEL(fc)(
#endif
);
} else {
#if USE_SLM && DYNAMIC_QUANTIZE
FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)(
OPTIONAL_SHAPE_INFO_TENSOR
input,
quantized_input,
de_quan_scale,
#if DECOMPRESSION_SCALE_TERM
decompression_scale,
#endif
#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR
decompression_zp,
#endif
output,
weights
, dq_wei_local_mem
#if BIAS_TERM
, biases
#endif
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_ARGS
#endif
);
#else
FUNC_CALL(fc_bf_tiled_kernel_default)(
OPTIONAL_SHAPE_INFO_TENSOR
input,
#if DECOMPRESSION_SCALE_TERM
decompression_scale,
#endif
#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR
decompression_zp,
#endif
output,
weights
#if USE_SLM
, wei_local_mem
#endif
#if BIAS_TERM
, biases
#endif
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_ARGS
#endif
);
#endif
}
#else
#if USE_SLM && DYNAMIC_QUANTIZE
FUNC_CALL(fc_bf_tiled_kernel_dyn_quan)(
OPTIONAL_SHAPE_INFO_TENSOR
input,
quantized_input,
de_quan_scale,
#if DECOMPRESSION_SCALE_TERM
decompression_scale,
#endif
#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR
decompression_zp,
#endif
output,
weights
, dq_wei_local_mem
#if BIAS_TERM
, biases
#endif
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_ARGS
#endif
);
#else
FUNC_CALL(fc_bf_tiled_kernel_default)(
OPTIONAL_SHAPE_INFO_TENSOR
input,
@ -865,31 +1306,10 @@ KERNEL(fc)(
, FUSED_OPS_ARGS
#endif
);
}
#else
FUNC_CALL(fc_bf_tiled_kernel_default)(
OPTIONAL_SHAPE_INFO_TENSOR
input,
#if DECOMPRESSION_SCALE_TERM
decompression_scale,
#endif
#if DECOMPRESSION_ZP_TERM && !DECOMPRESSION_ZP_SCALAR
decompression_zp,
#endif
output,
weights
#if USE_SLM
, wei_local_mem
#endif
#if BIAS_TERM
, biases
#endif
#if HAS_FUSED_OPS_DECLS
, FUSED_OPS_ARGS
#endif
);
#endif
}
#endif // !FC_KERNEL_DYNAMIC_QUANTIZE
#undef INPUT_VEC_TYPE
#undef ACCUMULATOR_VEC_TYPE

89
src/plugins/intel_gpu/src/kernel_selector/cl_kernels/include/batch_headers/int4_utils.cl
View File

@ -20,6 +20,12 @@ inline uchar2 cvt_uint4x2_to_uint8x2(uint4x2_t v) __attribute__((overloadable))
return (uchar2)(v0, v1);
}
inline char2 cvt_uint4x2_to_int8x2(uint4x2_t v) __attribute__((overloadable)) {
const char v0 = convert_char(v.s0 & 0x0F);
const char v1 = convert_char((v.s0 & 0xF0) >> 4);
return (char2)(v0, v1);
}
inline char2 cvt_int4x2_to_int8x2(int4x2_t v) __attribute__((overloadable)) {
const char s_bit = (v.s0 & convert_char(0x08));
const char mask = s_bit > 0 ? convert_char(0xF0) : convert_char(0x00);
@ -28,6 +34,68 @@ inline char2 cvt_int4x2_to_int8x2(int4x2_t v) __attribute__((overloadable)) {
return (char2)(v0, v1);
}
inline uchar2 unpack_to_uchar(uint4x2_t v) __attribute__((overloadable)) {
return cvt_uint4x2_to_uint8x2(v);
}
inline uchar8 unpack_to_uchar(uint4x8_t v) __attribute__((overloadable)) {
uchar2 v0 = unpack_to_uchar(v.s0);
uchar2 v1 = unpack_to_uchar(v.s1);
uchar2 v2 = unpack_to_uchar(v.s2);
uchar2 v3 = unpack_to_uchar(v.s3);
return (uchar8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1);
}
inline char2 unpack_to_char(uint4x2_t v) __attribute__((overloadable)) {
return cvt_uint4x2_to_int8x2(v);
}
inline char4 unpack_to_char(uint4x4_t v) __attribute__((overloadable)) {
char2 v0 = unpack_to_char(v.s0);
char2 v1 = unpack_to_char(v.s1);
return (char4)(v0.s0, v0.s1, v1.s0, v1.s1);
}
inline char8 unpack_to_char(uint4x8_t v) __attribute__((overloadable)) {
char2 v0 = unpack_to_char(v.s0);
char2 v1 = unpack_to_char(v.s1);
char2 v2 = unpack_to_char(v.s2);
char2 v3 = unpack_to_char(v.s3);
return (char8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1);
}
inline char8 unpack_transposed_to_char(uint4x8_t v) __attribute__((overloadable)) {
char2 v0 = unpack_to_char(v.s0);
char2 v1 = unpack_to_char(v.s1);
char2 v2 = unpack_to_char(v.s2);
char2 v3 = unpack_to_char(v.s3);
return (char8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1);
}
inline uchar8 unpack_transposed_to_uchar(uint4x8_t v) __attribute__((overloadable)) {
uchar2 v0 = unpack_to_uchar(v.s0);
uchar2 v1 = unpack_to_uchar(v.s1);
uchar2 v2 = unpack_to_uchar(v.s2);
uchar2 v3 = unpack_to_uchar(v.s3);
return (uchar8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1);
}
inline char8 unpack_transposed_to_char_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) {
char2 v0 = unpack_to_char(v.s0);
char2 v1 = unpack_to_char(v.s2);
char2 v2 = unpack_to_char(v.s1);
char2 v3 = unpack_to_char(v.s3);
return (char8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1);
}
inline uchar8 unpack_transposed_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) {
uchar2 v0 = unpack_to_uchar(v.s0);
uchar2 v1 = unpack_to_uchar(v.s2);
uchar2 v2 = unpack_to_uchar(v.s1);
uchar2 v3 = unpack_to_uchar(v.s3);
return (uchar8)(v0.s0, v1.s0, v2.s0, v3.s0, v0.s1, v1.s1, v2.s1, v3.s1);
}
inline float2 unpack_to_float(uint4x2_t v) __attribute__((overloadable)) {
return convert_float2(cvt_uint4x2_to_uint8x2(v));
}
@ -116,7 +184,28 @@ inline half8 unpack_to_half_osv32_isv2(int4x8_t v) __attribute__((overloadable))
half2 f3 = unpack_to_half(v.s3);
return (half8)(f0.s0, f0.s1, f1.s0, f1.s1, f2.s0, f2.s1, f3.s0, f3.s1);
}
inline char8 unpack_to_char_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) {
char2 v0 = unpack_to_char(v.s0);
char2 v1 = unpack_to_char(v.s2);
char2 v2 = unpack_to_char(v.s1);
char2 v3 = unpack_to_char(v.s3);
return (char8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1);
}
inline uchar8 unpack_to_uchar_osv32_isv2(uint4x8_t v) __attribute__((overloadable)) {
uchar2 v0 = unpack_to_uchar(v.s0);
uchar2 v1 = unpack_to_uchar(v.s2);
uchar2 v2 = unpack_to_uchar(v.s1);
uchar2 v3 = unpack_to_uchar(v.s3);
return (uchar8)(v0.s0, v0.s1, v1.s0, v1.s1, v2.s0, v2.s1, v3.s0, v3.s1);
}
#endif // defined(cl_khr_fp16)
#define UNPACK_INT4x2(target_type, value) CAT(unpack_to_, target_type)(value)
#define UNPACK_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_to_, target_type), _osv32_isv2)(value)
#define UNPACK_TRANSPOSED_INT4x2(target_type, value) CAT(unpack_transposed_to_, target_type)(value)
#define UNPACK_TRANSPOSED_INT4x2_OSV32_ISV2(target_type, value) CAT(CAT(unpack_transposed_to_, target_type), _osv32_isv2)(value)

379
src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.cpp
View File

@ -3,15 +3,79 @@
//
#include "fully_connected_kernel_bf_tiled.h"
#include "kernel_selector_utils.h"
#include <vector>
#include <functional>
#include "common_types.h"
static constexpr size_t simd = 16;
static constexpr size_t quantize_grp_size = 32;
static constexpr size_t min_slm_size = 256;
namespace kernel_selector {
static std::pair<size_t, size_t> get_input_bf_size(const fully_connected_params& params) {
size_t input_f = params.inputs[0].Feature().v;
size_t input_batch = params.inputs[0].Batch().v;
// 3D input
if (params.outputs[0].GetLayout() == DataLayout::bfyx) {
input_f = params.inputs[0].Y().v;
input_batch = params.inputs[0].Batch().v * params.inputs[0].Feature().v;
}
return {input_batch, input_f};
}
static std::pair<size_t, size_t> get_output_aligned_bf_size(const fully_connected_params& params,
bool needs_align,
uint32_t align_b = 1,
int32_t align_f = 1) {
size_t output_f = (needs_align == true) ? CeilDiv(params.outputs[0].Feature().v, align_f) : params.outputs[0].Feature().v;
size_t output_b = params.outputs[0].Batch().v;
// 3D output
if (params.outputs[0].GetLayout() == DataLayout::bfyx) {
output_f = (needs_align == true) ? CeilDiv(params.outputs[0].Y().v, align_f) : params.outputs[0].Y().v;
output_b = params.outputs[0].Batch().v * params.outputs[0].Feature().v;
}
output_b = (needs_align == true) ? CeilDiv(output_b, align_b) : output_b;
return {output_b, output_f};
}
// DYNAMIC_QUANTIZE
static bool should_dynamic_quantize(const fully_connected_params& params) {
auto dynamic_quantization_group_size = params.dynamic_quantization_group_size;
GPU_DEBUG_GET_INSTANCE(debug_config);
GPU_DEBUG_IF(debug_config->enable_dynamic_quantize) {
dynamic_quantization_group_size = quantize_grp_size;
}
if (params.inputs[0].GetFirstElementOffset() != 0)
return false;
if (dynamic_quantization_group_size < quantize_grp_size)
return false;
auto threads = get_input_bf_size(params);
auto input_b = threads.first;
auto input_f = threads.second;
const size_t scale_group_size = params.weights.IFM().v / params.decompression_scale.Feature().v;
if ((scale_group_size % simd == 0) && (input_f % quantize_grp_size == 0) &&
(params.is_shape_agnostic || (params.inputs[0].Batch().v > 1 && input_b > min_slm_size)) &&
params.inputs[0].GetDType() == Datatype::F16 &&
(params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) &&
(params.decompression_zero_point.Feature().v == 1)) {
GPU_DEBUG_TRACE_DETAIL << " Dynamic quantizing for FC : scale_group_size " << scale_group_size << ", Input (" <<
kernel_selector::toString(params.inputs[0].GetDType()) << ", " << kernel_selector::toString(params.outputs[0].GetLayout()) <<
") B: " << params.inputs[0].Batch().v << ", F: " << params.inputs[0].Feature().v << ", Y: " << params.inputs[0].Y().v << std ::endl;
return true;
}
return false;
}
FullyConnected_bf_tiled::FullyConnected_bf_tiled() : FullyConnectedKernelBase("fully_connected_gpu_bf_tiled") {
for (unsigned tile_b = 1; tile_b <= 32; ++tile_b)
for (unsigned tile_ofm = 1; tile_ofm <= 4; tile_ofm *= 2)
@ -154,12 +218,9 @@ struct TuneParamsSelector {
bool TuneParamsSelector::VerifyTuneParams(const fully_connected_params& params, const tune_params& tparams) {
// Check divisibility by dispatch tile sizes.
size_t output_f = params.outputs[0].Feature().v;
size_t output_b = params.outputs[0].Batch().v;
if (params.outputs[0].GetLayout() == DataLayout::bfyx) {
output_b *= params.outputs[0].Feature().v;
output_f = params.outputs[0].Y().v;
}
auto bf_size = get_output_aligned_bf_size(params, false);
size_t output_b = bf_size.first;
size_t output_f = bf_size.second;
if (params.compressed &&
(params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4) &&
@ -182,7 +243,7 @@ bool TuneParamsSelector::VerifyTuneParams(const fully_connected_params& params,
if (tparams.kernel_type == FullyConnected_bf_tiled::KernelType::SLM) {
bool is_i4_u4 = (params.weights.GetDType() == WeightsType::INT4 || params.weights.GetDType() == WeightsType::UINT4);
const auto required_batch_alignment = 64;
if (!params.is_shape_agnostic && (!IsAligned(output_b, required_batch_alignment) || output_b < 256))
if (!params.is_shape_agnostic && (!IsAligned(output_b, required_batch_alignment) || output_b < min_slm_size))
return false;
const auto required_tile_b = 8;
@ -228,14 +289,10 @@ FullyConnected_bf_tiled::GetAutoTuneParams(const fully_connected_params& params,
&& TuneParamsSelector::VerifyTuneParams(params, auto_tune_params[idx]))
return auto_tune_params[idx];
size_t batch = params.outputs[0].Batch().v;
size_t output_f = params.outputs[0].Feature().v;
auto bf_size = get_output_aligned_bf_size(params, false);
size_t batch = bf_size.first;
size_t output_f = bf_size.second;
// 3d output
if (params.outputs[0].GetLayout() == DataLayout::bfyx) {
batch *= params.outputs[0].Feature().v;
output_f = params.outputs[0].Y().v;
}
Datatype dtype = params.inputs[0].GetDType();
auto selector = TuneParamsSelector(params);
@ -259,7 +316,7 @@ FullyConnected_bf_tiled::GetAutoTuneParams(const fully_connected_params& params,
} else {
// Try to use SLM kernels if possible
if (preferred_kernel_type != KernelType::DEFAULT) {
if (params.is_shape_agnostic) {
if (params.is_shape_agnostic && !should_dynamic_quantize(params)) {
selector.Case(tune_params(16, 2, 2, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM))
.Case(tune_params(16, 2, 1, 4, 1, 1, EXE_MODE_DEFAULT, KernelType::SLM));
}
@ -344,14 +401,9 @@ FullyConnected_bf_tiled::SetDefault(const fully_connected_params& params, int au
auto tparams = GetAutoTuneParams(params, kernel_type, autoTuneIndex);
size_t feature_threads = CeilDiv(params.outputs[0].Feature().v, tparams.tile_ofm * simd);
size_t batch_threads = params.outputs[0].Batch().v;
if (params.outputs[0].GetLayout() == DataLayout::bfyx) {
feature_threads = CeilDiv(params.outputs[0].Y().v, tparams.tile_ofm * simd);
batch_threads = params.outputs[0].Batch().v * params.outputs[0].Feature().v;
}
batch_threads = CeilDiv(batch_threads, tparams.tile_b);
auto threads = get_output_aligned_bf_size(params, true, tparams.tile_b, tparams.tile_ofm * simd);
auto batch_threads = threads.first;
auto feature_threads = threads.second;
const size_t lws_batches = 8;
const size_t aligned_batch = Align(batch_threads, lws_batches); // Each WG calculates 8x8 batches (TILE_B x LWS[2] size)
@ -380,9 +432,7 @@ FullyConnected_bf_tiled::SetDefault(const fully_connected_params& params, int au
KernelsPriority FullyConnected_bf_tiled::GetKernelsPriority(const Params& params) const {
const auto& fc_params = static_cast<const fully_connected_params&>(params);
size_t output_b = fc_params.outputs[0].Batch().v;
if (fc_params.outputs[0].GetLayout() == DataLayout::bfyx)
output_b *= fc_params.outputs[0].Feature().v;
size_t output_b = get_output_aligned_bf_size(fc_params, false).first;
float estimated_time = FORCE_PRIORITY_9;
if (output_b > 1 && fc_params.inputs[0].GetDType() == Datatype::F32)
@ -445,10 +495,23 @@ JitConstants FullyConnected_bf_tiled::GetJitConstants(const fully_connected_para
jit.AddConstant(MakeJitConstant("FILTER_LOAD_BLOCK_SIZE", block_read_size));
jit.AddConstant(MakeJitConstant("FILTER_ELEMENTS_PER_LOAD", weights_elements_per_load));
jit.Merge(make_int4_packed_type_jit_constant("INT4_PACKED_TYPE_PRELOAD", params.weights.GetDType(), weights_elements_per_load));
} else {
jit.AddConstant(MakeJitConstant("USE_SLM", 0));
}
// Validated perf gain, Dynamic quantize force enable SCALE_POST_OP for char type multiplication
if (should_dynamic_quantize(params) && dispatchData.tile_m > 1 && dispatchData.tile_n == 2) {
jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 1));
jit.AddConstant(MakeJitConstant("DECOMPRESSION_SCALE_POST_OP", 1));
jit.AddConstant(MakeJitConstant("DQ_TYPE", "char"));
jit.AddConstant(MakeJitConstant("QUANTIZE_GROUP_SIZE", quantize_grp_size));
} else {
jit.AddConstant(MakeJitConstant("DYNAMIC_QUANTIZE", 0));
}
jit.AddConstant(MakeJitConstant("SIMD", simd));
jit.AddConstant(MakeJitConstant("TILE_B", dispatchData.tile_m));
jit.AddConstant(MakeJitConstant("HALF_TILE_B", dispatchData.tile_m/2));
jit.AddConstant(MakeJitConstant("TILE_OFM", dispatchData.tile_n));
jit.AddConstant(MakeJitConstant("TILE_IFM", dispatchData.tile_mk));
jit.AddConstant(MakeJitConstant("TILE_K", dispatchData.tile_nk));
@ -523,29 +586,53 @@ void FullyConnected_bf_tiled::GetUpdateDispatchDataFunc(KernelData& kd) const {
kd.update_dispatch_data_func = [this](const Params& params, KernelData& kd) {
const auto& prim_params = static_cast<const fully_connected_params&>(params);
OPENVINO_ASSERT(kd.kernels.size() == 2, "[GPU] Invalid kernels size for update dispatch data func, expected 2, got ", kd.kernels.size());
size_t output_batch = get_output_aligned_bf_size(prim_params, false).first;
size_t output_batch = prim_params.outputs[0].Batch().v;
if (prim_params.outputs[0].GetLayout() == DataLayout::bfyx)
output_batch *= prim_params.outputs[0].Feature().v;
// Get index of the added shape-agnostic kernel
int kernel_offset = 0;
if (kd.kernels.size() == 3)
kernel_offset = 1; // quantize kernel exists
// Choose one of the two shape agnostic kernels:
// - kd.kernels[0] for batches <= 240 (default version)
// - kd.kernels[1] for batches >= 256 (slm version)
// Choose one of the two shape agnostic kernels: N == added kernel number
// - kd.kernels[N-1] for batches <= 240 (default version)
// - kd.kernels[N] for batches >= 256 (slm version)
const auto default_alignment = 16;
// We can use SLM version if `output_batch + default_alignment > 256` because memory and batch are aligned (whether 16 or 64 elements)
const auto skip_kernel_idx = output_batch + default_alignment > 256 ? 0 : 1;
const auto execute_kernel_idx = 1 - skip_kernel_idx;
// We can use SLM version if `output_batch + default_alignment > min_slm_size(256)` because memory and batch are aligned (whether 16 or 64 elements)
const auto execute_type = (output_batch + default_alignment > min_slm_size) ? KernelType::SLM : KernelType::DEFAULT;
const auto execute_kernel_idx = ((execute_type == KernelType::SLM) ? 1 : 0) + kernel_offset;
const auto skip_kernel_idx = ((execute_type == KernelType::SLM) ? 0 : 1) + kernel_offset;
// Check default or SLM version FC, and disable remain version
kd.kernels[skip_kernel_idx].skip_execution = true;
GPU_DEBUG_TRACE_DETAIL << "FC bf tiled: " << (execute_kernel_idx == 1 ? "SLM" : "Default") << " shape-agnostic kernel version "
GPU_DEBUG_TRACE_DETAIL << "FC bf tiled: " << (execute_type == KernelType::SLM ? "SLM" : "Default") << " shape-agnostic kernel version "
<< "will be used for batch size = " << output_batch << "\n";
auto dispatchData = SetDefault(prim_params, -1, execute_kernel_idx);
auto dispatchData = SetDefault(prim_params, -1, static_cast<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);
if (!kd.internalBufferSizes.empty()) {
// Pre-quantizing kernel was generated. Update the kernel and intermediate buffers or disable it.
if (execute_type == KernelType::DEFAULT) {
kd.kernels[0].skip_execution = true;
} else {
kd.kernels[0].skip_execution = false;
size_t input_f = get_input_bf_size(prim_params).second;
size_t input_size = input_f * dispatchData.tile_m * dispatchData.gws[2];
if (kd.internalBufferSizes[0] < input_size) {
kd.internalBufferSizes.clear();
kd.internalBufferSizes.push_back(input_size); // quantized input is char type
kd.internalBufferSizes.push_back(input_size / quantize_grp_size * 2); // de_quan_scale is half type
}
kd.kernels[0].params.workGroups.global = {std::max((input_size / quantize_grp_size), (size_t)1), 1, 1};
kd.kernels[0].params.workGroups.local = {16, 1, 1};
}
}
};
}
}
@ -572,34 +659,51 @@ KernelsData FullyConnected_bf_tiled::GetTunedKernelsDataByIndex(const Params &pa
weights_layout = WeightsLayout::os_iyx_osv64;
}
auto kernels_data = GetCommonKernelsData(params,
fc_params.inputs[0].GetLayout(),
weights_layout,
tparams.exec_options,
autoTuneIndex);
KernelsData kernels_data;
if (should_dynamic_quantize(fc_params)) {
// Use seperate 2 kernels for dynamic quantizing : quantizing_kernel + fc_kernel
// 1st kernel : Dynamic quantizing by quantize_grp_size
// 2nd kernel : fully connected kernel with KernelType::DEFAULT. Quantized inputs and scale values could be used.
// 3rd kernel : (optional) fully connected shape_agnostic kernel with KernelType::SLM. Quantized inputs and scale values would be used.
kernels_data = GetMultiKernelsData(params,
fc_params.inputs[0].GetLayout(),
weights_layout,
tparams.exec_options,
autoTuneIndex);
OPENVINO_ASSERT(!kernels_data.empty() && !kernels_data[0].kernels.empty(), "[GPU] Error to create multi kernel for dynamic quantizing.");
// In case of dynamic params try to configure additional optimized SLM kernel for large batches
if (params.is_shape_agnostic) {
auto tparams = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex);
auto can_select_slm_kernel = tparams.kernel_type == KernelType::SLM;
if (params.is_shape_agnostic)
GetUpdateDispatchDataFunc(kernels_data[0]);
} else {
kernels_data = GetCommonKernelsData(params,
fc_params.inputs[0].GetLayout(),
weights_layout,
tparams.exec_options,
autoTuneIndex,
0);
if (!can_select_slm_kernel)
return kernels_data;
if (params.is_shape_agnostic) {
auto tparams = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex);
auto can_select_slm_kernel = tparams.kernel_type == KernelType::SLM;
auto slm_kernel = GetCommonKernelsData(params,
fc_params.inputs[0].GetLayout(),
weights_layout,
tparams.exec_options,
autoTuneIndex,
1);
if (!can_select_slm_kernel)
return kernels_data;
if (slm_kernel.empty() || slm_kernel[0].kernels.empty())
return kernels_data;
auto slm_kernel = GetCommonKernelsData(params,
fc_params.inputs[0].GetLayout(),
weights_layout,
tparams.exec_options,
autoTuneIndex,
1);
kernels_data[0].kernels.push_back(slm_kernel[0].kernels.back());
if (slm_kernel.empty() || slm_kernel[0].kernels.empty())
return kernels_data;
// Update default update_dispatch_data_func function
GetUpdateDispatchDataFunc(kernels_data[0]);
kernels_data[0].kernels.push_back(slm_kernel[0].kernels.back());
// Update default update_dispatch_data_func function
GetUpdateDispatchDataFunc(kernels_data[0]);
}
}
return kernels_data;
@ -630,4 +734,157 @@ KernelsData FullyConnected_bf_tiled::GetKernelsData(const Params& params) const
return res;
}
KernelsData FullyConnected_bf_tiled::GetMultiKernelsData(const Params &params,
DataLayout dl,
WeightsLayout wl,
const std::string exeMode,
int autoTuneIndex) const {
if (!Validate(params)) {
return KernelsData();
}
const auto& fc_params = static_cast<const fully_connected_params&>(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<fully_connected_params>(params, 2);
fully_connected_params& new_params = *static_cast<fully_connected_params*>(kd.params.get());
if (!bProperInput) {
new_params.inputs[0] = new_params.inputs[0].TransformIgnorePadding(dl);
kd.reorderInput = true;
}
bool succeed = UpdateWeightsParams(new_params,
wl,
kd.weightsReorderParams,
GetSupportedKey());
if (!succeed) {
return {};
}
int inputs_count = 1;
if (new_params.compressed) {
inputs_count++;
if (new_params.has_decompression_zp && !new_params.scalar_zp)
inputs_count++;
}
// Generate dispatch data for KernelType::DEFAULT
int kernel_number = 0;
const DispatchData dispatchData = SetDefault(new_params, autoTuneIndex, kernel_number);
// Dynamic-quantize kernel
{
auto& quan_kernel = kd.kernels[0];
DispatchData dyn_quan_dispatch = dispatchData;
dyn_quan_dispatch.gws = {std::max((fc_params.inputs[0].PhysicalSize() / quantize_grp_size), (size_t)1), 1, 1};
dyn_quan_dispatch.lws = {16, 1, 1};
quan_kernel.params.workGroups.global = dyn_quan_dispatch.gws;
quan_kernel.params.workGroups.local = dyn_quan_dispatch.lws;
quan_kernel.skip_execution = false;
auto quan_entry_point = GetEntryPoint(kernelName, fc_params.layerID, params, kernel_number);
auto quan_cldnn_jit = GetJitConstants(new_params, dyn_quan_dispatch);
quan_cldnn_jit.AddConstant(MakeJitConstant("FC_KERNEL_DYNAMIC_QUANTIZE", 1));
auto quan_jit = CreateJit(kernelName, quan_cldnn_jit, quan_entry_point);
FillCLKernelData(quan_kernel,
dyn_quan_dispatch,
params.engineInfo,
kernelName,
quan_jit,
quan_entry_point,
exeMode, // No exec mode
false,
false,
1, // Only INPUT_0 is used for quantizing
0, // No fused ops
0, // No output
fc_params.is_shape_agnostic);
quan_kernel.params.arguments.clear(); // Clear original output argument
quan_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INPUT, 0});
quan_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0});
quan_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1});
kd.internalBufferSizes.push_back(fc_params.inputs[0].PhysicalSize());
kd.internalBufferSizes.push_back(fc_params.inputs[0].PhysicalSize() / quantize_grp_size * 2);
kernel_number++;
}
kd.internalBufferDataType = Datatype::F16;
// FC kernel for dynamic quantized input with KernelType::DEFAULT
{
auto entry_point = GetEntryPoint(kernelName, fc_params.layerID, params, kernel_number);
auto cldnn_jit = GetJitConstants(new_params, dispatchData);
auto jit = CreateJit(kernelName, cldnn_jit, entry_point);
auto& fc_kernel = kd.kernels[1];
fc_kernel.params.workGroups.global = dispatchData.gws;
fc_kernel.params.workGroups.local = dispatchData.lws;
fc_kernel.skip_execution = false;
FillCLKernelData(fc_kernel,
dispatchData,
params.engineInfo,
kernelName,
jit,
entry_point,
exeMode,
true,
!fc_params.bias.empty(),
inputs_count,
GetFusedPrimitiveInputsCount(params),
1,
fc_params.is_shape_agnostic);
fc_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0});
fc_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1});
kernel_number++;
}
const DispatchData slm_Data = SetDefault(new_params, autoTuneIndex, kernel_number);
auto slm_params = GetAutoTuneParams(fc_params, KernelType::SLM, autoTuneIndex);
auto can_select_slm_kernel = slm_params.kernel_type == KernelType::SLM;
// FC kernel for dynamic quantized input with KernelType::SLM
if (params.is_shape_agnostic && can_select_slm_kernel) {
kd.kernels.resize(kernel_number + 1);
auto entry_point = GetEntryPoint(kernelName, fc_params.layerID, params, kernel_number);
auto cldnn_jit = GetJitConstants(new_params, slm_Data);
auto jit = CreateJit(kernelName, cldnn_jit, entry_point);
auto& sa_kernel = kd.kernels[2];
sa_kernel.params.workGroups.global = slm_Data.gws;
sa_kernel.params.workGroups.local = slm_Data.lws;
sa_kernel.skip_execution = false;
FillCLKernelData(sa_kernel,
slm_Data,
params.engineInfo,
kernelName,
jit,
entry_point,
slm_params.exec_options,
true,
!fc_params.bias.empty(),
inputs_count,
GetFusedPrimitiveInputsCount(params),
1,
fc_params.is_shape_agnostic);
sa_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 0});
sa_kernel.params.arguments.push_back({ArgumentDescriptor::Types::INTERNAL_BUFFER, 1});
}
kd.autoTuneIndex = autoTuneIndex;
return {kd};
}
} // namespace kernel_selector

6
src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_kernel_bf_tiled.h
View File

@ -31,6 +31,12 @@ public:
ParamsKey GetSupportedKey() const override;
DeviceFeaturesKey get_required_device_features_key(const Params& params) const override;
KernelsData GetMultiKernelsData(const Params &params,
DataLayout dl,
WeightsLayout wl,
const std::string exeMode = EXE_MODE_DEFAULT,
int autoTuneIndex = -1) const;
struct tune_params {
tune_params(unsigned tile_b,
unsigned tile_ofm,

1
src/plugins/intel_gpu/src/kernel_selector/kernels/fully_connected/fully_connected_params.h
View File

@ -15,6 +15,7 @@ struct fully_connected_params : public weight_bias_params {
fully_connected_params() : weight_bias_params(KernelType::FULLY_CONNECTED) {}
QuantizationType quantization = QuantizationType::NONE;
size_t dynamic_quantization_group_size = 0;
ParamsKey GetParamsKey() const override {
ParamsKey k = weight_bias_params::GetParamsKey();

3
src/plugins/intel_gpu/src/plugin/compiled_model.cpp
View File

@ -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};

1
src/plugins/intel_gpu/src/plugin/plugin.cpp
View File

@ -555,6 +555,7 @@ std::vector<ov::PropertyName> 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;

5
src/plugins/intel_gpu/src/runtime/debug_configuration.cpp
View File

@ -181,6 +181,7 @@ static void print_help_messages() {
message_list.emplace_back("OV_GPU_DisableRuntimeSkipReorder", "Disable runtime skip reorder.");
message_list.emplace_back("OV_GPU_DisablePrimitiveFusing", "Disable primitive fusing");
message_list.emplace_back("OV_GPU_DisableFakeAlignment", "Disable fake alignment");
message_list.emplace_back("OV_GPU_EnableDynamicQuantize", "Enable Dynamic quantization for fully connected primitive");
message_list.emplace_back("OV_GPU_DumpIteration", "Dump n-th execution of network, separated by space.");
message_list.emplace_back("OV_GPU_MemPreallocationOptions", "Controls buffer pre-allocation feature. Expects 4 values separated by space in "
"the following order: number of iterations for pre-allocation(int), max size of single iteration in bytes(int), "
@ -245,7 +246,8 @@ debug_configuration::debug_configuration()
, disable_build_time_weight_reorder_for_dynamic_nodes(0)
, disable_runtime_skip_reorder(0)
, disable_primitive_fusing(0)
, disable_fake_alignment(0) {
, disable_fake_alignment(0)
, enable_dynamic_quantize(0) {
#ifdef GPU_DEBUG_CONFIG
get_gpu_debug_env_var("Help", help);
get_common_debug_env_var("Verbose", verbose);
@ -296,6 +298,7 @@ debug_configuration::debug_configuration()
get_gpu_debug_env_var("DisableRuntimeSkipReorder", disable_runtime_skip_reorder);
get_gpu_debug_env_var("DisablePrimitiveFusing", disable_primitive_fusing);
get_gpu_debug_env_var("DisableFakeAlignment", disable_fake_alignment);
get_gpu_debug_env_var("EnableDynamicQuantize", enable_dynamic_quantize);
std::string dump_iteration_str;
get_gpu_debug_env_var("DumpIteration", dump_iteration_str);
std::string mem_preallocation_params_str;

1
src/plugins/intel_gpu/src/runtime/execution_config.cpp
View File

@ -56,6 +56,7 @@ void ExecutionConfig::set_default() {
std::make_tuple(ov::internal::exclusive_async_requests, false),
std::make_tuple(ov::internal::query_model_ratio, 1.0f),
std::make_tuple(ov::cache_mode, ov::CacheMode::OPTIMIZE_SPEED),
std::make_tuple(ov::hint::dynamic_quantization_group_size, 0),
// Legacy API properties
std::make_tuple(ov::intel_gpu::nv12_two_inputs, false),

144
src/plugins/intel_gpu/tests/unit/test_cases/fully_connected_gpu_test.cpp
View File

@ -1255,6 +1255,116 @@ public:
}
}
void test_compressed_int4_scale_dyn_quan(bool is_caching_test, bool is_dynamic, int batch = 1) {
tests::random_generator rg(GET_SUITE_NAME);
auto& engine = get_test_engine();
if (engine.get_device_info().dev_type == device_type::discrete_gpu)
GTEST_SKIP();
long int batch_num = batch;
long int ifm_num = 1024;
long int ofm_num = 4096;
long int scales_group_size = 32;
bool is_3d = true;
auto input_ps = is_3d ? ov::PartialShape{ batch_num, 1, ifm_num } : ov::PartialShape{ batch_num, ifm_num};
auto dyn_input_ps = is_3d ? ov::PartialShape{ -1, 1, ifm_num } : ov::PartialShape{ -1, ifm_num};
auto input_mem = engine.allocate_memory({ input_ps, data_types::f16, format::bfyx });
auto weights_mem = engine.allocate_memory({ {ofm_num, ifm_num}, data_types::u4, format::bfyx });
auto scale_mem = engine.allocate_memory({ {ofm_num, ifm_num / scales_group_size}, data_types::f16, format::bfyx });
auto input_data = rg.generate_random_1d<ov::float16>(batch_num * ifm_num, -16.0f, 16.0f);
set_values(input_mem, input_data);
auto weigths_data = rg.generate_random_1d<uint8_t>(ofm_num * ifm_num / 2, 0, 10);
set_values(weights_mem, weigths_data);
auto scale_data = rg.generate_random_1d<ov::float16>(ofm_num * ifm_num / scales_group_size, -4.0f, 4.0f);
set_values(scale_mem, scale_data);
auto in_layout = is_dynamic ? layout{ dyn_input_ps, data_types::f16, format::bfyx }
: layout{ input_ps, data_types::f16, format::bfyx };
auto fc_prim = fully_connected("fc_prim", input_info("input"), "weights", "", "scale", "", data_types::f16, padding(), is_3d ? 3 : 2, 2);
fc_prim.decompression_zero_point_scalar = 0;
// Implemented dynamic quantize kernel
auto get_ref_results = [&]() {
topology topology(
input_layout("input", in_layout),
data("weights", weights_mem),
data("scale", scale_mem),
fc_prim
);
auto config = get_test_default_config(engine);
config.set_property(ov::intel_gpu::allow_new_shape_infer(true));
config.set_property(ov::intel_gpu::optimize_data(true));
network network(engine, topology, config);
network.set_input_data("input", input_mem);
auto outputs = network.execute();
OPENVINO_ASSERT(outputs.size() == 1);
OPENVINO_ASSERT(outputs.begin()->first == "fc_prim");
auto output_layout = outputs.begin()->second.get_layout();
auto output_mem = outputs.begin()->second.get_memory();
return engine.reinterpret_buffer(*output_mem, output_layout);
};
topology topology(
input_layout("input", in_layout),
data("weights", weights_mem),
data("scale", scale_mem),
fc_prim
);
auto config = get_test_default_config(engine);
config.set_property(ov::intel_gpu::allow_new_shape_infer(true));
config.set_property(ov::intel_gpu::optimize_data(true));
config.set_property(ov::hint::dynamic_quantization_group_size(32));
network::ptr network = get_network(engine, topology, config, get_test_stream_ptr(), is_caching_test);
if (is_dynamic && !engine.get_device_info().supports_immad) {
auto inst = network->get_primitive("fc_prim");
auto impl = inst->get_impl();
ASSERT_TRUE(impl != NULL);
ASSERT_EQ(impl->get_kernels().size(), size_t((is_dynamic ? 3 : 2))); // shape-agnostic kernels
}
network->set_input_data("input", input_mem);
auto outputs = network->execute();
ASSERT_EQ(outputs.size(), size_t(1));
ASSERT_EQ(outputs.begin()->first, "fc_prim");
auto output_mem = outputs.begin()->second.get_memory();
cldnn::mem_lock<ov::float16> output_ptr (output_mem, get_test_stream());
auto ref_output_mem = get_ref_results();
cldnn::mem_lock<ov::float16> output_ptr_ref (ref_output_mem, get_test_stream());
size_t count = 0;
float max_diff = 0.f;
float avg = 0.f;
for (size_t i = 0; i < output_ptr_ref.size(); ++i) {
auto abs_diff = std::abs(output_ptr_ref[i] - output_ptr[i]);
if (max_diff < abs_diff)
max_diff = abs_diff;
avg = abs_diff;
count++;
OPENVINO_ASSERT(abs_diff < 256);
}
GPU_DEBUG_LOG << "---> count: " << count << ", max_diff:" << max_diff << ", avg_diff: " << (avg/count) << std::endl;
}
void test_compressed_int4_scale(bool is_caching_test, bool is_dynamic, long int batch_num, long int scales_group_size = 128) {
tests::random_generator rg(GET_SUITE_NAME);
auto& engine = get_test_engine();
@ -3158,6 +3268,40 @@ TEST_F(fully_connected_gpu_tests, compressed_int4_scale_b1g128) {
this->test_compressed_int4_scale(false, false, 1, 128);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_single_batch) {
this->test_compressed_int4_scale_dyn_quan(false, false);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan) {
this->test_compressed_int4_scale_dyn_quan(false, false, 512);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_unaligned) {
this->test_compressed_int4_scale_dyn_quan(false, false, 511);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_dynamic_single_batch) {
this->test_compressed_int4_scale_dyn_quan(false, true, 1);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_dynamic) {
this->test_compressed_int4_scale_dyn_quan(false, true, 512);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_quan_dynamic_unaligned) {
this->test_compressed_int4_scale_dyn_quan(false, true, 511);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_cache) {
this->test_compressed_int4_scale_dyn_quan(true, false, 512);
}
TEST_F(fully_connected_gpu_tests, compressed_int4_scale_dyn_cache_dynamic) {
this->test_compressed_int4_scale_dyn_quan(true, true, 512);
}
TEST_F(fully_connected_gpu_tests, compressed_scale_bias) {
this->test_compressed_scale_bias(false);
}