From 2605faef88714b9e9d5ad0eac7f17c387db4cc1c Mon Sep 17 00:00:00 2001 From: Tianlei Wu Date: Thu, 11 Jun 2020 14:19:55 -0700 Subject: [PATCH] Add past state support in Attention Op for GPT-2 (#4107) Update Attention op to allow past state input and output. Add fusion script and tests --- onnxruntime/contrib_ops/cpu/bert/attention.cc | 89 ++++- onnxruntime/contrib_ops/cpu/bert/attention.h | 14 +- .../contrib_ops/cpu/bert/attention_helper.h | 183 ++++++--- .../cpu/quantization/attention_quant.cc | 10 +- .../contrib_ops/cuda/bert/attention.cc | 25 +- .../contrib_ops/cuda/bert/attention_impl.cu | 188 +++++++-- .../contrib_ops/cuda/bert/attention_impl.h | 13 +- .../quantization/attention_quantization.cc | 13 +- .../core/graph/contrib_ops/contrib_defs.cc | 43 ++- .../tools/transformers/BertOnnxModel.py | 22 +- .../tools/transformers/Gpt2OnnxModel.py | 142 +------ .../python/tools/transformers/benchmark.py | 5 +- .../tools/transformers/benchmark_gpt2.py | 109 +++--- .../tools/transformers/bert_perf_test.py | 9 +- .../tools/transformers/bert_test_data.py | 6 +- .../tools/transformers/dev_benchmark.cmd | 2 +- .../tools/transformers/fusion_attention.py | 7 +- .../python/tools/transformers/fusion_base.py | 17 +- .../tools/transformers/fusion_biasgelu.py | 2 +- .../tools/transformers/fusion_embedlayer.py | 12 +- .../transformers/fusion_gelu_approximation.py | 2 +- .../transformers/fusion_gpt_attention.py | 196 ++++++++++ .../fusion_gpt_attention_no_past.py | 137 +++++++ .../tools/transformers/fusion_layernorm.py | 2 +- .../transformers/fusion_skiplayernorm.py | 2 +- .../python/tools/transformers/optimizer.py | 117 ++++-- .../python/tools/transformers/pytest.ini | 9 + .../generate_tiny_gpt2_model.py | 363 ++++++++++++++++++ .../gpt2_pytorch1.5_opset11/gpt2_past.onnx | Bin 0 -> 182687 bytes .../tools/transformers/test_optimizer.py | 61 ++- .../test/contrib_ops/attention_op_test.cc | 337 +++++++++++++++- 31 files changed, 1733 insertions(+), 404 deletions(-) create mode 100644 onnxruntime/python/tools/transformers/fusion_gpt_attention.py create mode 100644 onnxruntime/python/tools/transformers/fusion_gpt_attention_no_past.py create mode 100644 onnxruntime/python/tools/transformers/pytest.ini create mode 100644 onnxruntime/python/tools/transformers/test_data/gpt2_pytorch1.5_opset11/generate_tiny_gpt2_model.py create mode 100644 onnxruntime/python/tools/transformers/test_data/gpt2_pytorch1.5_opset11/gpt2_past.onnx diff --git a/onnxruntime/contrib_ops/cpu/bert/attention.cc b/onnxruntime/contrib_ops/cpu/bert/attention.cc index d6bb58300d..e6ed0c644b 100644 --- a/onnxruntime/contrib_ops/cpu/bert/attention.cc +++ b/onnxruntime/contrib_ops/cpu/bert/attention.cc @@ -27,18 +27,21 @@ AttentionBase::AttentionBase(const OpKernelInfo& info) { int64_t num_heads = 0; ORT_ENFORCE(info.GetAttr("num_heads", &num_heads).IsOK() && num_heads > 0); num_heads_ = static_cast(num_heads); + is_unidirectional_ = info.GetAttrOrDefault("unidirectional", 0) == 1; } Status AttentionBase::CheckInputs(const Tensor* input, const Tensor* weights, const Tensor* bias, - const Tensor* mask_index) const { - // Input and output shapes: + const Tensor* mask_index, + const Tensor* past) const { + // Input shapes: // input : (batch_size, sequence_length, hidden_size) // weights : (hidden_size, 3 * hidden_size) // bias : (3 * hidden_size) // mask_index : (batch_size) if presented + // past : (2, batch_size, num_heads, past_sequence_length, head_size) const auto dims = input->Shape().GetDims(); if (dims.size() != 3) { @@ -91,9 +94,59 @@ Status AttentionBase::CheckInputs(const Tensor* input, } } + if (past != nullptr) { // past is optional + if (!is_unidirectional_) { + return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Input 4 (past) is only allowed for unidirectional"); + } + + const auto past_dims = past->Shape().GetDims(); + if (past_dims.size() != 5) { + return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Input 4 is expected to have 5 dimension, got ", + past_dims.size()); + } + if (static_cast(past_dims[0]) != 2) { + return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Inputs 4 dimension 0 shall have length of 2"); + } + if (static_cast(past_dims[1]) != batch_size) { + return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Inputs 4 dimension 1 shall have same length as dimension 0 of input 0"); + } + if (static_cast(past_dims[2]) != num_heads_) { + return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Inputs 4 dimension 2 shall have length of num_heads", num_heads_); + } + if (static_cast(past_dims[4]) != hidden_size / num_heads_) { + return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Inputs 4 dimension 2 shall have length of ", hidden_size / num_heads_); + } + } + return Status::OK(); } +Tensor* AttentionBase::GetPresent(OpKernelContext* context, + const Tensor* past, + int batch_size, + int head_size, + int sequence_length, + int& past_sequence_length) const { + // Input and output shapes: + // past : (2, batch_size, num_heads, past_sequence_length, head_size) + // present : (2, batch_size, num_heads, past_sequence_length + sequence_length, head_size) + + std::vector present_dims{2, batch_size, num_heads_, sequence_length, head_size}; + if (nullptr != past) { + const auto past_dims = past->Shape().GetDims(); + past_sequence_length = static_cast(past_dims[3]); + present_dims[3] += past_dims[3]; + } + + TensorShape present_shape(present_dims); + Tensor* present = context->Output(1, present_shape); + if (nullptr != past && nullptr == present) { + ORT_THROW("Expect to have present state output when past state input is given"); + } + + return present; +} + template Attention::Attention(const OpKernelInfo& info) : OpKernel(info), AttentionBase(info) { } @@ -104,7 +157,9 @@ Status Attention::Compute(OpKernelContext* context) const { const Tensor* weights = context->Input(1); const Tensor* bias = context->Input(2); const Tensor* mask_index = context->Input(3); - ORT_RETURN_IF_ERROR(CheckInputs(input, weights, bias, mask_index)); + const Tensor* past = context->Input(4); + + ORT_RETURN_IF_ERROR(CheckInputs(input, weights, bias, mask_index, past)); const auto dims = input->Shape().GetDims(); const int batch_size = static_cast(dims[0]); @@ -115,6 +170,12 @@ Status Attention::Compute(OpKernelContext* context) const { TensorShape output_shape(dims); Tensor* output = context->Output(0, output_shape); + int past_sequence_length = 0; + Tensor* present = GetPresent(context, past, batch_size, head_size, sequence_length, past_sequence_length); + + // Total sequence length including that of past state: S* = S' + S + const int all_sequence_length = past_sequence_length + sequence_length; + constexpr size_t element_size = sizeof(T); AllocatorPtr allocator; @@ -182,18 +243,18 @@ Status Attention::Compute(OpKernelContext* context) const { } // STEP.2: compute the attention score. It does 2 things: - // I. attention_probs(B, N, S, S) = 1/sqrt(H) x Q(B, N, S, H) x K'(B, N, S, H -> B, N, H, S) + - // 1 x mask_data(B, N, S, S) - // II.attention_probs(B, N, S, S) = Softmax(attention_probs) - size_t attention_probs_bytes = SafeInt(batch_size) * num_heads_ * sequence_length * sequence_length * element_size; + // I. attention_probs(B, N, S, S*) = 1/sqrt(H) x Q(B, N, S, H) x K'(B, N, S*, H -> B, N, H, S*) + + // 1 x mask_data(B, N, S, S*) + // II.attention_probs(B, N, S, S*) = Softmax(attention_probs) + size_t attention_probs_bytes = SafeInt(batch_size) * num_heads_ * sequence_length * all_sequence_length * element_size; auto attention_probs = allocator->Alloc(attention_probs_bytes); BufferUniquePtr scratch_buffer(attention_probs, BufferDeleter(allocator)); size_t mask_data_bytes = 0; if (mask_index != nullptr) { - mask_data_bytes = SafeInt(batch_size) * sequence_length * sequence_length * element_size; + mask_data_bytes = SafeInt(batch_size) * sequence_length * all_sequence_length * element_size; } else if (is_unidirectional_) { - mask_data_bytes = SafeInt(sequence_length) * sequence_length * element_size; + mask_data_bytes = SafeInt(sequence_length) * all_sequence_length * element_size; } void* mask_data = nullptr; @@ -204,17 +265,21 @@ Status Attention::Compute(OpKernelContext* context) const { BufferUniquePtr mask_data_buffer(mask_data, BufferDeleter(allocator)); const int32_t* mask_index_data = mask_index != nullptr ? mask_index->template Data() : nullptr; + const T* past_data = past != nullptr ? past->template Data() : nullptr; + T* present_data = present != nullptr ? present->template MutableData() : nullptr; ComputeAttentionProbs(static_cast(attention_probs), Q, K, mask_index_data, static_cast(mask_data), - batch_size, sequence_length, head_size, num_heads_, is_unidirectional_, tp); + batch_size, sequence_length, past_sequence_length, head_size, num_heads_, is_unidirectional_, + past_data, present_data, tp); - // STEP.3: compute the attentionScore * Value. It does: out_tmp(B, N, S, H) = attention_probs(B, N, S, S) x V(B, N, S, H) + // STEP.3: compute the attentionScore * Value. It does: out_tmp(B, N, S, H) = attention_probs(B, N, S, S*) x V(B, N, S*, H) auto out_tmp_data = allocator->Alloc(SafeInt(batch_size) * num_heads_ * sequence_length * head_size * element_size); BufferUniquePtr out_tmp_buffer(out_tmp_data, BufferDeleter(allocator)); ComputeVxAttentionScore(output->template MutableData(), static_cast(out_tmp_data), static_cast(attention_probs), V, - batch_size, sequence_length, head_size, num_heads_, hidden_size, tp); + batch_size, sequence_length, past_sequence_length, head_size, num_heads_, hidden_size, + past_data, present_data, tp); return Status::OK(); } diff --git a/onnxruntime/contrib_ops/cpu/bert/attention.h b/onnxruntime/contrib_ops/cpu/bert/attention.h index e82d758cdb..fa686284ac 100644 --- a/onnxruntime/contrib_ops/cpu/bert/attention.h +++ b/onnxruntime/contrib_ops/cpu/bert/attention.h @@ -15,10 +15,18 @@ class AttentionBase { Status CheckInputs(const Tensor* input, const Tensor* weights, const Tensor* bias, - const Tensor* mask_index) const; + const Tensor* mask_index, + const Tensor* past) const; - int num_heads_; // number of attention heads - bool is_unidirectional_; // whether every token can only attend to previous tokens. + Tensor* GetPresent(OpKernelContext* context, + const Tensor* past, + int batch_size, + int head_size, + int sequence_length, + int& past_sequence_length) const; + + int num_heads_; // number of attention heads + bool is_unidirectional_; // whether every token can only attend to previous tokens. }; template diff --git a/onnxruntime/contrib_ops/cpu/bert/attention_helper.h b/onnxruntime/contrib_ops/cpu/bert/attention_helper.h index bed0513ff6..4a94d0b45b 100644 --- a/onnxruntime/contrib_ops/cpu/bert/attention_helper.h +++ b/onnxruntime/contrib_ops/cpu/bert/attention_helper.h @@ -59,111 +59,176 @@ inline void ComputeAttentionSoftmaxInplace(float* score, int N, int D, ThreadPoo MlasComputeSoftmax(score, score, N, D, false, tp); } -// Helper function to compute the attention probs. It does 2 things: -// I. attention_probs(B, N, S, S) = 1/sqrt(H) x Q(B, N, S, H) x K'(B, N, S, H -> B, N, H, S) + -// 1 x mask_data(B, N, S, S) -// II.attention_probs(B, N, S, S) = Softmax(attention_probs) template -void ComputeAttentionProbs(T* attention_probs, // output buffer for the attention probs. Its size is B*N*S*S - const T* Q, // Q data. Its size is B*N*S*H - const T* K, // k data. Its size is B*N*S*H +void PrepareMask(const int32_t* mask_index, + T* mask_data, + bool is_unidirectional, + int batch_size, + int sequence_length, + int past_sequence_length) { + const int all_sequence_length = past_sequence_length + sequence_length; + T* p_mask = mask_data; + if (is_unidirectional) { + // unidirectional mask has shape SxS* + for (int s_i = 0; s_i < sequence_length - 1; s_i++) { + for (int m_i = past_sequence_length + s_i + 1; m_i < all_sequence_length; m_i++) { + p_mask[s_i * all_sequence_length + m_i] = static_cast(-10000.0); + } + } + return; + } + + ORT_ENFORCE(mask_index, "mask index should not be null."); + for (int b_i = 0; b_i < batch_size; b_i++) { + // TODO: mask_index can be used in softmax to save some calculation. + // Convert mask_index to mask (-10000 means out of range, which will be 0 after softmax): B => BxS* + int valid_length = mask_index[b_i]; + for (int m_i = valid_length; m_i < all_sequence_length; m_i++) { + p_mask[m_i] = static_cast(-10000.0); + } + + // Broadcast mask from BxS* to BxSxS* + for (int s_i = 1; s_i < sequence_length; s_i++) { + memcpy(p_mask + s_i * all_sequence_length, p_mask, all_sequence_length * sizeof(T)); + } + p_mask += sequence_length * sequence_length; + } +} + +// Concatenate a past state chunk S'xH with input state chunk SxH into present state chunk S*xH +// Returns a pointer to the start of present state chunk. +template +T* ConcatStateChunk(const T* past, const T* chunk, T* present, size_t past_chunk_length, size_t present_chunk_length, std::ptrdiff_t i) { + T* start = present + i * present_chunk_length; + + T* p = start; + if (nullptr != past) { + const T* src_past = past + i * past_chunk_length; + memcpy(p, src_past, past_chunk_length * sizeof(T)); + p += past_chunk_length; + } + + memcpy(p, chunk, (present_chunk_length - past_chunk_length) * sizeof(T)); + return start; +} + +// Helper function to compute the attention probs. It does 2 things: +// I. attention_probs(B, N, S, S*) = 1/sqrt(H) x Q(B, N, S, H) x K'(B, N, S*, H -> B, N, H, S*) + +// 1 x mask_data(B, N, S, S*) +// II.attention_probs(B, N, S, S*) = Softmax(attention_probs) +template +void ComputeAttentionProbs(T* attention_probs, // output buffer for the attention probs. Its size is BxNxSxS + const T* Q, // Q data. Its size is BxNxSxH + const T* K, // k data. Its size is BxNxSxH const int32_t* mask_index, // mask index. nullptr if no mask or its size is B - T* mask_data, // buffer for mask data. Its size is: S*S if is_unidirectiona; B*S*S if mask_index; null otherwise + T* mask_data, // buffer for mask data. Its size is: SxS* if is_unidirectional; BxSxS* if mask_index; null otherwise int batch_size, // batch size of self-attention int sequence_length, // sequence length of self-attention + int past_sequence_length, // sequence length of past state int head_size, // head size of self-attention int num_heads, // number of heads of self-attention bool is_unidirectional, // indicate if it is unidrectional. + const T* past, // past state + T* present, // present state ThreadPool* tp) { + const int all_sequence_length = past_sequence_length + sequence_length; // S* = S' + S + const size_t past_chunk_length = static_cast(past_sequence_length * head_size); // S' x H + const size_t input_chunk_length = static_cast(sequence_length * head_size); // S x H + const size_t present_chunk_length = past_chunk_length + input_chunk_length; // S* x H + { if (mask_data != nullptr) { - if (is_unidirectional) { - for (int s_i = 0; s_i < sequence_length - 1; s_i++) { - for (int m_i = s_i + 1; m_i < sequence_length; m_i++) { - mask_data[s_i * sequence_length + m_i] = static_cast(-10000.0); - } - } - } else { - ORT_ENFORCE(mask_index, "mask index should not be null."); - T* p_mask = mask_data; - for (int b_i = 0; b_i < batch_size; b_i++) { - // TODO: mask_index can be used in softmax to save some calculation. - // Convert mask_index to mask (-10000 means out of range, which will be 0 after softmax): B => BxS - int valid_length = mask_index[b_i]; - for (int m_i = valid_length; m_i < sequence_length; m_i++) { - p_mask[m_i] = static_cast(-10000.0); - } - - // Broadcast mask from BxS to BxSxS - for (int s_i = 1; s_i < sequence_length; s_i++) { - memcpy(p_mask + s_i * sequence_length, p_mask, sequence_length * sizeof(T)); - } - p_mask += sequence_length * sequence_length; - } - } + PrepareMask(mask_index, mask_data, is_unidirectional, batch_size, sequence_length, past_sequence_length); } else { // no any mask - memset(attention_probs, 0, batch_size * num_heads * sequence_length * sequence_length * sizeof(T)); + memset(attention_probs, 0, batch_size * num_heads * sequence_length * all_sequence_length * sizeof(T)); } const int loop_len = batch_size * num_heads; const float alpha = 1.0f / sqrt(static_cast(head_size)); // The cost of Gemm - const double cost = - static_cast(head_size) * static_cast(sequence_length) * static_cast(sequence_length); + const double cost = static_cast(head_size * sequence_length * all_sequence_length); + ThreadPool::TryParallelFor(tp, loop_len, cost, [&](std::ptrdiff_t begin, std::ptrdiff_t end) { for (std::ptrdiff_t i = begin; i != end; ++i) { const std::ptrdiff_t batch_index = i / num_heads; - // broadcast mask data: SxS or (Bx)SxS -> (BxNx)SxS + // broadcast mask data: SxS* or (Bx)SxS* -> (BxNx)SxS* if (mask_data != nullptr) { - const T* broadcast_data_src = is_unidirectional ? reinterpret_cast(mask_data) : reinterpret_cast(mask_data) + batch_index * sequence_length * sequence_length; - T* broadcast_data_dest = reinterpret_cast(attention_probs) + sequence_length * sequence_length * i; - memcpy(broadcast_data_dest, broadcast_data_src, sequence_length * sequence_length * sizeof(T)); + const T* broadcast_data_src = is_unidirectional ? reinterpret_cast(mask_data) : reinterpret_cast(mask_data) + batch_index * sequence_length * all_sequence_length; + T* broadcast_data_dest = reinterpret_cast(attention_probs) + sequence_length * all_sequence_length * i; + memcpy(broadcast_data_dest, broadcast_data_src, sequence_length * all_sequence_length * sizeof(T)); + } + + const T* k = K + input_chunk_length * i; + if (nullptr != present) { + // concatenate past_K and K : (BxNx)S'xH, (BxNx)SxH -> (BxNx)S*xH + k = ConcatStateChunk(past, k, present, past_chunk_length, present_chunk_length, i); } // gemm - - // original transposed iteration - // A: Q (BxNxSxH) (B.N.)S x H S x H - // B: K' (BxNxSxH) (B.N.)H x S H x S - // C: attention_probs (BxNxSxS) (B.N.)S x S S x S - - math::Gemm(CblasNoTrans, CblasTrans, sequence_length, sequence_length, head_size, alpha, - Q + sequence_length * head_size * i, K + sequence_length * head_size * i, 1.0, - reinterpret_cast(attention_probs) + sequence_length * sequence_length * i, nullptr); + // original transposed each iteration + // A: Q (B x N x) S x H (B x N x) S x H S x H + // B: K' (B x N x) S* x H (B x N x) H x S* H x S* + // C: attention_probs (B x N x) S x S* (B x N x) S x S* S x S* + math::Gemm(CblasNoTrans, CblasTrans, sequence_length, all_sequence_length, head_size, alpha, + Q + input_chunk_length * i, k, 1.0, + reinterpret_cast(attention_probs) + sequence_length * all_sequence_length * i, nullptr); } }); } - // attention_probs(B, N, S, S) = Softmax(attention_probs) + // attention_probs(B, N, S, S*) = Softmax(attention_probs) { const int N = batch_size * num_heads * sequence_length; - const int D = sequence_length; + const int D = all_sequence_length; ComputeAttentionSoftmaxInplace(attention_probs, N, D, tp); } } template -void ComputeVxAttentionScore(T* output, // buffer for the result with size B*S*N*H - T* tmp_buffer, // buffer for temp use with size is B*N*S*H - const T* attention_probs, // Attention probs with size B*N*S*S - const T* V, // V valuee with size B*N*S*H +void ComputeVxAttentionScore(T* output, // buffer for the result with size BxSxNxH + T* tmp_buffer, // buffer for temp use with size is BxNxSxH + const T* attention_probs, // Attention probs with size BxNxSxS* + const T* V, // V value with size BxNxSxH int batch_size, // batch size int sequence_length, // sequence length + int past_sequence_length, // sequence length in past state int head_size, // head size int num_heads, // number of heads int hidden_size, // hidden size + const T* past, // past state + T* present, // present state ThreadPool* tp) { + const int all_sequence_length = past_sequence_length + sequence_length; // S* = S' + S + const size_t past_chunk_length = static_cast(past_sequence_length * head_size); // S' x H + const size_t input_chunk_length = static_cast(sequence_length * head_size); // S x H + const size_t present_chunk_length = past_chunk_length + input_chunk_length; // S* x H + + // Move the pointer of past and present to start of v values. + if (nullptr != past) { + past += batch_size * num_heads * past_sequence_length * head_size; + } + if (nullptr != present) { + present += batch_size * num_heads * all_sequence_length * head_size; + } + const double cost = static_cast(sequence_length) * static_cast(head_size) * static_cast(sequence_length); + ThreadPool::TryParallelFor(tp, batch_size * num_heads, cost, [&](std::ptrdiff_t begin, std::ptrdiff_t end) { - const int sequence_length_mul_head_size = sequence_length * head_size; for (std::ptrdiff_t i = begin; i != end; ++i) { - T* current_tmp_data = tmp_buffer + sequence_length_mul_head_size * i; - math::MatMul(sequence_length, head_size, sequence_length, - attention_probs + sequence_length * sequence_length * i, - V + sequence_length_mul_head_size * i, current_tmp_data, nullptr); + + const T* v = V + input_chunk_length * i; + if (nullptr != present) { + // concatenate past_V and V: (BxNx)S'xH, (BxNx)SxH -> (BxNx)S*xH + v = ConcatStateChunk(past, v, present, past_chunk_length, present_chunk_length, i); + } + + T* current_tmp_data = reinterpret_cast(tmp_buffer) + input_chunk_length * i; + math::MatMul(sequence_length, head_size, all_sequence_length, + attention_probs + sequence_length * all_sequence_length * i, + v, current_tmp_data, nullptr); // transpose: out(B, S, N, H) = transpose out_tmp(B, N, S, H) const int batch_index = static_cast(i / num_heads); diff --git a/onnxruntime/contrib_ops/cpu/quantization/attention_quant.cc b/onnxruntime/contrib_ops/cpu/quantization/attention_quant.cc index 8f186e23a4..145276c73c 100644 --- a/onnxruntime/contrib_ops/cpu/quantization/attention_quant.cc +++ b/onnxruntime/contrib_ops/cpu/quantization/attention_quant.cc @@ -59,7 +59,7 @@ Status QAttention::Compute(OpKernelContext* context) const { const Tensor* i_zp_tensor = context->Input(6); const Tensor* w_zp_tensor = context->Input(7); - ORT_RETURN_IF_ERROR(AttentionBase::CheckInputs(input, weights, bias, mask_index)); + ORT_RETURN_IF_ERROR(AttentionBase::CheckInputs(input, weights, bias, mask_index, nullptr)); ORT_RETURN_IF_NOT(IsScalarOr1ElementVector(input_scale_tensor), "input scale must be a scalar or 1D tensor of size 1"); @@ -193,8 +193,12 @@ Status QAttention::Compute(OpKernelContext* context) const { const int32_t* mask_index_data = mask_index != nullptr ? mask_index->template Data() : nullptr; + int past_sequence_length = 0; + const T* past_data = nullptr; + T* present_data = nullptr; ComputeAttentionProbs(static_cast(attention_probs), Q, K, mask_index_data, static_cast(mask_data), - batch_size, sequence_length, head_size, num_heads_, is_unidirectional_, tp); + batch_size, sequence_length, past_sequence_length, head_size, num_heads_, is_unidirectional_, + past_data, present_data, tp); // STEP.3: compute the attentionScore * Value. It does: out_tmp(B, N, S, H) = attention_probs(B, N, S, S) x V(B, N, S, H) auto out_tmp_data = @@ -202,7 +206,7 @@ Status QAttention::Compute(OpKernelContext* context) const { BufferUniquePtr out_tmp_buffer(out_tmp_data, BufferDeleter(allocator)); ComputeVxAttentionScore(output->template MutableData(), static_cast(out_tmp_data), static_cast(attention_probs), V, - batch_size, sequence_length, head_size, num_heads_, hidden_size, tp); + batch_size, sequence_length, past_sequence_length, head_size, num_heads_, hidden_size, past_data, present_data, tp); return Status::OK(); } diff --git a/onnxruntime/contrib_ops/cuda/bert/attention.cc b/onnxruntime/contrib_ops/cuda/bert/attention.cc index 5d679fd335..b203a08345 100644 --- a/onnxruntime/contrib_ops/cuda/bert/attention.cc +++ b/onnxruntime/contrib_ops/cuda/bert/attention.cc @@ -34,18 +34,16 @@ Attention::Attention(const OpKernelInfo& info) : CudaKernel(info), AttentionB template Status Attention::ComputeInternal(OpKernelContext* context) const { - // Input and output shapes: - // Input 0 - input : (batch_size, sequence_length, hidden_size) - // Input 1 - weights : (hidden_size, 3 * hidden_size) - // Input 2 - bias : (3 * hidden_size) - // Input 3 - mask_index : (batch_size) if presented - // Output : (batch_size, sequence_length, hidden_size) const Tensor* input = context->Input(0); const Tensor* weights = context->Input(1); const Tensor* bias = context->Input(2); const Tensor* mask_index = context->Input(3); - ORT_RETURN_IF_ERROR(CheckInputs(input, weights, bias, mask_index)); + const Tensor* past = context->Input(4); + ORT_RETURN_IF_ERROR(CheckInputs(input, weights, bias, mask_index, past)); + // Input and output shapes: + // Input 0 - input : (batch_size, sequence_length, hidden_size) + // Output 0 - output : (batch_size, sequence_length, hidden_size) const auto dims = input->Shape().GetDims(); int batch_size = static_cast(dims[0]); int sequence_length = static_cast(dims[1]); @@ -55,8 +53,11 @@ Status Attention::ComputeInternal(OpKernelContext* context) const { TensorShape output_shape(dims); Tensor* output = context->Output(0, output_shape); + int past_sequence_length = 0; + Tensor* present = GetPresent(context, past, batch_size, head_size, sequence_length, past_sequence_length); + cublasHandle_t cublas = CublasHandle(); - const size_t element_size = sizeof(T); + constexpr size_t element_size = sizeof(T); // Use GEMM for fully connection. int m = batch_size * sequence_length; @@ -84,7 +85,7 @@ Status Attention::ComputeInternal(OpKernelContext* context) const { reinterpret_cast(input->template Data()), k, &one, reinterpret_cast(gemm_buffer.get()), n, device_prop)); - size_t workSpaceSize = GetAttentionWorkspaceSize(element_size, batch_size, num_heads_, head_size, sequence_length); + size_t workSpaceSize = GetAttentionWorkspaceSize(element_size, batch_size, num_heads_, head_size, sequence_length, past_sequence_length); auto temp_buffer = GetScratchBuffer(workSpaceSize); if (!LaunchAttentionKernel( reinterpret_cast(gemm_buffer.get()), @@ -97,7 +98,11 @@ Status Attention::ComputeInternal(OpKernelContext* context) const { temp_buffer.get(), cublas, element_size, - is_unidirectional_)) { + is_unidirectional_, + past_sequence_length, + nullptr == past ? nullptr : past->template Data(), + nullptr == present ? nullptr : present->template MutableData() + )) { // Get last error to reset it to cudaSuccess. CUDA_CALL(cudaGetLastError()); return Status(common::ONNXRUNTIME, common::FAIL); diff --git a/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu b/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu index 6f583bf472..9e7a61b518 100644 --- a/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu +++ b/onnxruntime/contrib_ops/cuda/bert/attention_impl.cu @@ -40,8 +40,8 @@ static size_t AlignTo(size_t a, size_t b) { return CeilDiv(a, b) * b; } -size_t ScratchSize(size_t element_size, int batch_size, int num_heads, int sequence_length) { - const size_t len = batch_size * num_heads * sequence_length * sequence_length; +size_t ScratchSize(size_t element_size, int batch_size, int num_heads, int sequence_length, int past_sequence_length) { + const size_t len = batch_size * num_heads * sequence_length * (sequence_length + past_sequence_length); const size_t bytes = len * element_size; const size_t alignment = 256; @@ -49,13 +49,19 @@ size_t ScratchSize(size_t element_size, int batch_size, int num_heads, int seque return bytesAligned; } -size_t GetAttentionWorkspaceSize(size_t element_size, int batch_size, int num_heads, int head_size, int sequence_length) { +size_t GetAttentionWorkspaceSize( + size_t element_size, + int batch_size, + int num_heads, + int head_size, + int sequence_length, + int past_sequence_length) { size_t qkv_size = 3 * batch_size * sequence_length * num_heads * head_size * element_size; - return qkv_size + 2 * ScratchSize(element_size, batch_size, num_heads, sequence_length); + return qkv_size + 2 * ScratchSize(element_size, batch_size, num_heads, sequence_length, past_sequence_length); } template -__device__ inline void Softmax(const int sequence_length, const int valid_length, const T* input, T* output, bool is_unidirectional) { +__device__ inline void Softmax(const int past_sequence_length, const int sequence_length, const int valid_length, const T* input, T* output, bool is_unidirectional) { using BlockReduce = cub::BlockReduce; __shared__ typename BlockReduce::TempStorage tmp_storage; @@ -64,13 +70,14 @@ __device__ inline void Softmax(const int sequence_length, const int valid_length float thread_data_max(-CUDART_INF_F); - const int num_valid = is_unidirectional ? (blockIdx.x % sequence_length) + 1 : valid_length; + const int num_valid = is_unidirectional ? past_sequence_length + (blockIdx.x % sequence_length) + 1 : valid_length; // e^x is represented as infinity if x is large enough, like 100.f. // Infinity divided by Infinity is a NAN. Thus, softmax gets a NAN if one or more item are large enough. // a math transform as below is leveraged to get a stable softmax: // e^xi/(e^x1 + ...e^xn) = e^(xi - max) / (e^(x1 - max) + ... + e^(xn - max)) - const int offset = (blockIdx.y * gridDim.x + blockIdx.x) * sequence_length; + const int all_sequence_length = past_sequence_length + sequence_length; + const int offset = (blockIdx.y * gridDim.x + blockIdx.x) * all_sequence_length; for (int i = threadIdx.x; i < num_valid; i += TPB) { const int index = offset + i; if (thread_data_max < float(input[index])) { @@ -99,7 +106,7 @@ __device__ inline void Softmax(const int sequence_length, const int valid_length } __syncthreads(); - for (int i = threadIdx.x; i < sequence_length; i += TPB) { + for (int i = threadIdx.x; i < all_sequence_length; i += TPB) { const int index = offset + i; const float val = (i < num_valid) ? expf(float(input[index]) - max_block) * sum_reverse_block : 0.f; output[index] = T(val); @@ -107,17 +114,19 @@ __device__ inline void Softmax(const int sequence_length, const int valid_length } template -__device__ inline void SoftmaxSmall(const int sequence_length, const int valid_length, const T* input, T* output, bool is_unidirectional) { +__device__ inline void SoftmaxSmall(const int past_sequence_length, const int sequence_length, const int valid_length, const T* input, T* output, bool is_unidirectional) { using BlockReduce = cub::BlockReduce; __shared__ typename BlockReduce::TempStorage tmp_storage; __shared__ float sum_reverse_block; __shared__ float max_block; - const int offset = (blockIdx.y * gridDim.x + blockIdx.x) * sequence_length; + // Input dimension is BxNxSxS*; blockIdx.y is batch index b; gridDim.x=N*S; blockIdx.x is index within N*S; + const int all_sequence_length = past_sequence_length + sequence_length; + const int offset = (blockIdx.y * gridDim.x + blockIdx.x) * all_sequence_length; const int index = offset + threadIdx.x; - const int num_valid = is_unidirectional ? (blockIdx.x % sequence_length) + 1 : valid_length; + const int num_valid = is_unidirectional ? past_sequence_length + (blockIdx.x % sequence_length) + 1 : valid_length; // e^x is represented as infinity if x is large enough, like 100.f. // Infinity divided by Infinity is a NAN. Thus, softmax gets a NAN if one or more item are large enough. @@ -146,43 +155,44 @@ __device__ inline void SoftmaxSmall(const int sequence_length, const int valid_l // Store max value if (threadIdx.x == 0) { - sum_reverse_block = (1.f) / sum; + sum_reverse_block = (num_valid == 0) ? 0.f : (1.f) / sum; } __syncthreads(); - if (threadIdx.x < sequence_length) { + // threadIdx.x might be larger than all_sequence_length due to alignment to 32x. + if (threadIdx.x < all_sequence_length) { // this will be 0 for threadIdx.x >= num_valid output[index] = T(thread_data_exp * sum_reverse_block); } } template -__global__ void SoftmaxKernelSmall(const int sequence_length, const T* input, T* output, bool is_unidirectional) { - SoftmaxSmall(sequence_length, sequence_length, input, output, is_unidirectional); +__global__ void SoftmaxKernelSmall(const int past_sequence_length, const int sequence_length, const T* input, T* output, bool is_unidirectional) { + SoftmaxSmall(past_sequence_length, sequence_length, sequence_length, input, output, is_unidirectional); } template -__global__ void SoftmaxKernel(const int sequence_length, const T* input, T* output, bool is_unidirectional) { - Softmax(sequence_length, sequence_length, input, output, is_unidirectional); +__global__ void SoftmaxKernel(const int past_sequence_length, const int sequence_length, const T* input, T* output, bool is_unidirectional) { + Softmax(past_sequence_length, sequence_length, sequence_length, input, output, is_unidirectional); } template bool ComputeSoftmax( - cudaStream_t stream, const int sequence_length, const int batch_size, const int num_heads, + cudaStream_t stream, const int past_sequence_length, const int sequence_length, const int batch_size, const int num_heads, const T* input, T* output, bool is_unidirectional) { const dim3 grid(sequence_length * num_heads, batch_size, 1); if (sequence_length <= 32) { const int blockSize = 32; - SoftmaxKernelSmall<<>>(sequence_length, input, output, is_unidirectional); + SoftmaxKernelSmall<<>>(past_sequence_length, sequence_length, input, output, is_unidirectional); } else if (sequence_length <= 128) { const int blockSize = 128; - SoftmaxKernelSmall<<>>(sequence_length, input, output, is_unidirectional); + SoftmaxKernelSmall<<>>(past_sequence_length, sequence_length, input, output, is_unidirectional); } else if (sequence_length == 384) { const int blockSize = 384; - SoftmaxKernelSmall<<>>(sequence_length, input, output, is_unidirectional); + SoftmaxKernelSmall<<>>(past_sequence_length, sequence_length, input, output, is_unidirectional); } else { const int blockSize = 256; - SoftmaxKernel<<>>(sequence_length, input, output, is_unidirectional); + SoftmaxKernel<<>>(past_sequence_length, sequence_length, input, output, is_unidirectional); } return CUDA_CALL(cudaPeekAtLastError()); @@ -197,7 +207,7 @@ __global__ void MaskedSoftmaxKernelSmall(const int sequence_length, const int* m } __syncthreads(); - SoftmaxSmall(sequence_length, num_valid, input, output, false); + SoftmaxSmall(0, sequence_length, num_valid, input, output, false); } template @@ -209,7 +219,7 @@ __global__ void MaskedSoftmaxKernel(const int sequence_length, const int* mask_i } __syncthreads(); - Softmax(sequence_length, num_valid, input, output, false); + Softmax(0, sequence_length, num_valid, input, output, false); } template @@ -370,6 +380,89 @@ bool LaunchTransQkv(cudaStream_t stream, return CUDA_CALL(cudaPeekAtLastError()); } +template +__global__ void ConcatPastToPresent(const int sequence_length, + const T* past, + const T* k_v, + T* present) { + const int h = threadIdx.x; + const int n = threadIdx.y; + const int s = blockIdx.x; + const int b = blockIdx.y; + const int is_v = blockIdx.z; // 0 for k, 1 for v + + const int all_sequence_length = gridDim.x; + const int batch_size = gridDim.y; + const int num_heads = blockDim.y; + const int H = blockDim.x; + + // past: 2 x BxNxS'xH (past_k and past_v) + // k_v: 2 x BxNxSxH (k and v) + // present: 2 x BxNxS*xH (present_k and present_v) + const int past_sequence_length = all_sequence_length - sequence_length; + + const int present_SH = all_sequence_length * H; + const int present_NSH = num_heads * present_SH; + int out_offset = b * present_NSH + n * present_SH + s * H + h + is_v * (present_NSH * batch_size); + if (s < past_sequence_length) { + const int past_SH = past_sequence_length * H; + const int past_NSH = num_heads * past_SH; + const int in_offset = b * past_NSH + n * past_SH + s * H + h + is_v * (past_NSH * batch_size); + present[out_offset] = past[in_offset]; +} else if (s < all_sequence_length) { + const int SH = sequence_length * H; + const int NSH = num_heads * SH; + const int in_offset = b * NSH + n * SH + (s - past_sequence_length) * H + h + is_v * (NSH * batch_size); + present[out_offset] = k_v[in_offset]; + } +} + +bool LaunchConcatPastToPresent(cudaStream_t stream, + const int past_sequence_length, + const int sequence_length, + const int batch_size, + const int head_size, + const int num_heads, + const float* past, + const float* k_v, + float* present) { + const int all_sequence_length = past_sequence_length + sequence_length; + const dim3 grid(all_sequence_length, batch_size, 2); + if (0 == (head_size & 1)) { + const dim3 block(head_size / 2, num_heads, 1); + ConcatPastToPresent<<>>(sequence_length, reinterpret_cast(past), reinterpret_cast(k_v), reinterpret_cast(present)); + } else + { + const dim3 block(head_size, num_heads, 1); + ConcatPastToPresent<<>>(sequence_length, past, k_v, present); + } + return CUDA_CALL(cudaPeekAtLastError()); +} + +bool LaunchConcatPastToPresent(cudaStream_t stream, + const int past_sequence_length, + const int sequence_length, + const int batch_size, + const int head_size, + const int num_heads, + const half* past, + const half* k_v, + half* present) { + const int all_sequence_length = past_sequence_length + sequence_length; + const dim3 grid(all_sequence_length, batch_size, 2); + if (0 == (head_size % 4)) { + const dim3 block(head_size / 4, num_heads, 1); + ConcatPastToPresent<<>>(sequence_length, reinterpret_cast(past), reinterpret_cast(k_v), reinterpret_cast(present)); + } else if (0 == (head_size & 1)) { + const dim3 block(head_size / 2, num_heads, 1); + ConcatPastToPresent<<>>(sequence_length, reinterpret_cast(past), reinterpret_cast(k_v), reinterpret_cast(present)); + } else { // this should be an "odd" case. probably not worth catching it in the half2 kernel. + const dim3 block(head_size, num_heads, 1); + ConcatPastToPresent<<>>(sequence_length, past, k_v, present); + } + return CUDA_CALL(cudaPeekAtLastError()); +} + cublasStatus_t inline CublasGemmStridedBatched( cublasHandle_t handle, cublasOperation_t transa, cublasOperation_t transb, int m, int n, int k, const float alpha, @@ -394,8 +487,8 @@ bool QkvToContext( const int batch_size, const int sequence_length, const int num_heads, const int head_size, const size_t element_size, const T* input, T* output, T* workspace, const int* mask_index, - bool is_unidirectional) { - const size_t bytes = ScratchSize(element_size, batch_size, num_heads, sequence_length); + bool is_unidirectional, int past_sequence_length, const T* past, T* present) { + const size_t bytes = ScratchSize(element_size, batch_size, num_heads, sequence_length, past_sequence_length); T* scratch1 = workspace; T* scratch2 = scratch1 + (bytes / element_size); T* scratch3 = scratch2 + (bytes / element_size); @@ -409,7 +502,6 @@ bool QkvToContext( const int batches = batch_size * num_heads; const int size_per_batch = sequence_length * head_size; const int total_size = batches * size_per_batch; - const int temp_matrix_size = sequence_length * sequence_length; const T* q = scratch3; const T* k = q + total_size; @@ -418,29 +510,46 @@ bool QkvToContext( cublasSetStream(cublas, stream); CublasMathModeSetter helper(cublas, CUBLAS_TENSOR_OP_MATH); - // compute Q*K' (as K'*Q), scaled by 1/sqrt(H) and store in scratch1: BxNxSxS + // Concat past (2xBxNxS'xH) to present (2xBxNxS*xH): + // past_k (BxNxS'xH) + k (BxNxSxH) => present_k (BxNxS*xH) + // past_v (BxNxS'xH) + v (BxNxSxH) => present_v (BxNxS*xH) + const int present_size_per_batch = (past_sequence_length + sequence_length) * head_size; + if (nullptr != present) { + if (!LaunchConcatPastToPresent(stream, past_sequence_length, sequence_length, batch_size, head_size, num_heads, past, k, present)) { + return false; + } + + // update pointers to present_k and present_v. + k = present; + v = present + batches * present_size_per_batch; + } + + // compute Q*K' (as K'*Q), scaled by 1/sqrt(H) and store in scratch1: BxNxSxS* + // Q: BxNxSxH, K (present_k): BxNxS*xH, Q*K': BxNxSxS* const float rsqrt_head_size = 1.f / sqrt(static_cast(head_size)); + const int all_sequence_length = past_sequence_length + sequence_length; + const int temp_matrix_size = sequence_length * all_sequence_length; if (!CUBLAS_CALL(CublasGemmStridedBatched( - cublas, CUBLAS_OP_T, CUBLAS_OP_N, sequence_length, sequence_length, head_size, rsqrt_head_size, k, head_size, size_per_batch, - q, head_size, size_per_batch, 0.f, scratch1, sequence_length, temp_matrix_size, batches))) { + cublas, CUBLAS_OP_T, CUBLAS_OP_N, all_sequence_length, sequence_length, head_size, rsqrt_head_size, k, head_size, present_size_per_batch, + q, head_size, size_per_batch, 0.f, scratch1, all_sequence_length, temp_matrix_size, batches))) { return false; } - // apply softmax and store result P to scratch2: BxNxSxS + // apply softmax and store result P to scratch2: BxNxSxS* if (nullptr != mask_index) { if (!ComputeMaskedSoftmax(stream, sequence_length, batch_size, num_heads, mask_index, scratch1, scratch2)) { return false; } } else { - if (!ComputeSoftmax(stream, sequence_length, batch_size, num_heads, scratch1, scratch2, is_unidirectional)) { + if (!ComputeSoftmax(stream, past_sequence_length, sequence_length, batch_size, num_heads, scratch1, scratch2, is_unidirectional)) { return false; } } // compute P*V (as V*P), and store in scratch3: BxNxSxH if (!CUBLAS_CALL(CublasGemmStridedBatched( - cublas, CUBLAS_OP_N, CUBLAS_OP_N, head_size, sequence_length, sequence_length, 1.f, v, head_size, size_per_batch, - scratch2, sequence_length, temp_matrix_size, 0.f, scratch3, head_size, size_per_batch, batches))) { + cublas, CUBLAS_OP_N, CUBLAS_OP_N, head_size, sequence_length, all_sequence_length, 1.f, v, head_size, present_size_per_batch, + scratch2, all_sequence_length, temp_matrix_size, 0.f, scratch3, head_size, size_per_batch, batches))) { return false; } @@ -459,7 +568,10 @@ bool LaunchAttentionKernel( void* workspace, cublasHandle_t& cublas, const size_t element_size, - bool is_unidirectional) { + bool is_unidirectional, + int past_sequence_length, + const void* past, + void* present) { // use default stream const cudaStream_t stream = nullptr; @@ -467,12 +579,14 @@ bool LaunchAttentionKernel( return QkvToContext(cublas, stream, batch_size, sequence_length, num_heads, head_size, element_size, reinterpret_cast(input), reinterpret_cast(output), reinterpret_cast(workspace), - mask_index, is_unidirectional); + mask_index, is_unidirectional, + past_sequence_length, reinterpret_cast(past), reinterpret_cast(present)); } else { return QkvToContext(cublas, stream, batch_size, sequence_length, num_heads, head_size, element_size, reinterpret_cast(input), reinterpret_cast(output), reinterpret_cast(workspace), - mask_index, is_unidirectional); + mask_index, is_unidirectional, + past_sequence_length, reinterpret_cast(past), reinterpret_cast(present)); } } diff --git a/onnxruntime/contrib_ops/cuda/bert/attention_impl.h b/onnxruntime/contrib_ops/cuda/bert/attention_impl.h index e106d8f266..6e58e73072 100644 --- a/onnxruntime/contrib_ops/cuda/bert/attention_impl.h +++ b/onnxruntime/contrib_ops/cuda/bert/attention_impl.h @@ -7,7 +7,13 @@ namespace onnxruntime { namespace contrib { namespace cuda { -size_t GetAttentionWorkspaceSize(size_t element_size, int batchsize, int num_heads, int head_size, int sequence_length); +size_t GetAttentionWorkspaceSize( + size_t element_size, + int batchsize, + int num_heads, + int head_size, + int sequence_length, + int past_sequence_length); bool LaunchAttentionKernel( const void* input, // Input tensor @@ -20,7 +26,10 @@ bool LaunchAttentionKernel( void* workspace, // Temporary buffer cublasHandle_t& cublas, // Cublas handle const size_t element_size, // Element size of input tensor - bool is_unidirectional // Whether there is unidirecitonal mask. + bool is_unidirectional, // Whether there is unidirecitonal mask. + int past_sequence_length, // Sequence length in past state + const void* past, // Past state input + void* present // Present state output ); } // namespace cuda diff --git a/onnxruntime/contrib_ops/cuda/quantization/attention_quantization.cc b/onnxruntime/contrib_ops/cuda/quantization/attention_quantization.cc index 19b4d63723..6d4ed08d38 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/attention_quantization.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/attention_quantization.cc @@ -59,7 +59,7 @@ Status QAttention::CheckInputs(const Tensor* input, // Input 7 - weight_zero_point : scalar // Output : (batch_size, sequence_length, hidden_size) - ORT_RETURN_IF_ERROR(AttentionBase::CheckInputs(input, weights, bias, mask_index)); + ORT_RETURN_IF_ERROR(AttentionBase::CheckInputs(input, weights, bias, mask_index, nullptr)); ORT_RETURN_IF_NOT(IsScalarOr1ElementVector(input_scale_tensor), "input scale must be a scalar or 1D tensor of size 1"); @@ -161,7 +161,10 @@ Status QAttention::ComputeInternal(OpKernelContext* context) const { m, n); - size_t workSpaceSize = GetAttentionWorkspaceSize(element_size, batch_size, num_heads_, head_size, sequence_length); + const int past_sequence_length = 0; + const T* past_data = nullptr; + T* present_data = nullptr; + size_t workSpaceSize = GetAttentionWorkspaceSize(element_size, batch_size, num_heads_, head_size, sequence_length, past_sequence_length); auto temp_buffer = GetScratchBuffer(workSpaceSize); if (!LaunchAttentionKernel( reinterpret_cast(gemm_buffer.get()), @@ -174,7 +177,11 @@ Status QAttention::ComputeInternal(OpKernelContext* context) const { temp_buffer.get(), cublas, element_size, - is_unidirectional_)) { + is_unidirectional_, + past_sequence_length, + past_data, + present_data + )) { // Get last error to reset it to cudaSuccess. CUDA_CALL(cudaGetLastError()); return Status(common::ONNXRUNTIME, common::FAIL); diff --git a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc index fc0be49fa7..af85126923 100644 --- a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc +++ b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc @@ -309,10 +309,49 @@ mask_index shall not be provided.)DOC"; .Input(1, "weight", "2D input tensor with shape (hidden_size, 3 * hidden_size)", "T") .Input(2, "bias", "1D input tensor with shape (3 * hidden_size)", "T") .Input(3, "mask_index", "Attention mask index with shape (batch_size).", "M", OpSchema::Optional) - .Output(0, "output", "3D output tensor with shape (batch_size, sequence_length, hidden_size)", "T") + .Input(4, "past", "past state for key and value with shape (2, batch_size, num_heads, past_sequence_length, head_size).", "T", OpSchema::Optional) + .Output(0, "output", "3D output tensor with shape (batch_size, append_length, hidden_size)", "T") + .Output(1, "present", "present state for key and value with shape (2, batch_size, num_heads, past_sequence_length + sequence_length, head_size)", "T", OpSchema::Optional) .TypeConstraint("T", {"tensor(float)", "tensor(float16)"}, "Constrain input and output types to float tensors.") .TypeConstraint("M", {"tensor(int32)"}, "Constrain mask index to integer types") - .TypeAndShapeInferenceFunction(ONNX_NAMESPACE::propagateShapeAndTypeFromFirstInput); + .TypeAndShapeInferenceFunction([](ONNX_NAMESPACE::InferenceContext& ctx) { + propagateElemTypeFromInputToOutput(ctx, 0, 0); + if (ctx.getNumOutputs() > 1) { + propagateElemTypeFromInputToOutput(ctx, 0, 1); + } + + if (hasInputShape(ctx, 0)) { + propagateShapeFromInputToOutput(ctx, 0, 0); + + if (ctx.getNumOutputs() > 1) { + auto& input_shape = getInputShape(ctx, 0); + auto& input_dims = input_shape.dim(); + if (input_dims.size() != 3) { + fail_shape_inference("Inputs 0 shall be 3 dimensions"); + } + + if (hasInputShape(ctx, 4)) { + auto& past_shape = getInputShape(ctx, 4); + auto& past_dims = past_shape.dim(); + if (past_dims.size() != 5) { + fail_shape_inference("Inputs 4 shall be 5 dimensions"); + } + + if (past_dims[3].has_dim_value() && input_dims[1].has_dim_value()) { + auto all_sequence_length = past_shape.dim(3).dim_value() + input_shape.dim(1).dim_value(); + + ONNX_NAMESPACE::TensorShapeProto present_shape; + for (auto& dim : past_dims) { + *present_shape.add_dim() = dim; + } + present_shape.mutable_dim(3)->set_dim_value(all_sequence_length); + + updateOutputShape(ctx, 1, present_shape); + } + } + } + } + }); ONNX_CONTRIB_OPERATOR_SCHEMA(QAttention) .SetDomain(kMSDomain) diff --git a/onnxruntime/python/tools/transformers/BertOnnxModel.py b/onnxruntime/python/tools/transformers/BertOnnxModel.py index 50ddd29263..1dcaad65f2 100644 --- a/onnxruntime/python/tools/transformers/BertOnnxModel.py +++ b/onnxruntime/python/tools/transformers/BertOnnxModel.py @@ -183,7 +183,7 @@ class BertOnnxModel(OnnxModel): for node in self.nodes(): # Before: # input_ids --> Shape --> Gather(indices=0) --> Unsqueeze ------+ - # | | + # | | # | v # +----> Shape --> Gather(indices=1) --> Unsqueeze---> Concat --> ConstantOfShape -->Cast --> EmbedLayerNormaliation/ReduceSum # After: @@ -292,8 +292,18 @@ class BertOnnxModel(OnnxModel): attention = op_count['Attention'] gelu = op_count['Gelu'] + op_count['BiasGelu'] + op_count['FastGelu'] layer_norm = op_count['LayerNormalization'] + op_count['SkipLayerNormalization'] - is_optimized = (embed > 0) and (attention > 0) and (attention == gelu) and (layer_norm >= 2 * attention) - logger.info( - f"EmbedLayer={embed}, Attention={attention}, Gelu={gelu}, LayerNormalization={layer_norm}, Successful={is_optimized}" - ) - return is_optimized + is_perfect = (embed > 0) and (attention > 0) and (attention == gelu) and (layer_norm >= 2 * attention) + + if layer_norm == 0: + logger.debug("Layer Normalization not fused") + + if gelu == 0: + logger.debug("Gelu/FastGelu not fused") + + if embed == 0: + logger.debug("Embed Layer not fused") + + if attention == 0: + logger.debug("Attention not fused") + + return is_perfect diff --git a/onnxruntime/python/tools/transformers/Gpt2OnnxModel.py b/onnxruntime/python/tools/transformers/Gpt2OnnxModel.py index e6922bce3b..02824451bf 100644 --- a/onnxruntime/python/tools/transformers/Gpt2OnnxModel.py +++ b/onnxruntime/python/tools/transformers/Gpt2OnnxModel.py @@ -10,6 +10,8 @@ import numpy as np from collections import deque from onnx import ModelProto, TensorProto, numpy_helper from BertOnnxModel import BertOnnxModel +from fusion_gpt_attention_no_past import FusionGptAttentionNoPast +from fusion_gpt_attention import FusionGptAttention logger = logging.getLogger(__name__) @@ -19,135 +21,12 @@ class Gpt2OnnxModel(BertOnnxModel): super().__init__(model, num_heads, hidden_size) def fuse_attention(self): - """ - Fuse Attention subgraph into one Attention node. - """ - logger.debug(f"start attention fusion...") - - input_name_to_nodes = self.input_name_to_nodes() - output_name_to_node = self.output_name_to_node() - - attention_count = 0 - - for normalize_node in self.get_nodes_by_op_type("LayerNormalization"): - return_indice = [] - qkv_nodes = self.match_parent_path( - normalize_node, - ['Add', 'Reshape', 'Gemm', 'Reshape', 'Reshape', 'Transpose', 'MatMul'], - [0, None, 0, 0, 0, 0, 0], - output_name_to_node=output_name_to_node, - return_indice=return_indice - ) # yapf: disable - if qkv_nodes is None: - continue - (add_qkv, reshape_qkv, gemm_qkv, reshape_1, reshape_2, transpose_qkv, matmul_qkv) = qkv_nodes - - another_input = add_qkv.input[1 - return_indice[0]] - - v_nodes = self.match_parent_path( - matmul_qkv, - ['Transpose', 'Reshape', 'Split', 'Reshape', 'Gemm', 'Reshape'], - [1, 0, 0, 0, 0, 0]) # yapf: disable - if v_nodes is None: - logger.debug("fuse_attention: failed to match v path") - continue - (transpose_v, reshape_v, split_v, reshape_after_gemm, gemm, reshape_before_gemm) = v_nodes - - layernorm_before_attention = self.get_parent(reshape_before_gemm, 0, output_name_to_node) - if layernorm_before_attention is None or layernorm_before_attention.op_type != 'LayerNormalization': - logger.debug(f"failed to get layernorm before gemm. Got {layernorm_before_attention.op_type}") - continue - - if not another_input in layernorm_before_attention.input: - logger.debug("Add and LayerNormalization shall have one same input") - continue - - qk_nodes = self.match_parent_path(matmul_qkv, ['Softmax', 'Sub', 'Mul', 'Div', 'MatMul'], [0, 0, 0, 0, 0]) - if qk_nodes is not None: - (softmax_qk, sub_qk, mul_qk, div_qk, matmul_qk) = qk_nodes - mask_nodes = self.match_parent_path( - sub_qk, - ['Mul', 'Sub', 'Slice', 'Slice', 'Unsqueeze', 'Sub', 'Squeeze', 'Slice', 'Shape', 'Div'], - [1, 0, 1, 0, 1, 0, 0, 0, 0, 0]) # yapf: disable - if mask_nodes is None: - logger.debug("fuse_attention: failed to match mask path") - continue - div_mask = mask_nodes[-1] - - if div_qk != div_mask: - logger.debug("fuse_attention: skip since div_qk != div_mask") - continue - else: - # New pattern for gpt2 from PyTorch 1.5.0 and Transformers 2.9.0. - qk_nodes = self.match_parent_path(matmul_qkv, ['Softmax', 'Where', 'Div', 'MatMul'], [0, 0, 1, 0]) - if qk_nodes is None: - logger.debug("fuse_attention: failed to match qk path") - continue - (softmax_qk, where_qk, div_qk, matmul_qk) = qk_nodes - mask_nodes = self.match_parent_path( - where_qk, - ['Cast', 'Slice', 'Slice', 'Unsqueeze', 'Sub', 'Squeeze', 'Slice', 'Shape', 'Div'], - [ 0, 0, 0, 1, 0, 0, 0, 0, 0]) # yapf: disable - if mask_nodes is None: - logger.debug("fuse_attention: failed to match mask path") - continue - div_mask = mask_nodes[-1] - - if div_qk != div_mask: - logger.debug("fuse_attention: skip since div_qk != div_mask") - continue - - q_nodes = self.match_parent_path(matmul_qk, ['Transpose', 'Reshape', 'Split'], [0, 0, 0]) - if q_nodes is None: - logger.debug("fuse_attention: failed to match q path") - continue - (transpose_q, reshape_q, split_q) = q_nodes - if split_v != split_q: - logger.debug("fuse_attention: skip since split_v != split_q") - continue - - k_nodes = self.match_parent_path(matmul_qk, ['Transpose', 'Reshape', 'Split'], [1, 0, 0]) - if k_nodes is None: - logger.debug("fuse_attention: failed to match k path") - continue - (transpose_k, reshape_k, split_k) = k_nodes - if split_v != split_k: - logger.debug("fuse_attention: skip since split_v != split_k") - continue - - self.create_attention_node(gemm, gemm_qkv, layernorm_before_attention.output[0], reshape_qkv.output[0], - attention_count == 0) - # we rely on prune_graph() to clean old subgraph nodes: - # qk_nodes + q_nodes + k_nodes + v_nodes + mask_nodes + [reshape_qkv, transpose_qkv, matmul_qkv] - attention_count += 1 - - self.prune_graph() - logger.info(f"Fused Attention count:{attention_count}") - - def create_attention_node(self, gemm, gemm_qkv, input, output, add_graph_input): - attention_node_name = self.create_node_name('Attention') - attention_node = onnx.helper.make_node('Attention', - inputs=[input, gemm.input[1], gemm.input[2]], - outputs=[attention_node_name + "_output"], - name=attention_node_name) - attention_node.domain = "com.microsoft" - attention_node.attribute.extend( - [onnx.helper.make_attribute("num_heads", self.num_heads), - onnx.helper.make_attribute("unidirectional", 1)]) - - matmul_node = onnx.helper.make_node('MatMul', - inputs=[attention_node_name + "_output", gemm_qkv.input[1]], - outputs=[attention_node_name + "_matmul_output"], - name=attention_node_name + "_matmul") - - add_node = onnx.helper.make_node('Add', - inputs=[attention_node_name + "_matmul_output", gemm_qkv.input[2]], - outputs=[output], - name=attention_node_name + "_add") - - self.add_node(attention_node) - self.add_node(matmul_node) - self.add_node(add_node) + if len(self.model.graph.input) == 1 or len(self.model.graph.output) == 1: + fusion = FusionGptAttentionNoPast(self, self.num_heads) + fusion.apply() + else: + fusion = FusionGptAttention(self, self.num_heads) + fusion.apply() def postprocess(self): """ @@ -186,7 +65,12 @@ class Gpt2OnnxModel(BertOnnxModel): outputs=[add_node_name + "_output"], name=add_node_name) + self.replace_input_of_all_nodes(reshape_after_gemm.output[0], add_node_name + "_output") + + # Link root node output with MatMul + self.replace_input_of_all_nodes(root_node.output[0], matmul_node_name + "_input") root_node.output[0] = matmul_node_name + "_input" + self.replace_input_of_all_nodes(reshape_after_gemm.output[0], add_node_name + "_output") self.add_node(matmul_node) diff --git a/onnxruntime/python/tools/transformers/benchmark.py b/onnxruntime/python/tools/transformers/benchmark.py index ab16393a4d..4f321e93f0 100644 --- a/onnxruntime/python/tools/transformers/benchmark.py +++ b/onnxruntime/python/tools/transformers/benchmark.py @@ -222,12 +222,13 @@ def get_onnx_file_path(onnx_dir: str, model_name: str, input_count: int, optimiz def optimize_onnx_model_by_ort(onnx_model_path, ort_model_path, use_gpu, overwrite): if overwrite or not os.path.exists(ort_model_path): + from optimizer import optimize_model, get_fusion_statistics # Use onnxruntime to optimize model, which will be saved to *_ort.onnx opt_model = optimize_by_onnxruntime(onnx_model_path, use_gpu=use_gpu, optimized_model_path=ort_model_path, opt_level=99) - model_fusion_statistics[ort_model_path] = opt_model.get_fused_operator_statistics() + model_fusion_statistics[ort_model_path] = get_fusion_statistics(ort_model_path) else: logger.info(f"Skip optimization since model existed: {ort_model_path}") @@ -235,7 +236,7 @@ def optimize_onnx_model_by_ort(onnx_model_path, ort_model_path, use_gpu, overwri def optimize_onnx_model(onnx_model_path, optimized_model_path, model_type, num_attention_heads, hidden_size, use_gpu, fp16, overwrite): if overwrite or not os.path.exists(optimized_model_path): - from optimizer import optimize_model, optimize_by_onnxruntime + from optimizer import optimize_model from BertOnnxModel import BertOptimizationOptions optimization_options = BertOptimizationOptions(model_type) if fp16: diff --git a/onnxruntime/python/tools/transformers/benchmark_gpt2.py b/onnxruntime/python/tools/transformers/benchmark_gpt2.py index 2721f42b49..4722856f64 100644 --- a/onnxruntime/python/tools/transformers/benchmark_gpt2.py +++ b/onnxruntime/python/tools/transformers/benchmark_gpt2.py @@ -18,10 +18,10 @@ from transformers import GPT2Model, GPT2LMHeadModel, GPT2Tokenizer, AutoConfig logger = logging.getLogger('') -# Map alias to a tuple of Model Class and pretrained model name +# Map alias to a tuple of Model Class, Tokenizer, pretrained model name, use LMHead or not, use attention mask or not MODEL_CLASSES = { - "gpt2": (GPT2Model, GPT2Tokenizer, "gpt2"), - "distilgpt2": (GPT2LMHeadModel, GPT2Tokenizer, "distilgpt2"), + "gpt2": (GPT2Model, GPT2Tokenizer, "gpt2", False, False), + "distilgpt2": (GPT2LMHeadModel, GPT2Tokenizer, "distilgpt2", True, True), } @@ -44,12 +44,12 @@ def setup_environment(): dump_environment() -def pytorch_inference(model, input_ids, past=None, total_runs=100): +def pytorch_inference(model, input_ids, past=None, attention_mask=None, total_runs=100): latency = [] with torch.no_grad(): for _ in range(total_runs): start = time.time() - outputs = model(input_ids=input_ids, past=past) + outputs = model(input_ids=input_ids, past=past, attention_mask=attention_mask) latency.append(time.time() - start) average_latency = sum(latency) * 1000 / len(latency) @@ -57,10 +57,12 @@ def pytorch_inference(model, input_ids, past=None, total_runs=100): return outputs, average_latency -def onnxruntime_inference(ort_session, input_ids, past=None, total_runs=100): +def onnxruntime_inference(ort_session, input_ids, past=None, attention_mask=None, total_runs=100): ort_inputs = {'input_ids': numpy.ascontiguousarray(input_ids.cpu().numpy())} - # TODO: pass input tensor stored in GPU + if attention_mask is not None: + ort_inputs['attention_mask'] = numpy.ascontiguousarray(attention_mask.cpu().numpy()) + if past is not None: for i, past_i in enumerate(past): ort_inputs[f'past_{i}'] = numpy.ascontiguousarray(past[i].cpu().numpy()) @@ -82,6 +84,7 @@ def onnxruntime_inference_with_binded_io(ort_session, last_state, last_state_shape, past=None, + attention_mask=None, present=None, present_shape=None, total_runs=100): @@ -90,13 +93,17 @@ def onnxruntime_inference_with_binded_io(ort_session, # Bind inputs io_binding.bind_input('input_ids', input_ids.device.type, 0, numpy.longlong, list(input_ids.size()), input_ids.data_ptr()) + if attention_mask is not None: + io_binding.bind_input('attention_mask', attention_mask.device.type, 0, numpy.float32, list(attention_mask.size()), + attention_mask.data_ptr()) + if past is not None: for i, past_i in enumerate(past): io_binding.bind_input(f'past_{i}', past[i].device.type, 0, numpy.float32, list(past[i].size()), past[i].data_ptr()) # Bind outputs - io_binding.bind_output("last_state", last_state.device.type, 0, numpy.float32, last_state_shape, + io_binding.bind_output(ort_session.get_outputs()[0].name, last_state.device.type, 0, numpy.float32, last_state_shape, last_state.data_ptr()) if present is not None: for i, present_i in enumerate(present): @@ -125,6 +132,7 @@ def inference(model, ort_session, input_ids, past=None, + attention_mask=None, last_state=None, present=None, last_state_shape=None, @@ -132,12 +140,12 @@ def inference(model, total_runs=100, verify_outputs=True, with_io_binding=False): - outputs, torch_latency = pytorch_inference(model, input_ids, past, total_runs) - ort_outputs, ort_latency = onnxruntime_inference(ort_session, input_ids, past, total_runs) + outputs, torch_latency = pytorch_inference(model, input_ids, past, attention_mask, total_runs) + ort_outputs, ort_latency = onnxruntime_inference(ort_session, input_ids, past, attention_mask, total_runs) latencies = [torch_latency, ort_latency] if with_io_binding: ort_io_outputs, ort_io_latency = onnxruntime_inference_with_binded_io(ort_session, input_ids, last_state, - last_state_shape, past, present, + last_state_shape, past, attention_mask, present, present_shape, total_runs) latencies.append(ort_io_latency) if verify_outputs: @@ -203,7 +211,8 @@ def parse_arguments(): help='Use optimizer.py to optimize onnx model') parser.set_defaults(optimize_onnx=False) - parser.add_argument('--with_io_binding', + parser.add_argument('-i', + '--with_io_binding', required=False, action='store_true', help='Run ONNX Runtime with binded inputs and outputs. ') @@ -243,7 +252,9 @@ def setup_logger(verbose=True): logger.setLevel(logging_level) -def export_onnx(model, config, tokenizer, device, output_dir): +def export_onnx(model, config, tokenizer, device, output_dir, use_LMHead=False, use_attention_mask=False): + """ Export GPT-2 model with past state to ONNX model + """ model.to(device) inputs = tokenizer.encode_plus("Here is an example input for GPT2 model", @@ -251,56 +262,59 @@ def export_onnx(model, config, tokenizer, device, output_dir): return_tensors='pt') input_ids = inputs['input_ids'].to(device) logger.debug(f"input_ids={input_ids}") + + # Use example input to generate an example of past state. outputs = model(input_ids=input_ids, past=None) assert len(outputs) == 2 logger.debug(f"output 0 shape={outputs[0].shape}") logger.debug(f"outputs[1][0] shape={outputs[1][0].shape}") num_layer = model.config.n_layer - present_names = [f'present_{i}' for i in range(num_layer)] - output_names = ["last_state"] + present_names - - input_names = ['input_ids'] - - # input_ids has only one word for model with past state. - # Shape of input tensors: - # input_ids: (batch_size, 1) - # past_{i}: (2, batch_size, num_heads, seq_len, hidden_size/num_heads) - # Shape of output tensors: - # last_state: (batch_size, seq_len + 1, hidden_size) - # present_{i}: (2, batch_size, num_heads, seq_len + 1, hidden_size/num_heads) - dynamic_axes = {'input_ids': {0: 'batch_size'}, 'last_state': {0: 'batch_size', 1: 'seq_len_plus_1'}} - - for name in present_names: - dynamic_axes[name] = {1: 'batch_size', 3: 'seq_len_plus_1'} - past_names = [f'past_{i}' for i in range(num_layer)] - input_names = ['input_ids'] + past_names - dummy_past = [torch.zeros(list(outputs[1][0].shape), dtype=torch.float32, device=device) for _ in range(num_layer)] + present_names = [f'present_{i}' for i in range(num_layer)] + + # GPT2Model output last_state has shape (batch_size, all_seq_len, hidden_size) + # GPT2LMHeadModel output prediction_scores has shape (batch_size, all_seq_len, vocab_size) + # where all_seq_len = past_seq_len + seq_len + output_names = ["prediction_scores" if use_LMHead else "last_state"] + present_names + + # Shape of input tensors: + # input_ids: (batch_size, seq_len) + # past_{i}: (2, batch_size, num_heads, past_seq_len, hidden_size/num_heads) + # attention_mask: (batch_size, seq_len) + # Shape of output tensors: + # last_state: (batch_size, all_seq_len, hidden_size) + # present_{i}: (2, batch_size, num_heads, all_seq_len, hidden_size/num_heads) + dynamic_axes = {'input_ids': {0: 'batch_size', 1 : 'seq_len'}, output_names[0]: {0: 'batch_size', 1: 'all_seq_len'}} for name in past_names: - dynamic_axes[name] = {1: 'batch_size', 3: 'seq_len'} - logger.debug(f"vocab_size:{model.config.vocab_size}") + dynamic_axes[name] = {1: 'batch_size', 3: 'past_seq_len'} + for name in present_names: + dynamic_axes[name] = {1: 'batch_size', 3: 'all_seq_len'} + + if use_attention_mask: + dynamic_axes['attention_mask'] = {0: 'batch_size', 1: 'seq_len'} dummy_input_ids = torch.randint(low=0, high=model.config.vocab_size - 1, size=(1, 1), dtype=torch.int64, device=device) - logger.debug(f"dummy_input_ids={dummy_input_ids}") - export_inputs = (dummy_input_ids, tuple(dummy_past)) + # Use the example past state to create dummy past state inputs. + dummy_past = [torch.zeros(list(outputs[1][0].shape), dtype=torch.float32, device=device) for _ in range(num_layer)] - export_model_path = os.path.join(output_dir, 'gpt2_past.onnx') + dummy_mask = torch.ones([1, 1], dtype=torch.float32, device=device) if use_attention_mask else None + + model_name = "gpt2{}_past{}.onnx".format("_lm" if use_LMHead else "", "_mask" if use_attention_mask else "") + export_model_path = os.path.join(output_dir, model_name) - # Let's run performance test on PyTorch before updating environment variable. with torch.no_grad(): - outputs = model(input_ids=dummy_input_ids, past=dummy_past) - + outputs = model(input_ids=dummy_input_ids, past=dummy_past, attention_mask=dummy_mask) logger.debug(f"present_0 shape={outputs[1][0].shape}") torch.onnx.export(model, - args=export_inputs, + args=(dummy_input_ids, tuple(dummy_past), dummy_mask) if use_attention_mask else (dummy_input_ids, tuple(dummy_past)), f=export_model_path, - input_names=input_names, + input_names=['input_ids'] + past_names + (['attention_mask'] if use_attention_mask else []), output_names=output_names, example_outputs=outputs, dynamic_axes=dynamic_axes, @@ -324,7 +338,7 @@ def main(): os.makedirs(output_dir) use_torchscript = False - (model_class, tokenizer_class, model_name) = MODEL_CLASSES[args.model_type] + (model_class, tokenizer_class, model_name, use_LMHead, use_attention_mask) = MODEL_CLASSES[args.model_type] config = AutoConfig.from_pretrained(model_name, torchscript=use_torchscript, cache_dir=cache_dir) model = model_class.from_pretrained(model_name, config=config, cache_dir=cache_dir) tokenizer = tokenizer_class.from_pretrained(model_name, cache_dir=cache_dir) @@ -332,7 +346,7 @@ def main(): # model = torch.jit.trace(model, (input_ids, past)) device = torch.device("cuda:0" if args.use_gpu else "cpu") - export_model_path = export_onnx(model, config, tokenizer, device, output_dir) + export_model_path = export_onnx(model, config, tokenizer, device, output_dir, use_LMHead, use_attention_mask) # setup environment variables before importing onnxruntime. setup_environment() @@ -379,7 +393,10 @@ def main(): ] # dummy last state - last_state_size = numpy.prod([max_batch_size, 1, config.hidden_size]) + if use_LMHead: + last_state_size = numpy.prod([max_batch_size, 1, config.vocab_size]) + else: + last_state_size = numpy.prod([max_batch_size, 1, config.hidden_size]) dummy_last_state = torch.empty(last_state_size).to(device) for batch_size in args.batch_sizes: @@ -394,6 +411,7 @@ def main(): size=(batch_size, 1), dtype=torch.int64, device=device) + dummy_mask = torch.ones([batch_size, 1], dtype=torch.float32, device=device) if use_attention_mask else None # Calculate the expected output shapes last_state_shape = [batch_size, 1, config.hidden_size] @@ -406,6 +424,7 @@ def main(): session, dummy_input_ids, dummy_past, + dummy_mask, dummy_last_state, dummy_present, last_state_shape, diff --git a/onnxruntime/python/tools/transformers/bert_perf_test.py b/onnxruntime/python/tools/transformers/bert_perf_test.py index 124a750b3e..16d8a04939 100644 --- a/onnxruntime/python/tools/transformers/bert_perf_test.py +++ b/onnxruntime/python/tools/transformers/bert_perf_test.py @@ -411,7 +411,11 @@ def parse_arguments(): parser.add_argument('--input_ids_name', required=False, type=str, default=None, help="input name for input ids") parser.add_argument('--segment_ids_name', required=False, type=str, default=None, help="input name for segment ids") - parser.add_argument('--input_mask_name', required=False, type=str, default=None, help="input name for attention mask") + parser.add_argument('--input_mask_name', + required=False, + type=str, + default=None, + help="input name for attention mask") args = parser.parse_args() return args @@ -430,7 +434,8 @@ def main(): if not min(batch_size_set) >= 1 and max(batch_size_set) <= 128: raise Exception("batch_size not in range [1, 128]") - model_setting = ModelSetting(args.model, args.input_ids_name, args.segment_ids_name, args.input_mask_name, args.opt_level) + model_setting = ModelSetting(args.model, args.input_ids_name, args.segment_ids_name, args.input_mask_name, + args.opt_level) for batch_size in batch_size_set: test_setting = TestSetting( diff --git a/onnxruntime/python/tools/transformers/bert_test_data.py b/onnxruntime/python/tools/transformers/bert_test_data.py index 4dc7f74814..cb2ccb16e9 100644 --- a/onnxruntime/python/tools/transformers/bert_test_data.py +++ b/onnxruntime/python/tools/transformers/bert_test_data.py @@ -237,7 +237,11 @@ def parse_arguments(): parser.add_argument('--input_ids_name', required=False, type=str, default=None, help="input name for input ids") parser.add_argument('--segment_ids_name', required=False, type=str, default=None, help="input name for segment ids") - parser.add_argument('--input_mask_name', required=False, type=str, default=None, help="input name for attention mask") + parser.add_argument('--input_mask_name', + required=False, + type=str, + default=None, + help="input name for attention mask") parser.add_argument('--samples', required=False, type=int, default=1, help="number of test cases to be generated") diff --git a/onnxruntime/python/tools/transformers/dev_benchmark.cmd b/onnxruntime/python/tools/transformers/dev_benchmark.cmd index 7ef3f7dd00..192b5c8d4d 100644 --- a/onnxruntime/python/tools/transformers/dev_benchmark.cmd +++ b/onnxruntime/python/tools/transformers/dev_benchmark.cmd @@ -48,7 +48,7 @@ REM This script will generate a logs file with a list of commands used in tests. >benchmark.log echo echo ort=%run_ort% torch=%run_torch% torchscript=%run_torchscript% gpu_fp32=%run_gpu_fp32% gpu_fp16=%run_gpu_fp16% cpu=%run_cpu% optimizer=%use_optimizer% batch="%batch_sizes%" sequence="%sequence_length%" models="%models_to_test%" input_counts="%input_counts%" REM Set it to false to skip testing. You can use it to dry run this script with the benchmark.log file. -set run_tests=false +set run_tests=true REM ------------------------------------------- if %run_cpu% == true if %run_gpu_fp32% == true echo cannot test cpu and gpu at same time & goto :EOF diff --git a/onnxruntime/python/tools/transformers/fusion_attention.py b/onnxruntime/python/tools/transformers/fusion_attention.py index 4ada0f690d..67a73b91df 100644 --- a/onnxruntime/python/tools/transformers/fusion_attention.py +++ b/onnxruntime/python/tools/transformers/fusion_attention.py @@ -188,10 +188,8 @@ class FusionAttention(Fusion): # Note that Cast might be removed by OnnxRuntime so we match two patterns here. _, mask_nodes, _ = self.model.match_parent_paths( - add_qk, - [(['Mul', 'Sub', 'Cast', 'Unsqueeze', 'Unsqueeze'], [1, 0, 1, 0, 0]), - (['Mul', 'Sub', 'Unsqueeze', 'Unsqueeze'], [1, 0, 1, 0])], - output_name_to_node) + add_qk, [(['Mul', 'Sub', 'Cast', 'Unsqueeze', 'Unsqueeze'], [1, 0, 1, 0, 0]), + (['Mul', 'Sub', 'Unsqueeze', 'Unsqueeze'], [1, 0, 1, 0])], output_name_to_node) if mask_nodes is None: logger.debug("fuse_attention: failed to match mask path") return @@ -215,4 +213,3 @@ class FusionAttention(Fusion): # Use prune graph to remove mask nodes since they are shared by all attention nodes. #self.nodes_to_remove.extend(mask_nodes) self.prune_graph = True - \ No newline at end of file diff --git a/onnxruntime/python/tools/transformers/fusion_base.py b/onnxruntime/python/tools/transformers/fusion_base.py index c5f7ac1521..94b19eb7f7 100644 --- a/onnxruntime/python/tools/transformers/fusion_base.py +++ b/onnxruntime/python/tools/transformers/fusion_base.py @@ -10,17 +10,21 @@ logger = getLogger(__name__) class Fusion: - def __init__(self, model: OnnxModel, name: str, search_op_types: Union[str, List[str]]): + def __init__(self, + model: OnnxModel, + fused_op_type: str, + search_op_types: Union[str, List[str]], + description: str = None): self.search_op_types: List[str] = [search_op_types] if isinstance(search_op_types, str) else search_op_types - self.name: str = name + self.fused_op_type: str = fused_op_type + self.description: str = f"{fused_op_type}({description})" if description else fused_op_type self.model: OnnxModel = model self.nodes_to_remove: List = [] self.nodes_to_add: List = [] self.prune_graph: bool = False def apply(self): - logger.debug(f"start {self.name} fusion...") - + logger.debug(f"start {self.description} fusion...") input_name_to_nodes = self.model.input_name_to_nodes() output_name_to_node = self.model.output_name_to_node() @@ -29,7 +33,10 @@ class Fusion: for node in self.model.get_nodes_by_op_type(search_op_type): self.fuse(node, input_name_to_nodes, output_name_to_node) - logger.info(f"Fused {self.name} count: {len(self.nodes_to_add)}") + op_list = [node.op_type for node in self.nodes_to_add] + count = op_list.count(self.fused_op_type) + if count > 0: + logger.info(f"Fused {self.description} count: {count}") self.model.remove_nodes(self.nodes_to_remove) self.model.add_nodes(self.nodes_to_add) diff --git a/onnxruntime/python/tools/transformers/fusion_biasgelu.py b/onnxruntime/python/tools/transformers/fusion_biasgelu.py index 692c82f947..2d7fd3d45a 100644 --- a/onnxruntime/python/tools/transformers/fusion_biasgelu.py +++ b/onnxruntime/python/tools/transformers/fusion_biasgelu.py @@ -14,7 +14,7 @@ logger = getLogger(__name__) class FusionBiasGelu(Fusion): def __init__(self, model: OnnxModel, is_fastgelu): if is_fastgelu: - super().__init__(model, 'FastGelu(add bias)', 'FastGelu') + super().__init__(model, 'FastGelu', 'FastGelu', 'add bias') else: super().__init__(model, 'BiasGelu', 'Gelu') diff --git a/onnxruntime/python/tools/transformers/fusion_embedlayer.py b/onnxruntime/python/tools/transformers/fusion_embedlayer.py index f8f5c04839..92af1a749d 100644 --- a/onnxruntime/python/tools/transformers/fusion_embedlayer.py +++ b/onnxruntime/python/tools/transformers/fusion_embedlayer.py @@ -38,9 +38,8 @@ class FusionEmbedLayerNoMask(Fusion): """ def __init__(self, model: OnnxModel, - name: str = "EmbedLayerNormalization(no mask)", - search_op_types="SkipLayerNormalization"): - super().__init__(model, name, search_op_types) + description='no mask'): + super().__init__(model, "EmbedLayerNormalization", "SkipLayerNormalization", description) self.utils = FusionUtils(model) def fuse(self, node, input_name_to_nodes, output_name_to_node): @@ -141,10 +140,9 @@ class FusionEmbedLayerNoMask(Fusion): # Cast might be removed by OnnxRuntime. _, segment_id_path, _ = self.model.match_parent_paths( - segment_ids_cast_node, + segment_ids_cast_node, [(['ConstantOfShape', 'Concat', 'Unsqueeze', 'Gather', 'Shape', 'Cast'], [0, 0, 1, 0, 0, 0]), - (['ConstantOfShape', 'Concat', 'Unsqueeze', 'Gather', 'Shape'], [0, 0, 1, 0, 0])], - output_name_to_node) + (['ConstantOfShape', 'Concat', 'Unsqueeze', 'Gather', 'Shape'], [0, 0, 1, 0, 0])], output_name_to_node) if segment_id_path and input_ids_cast_node and input_ids_cast_node.input[0] == segment_id_path[-1].input[0]: logger.debug("Simplify semgent id path...") @@ -187,7 +185,7 @@ class FusionEmbedLayerNoMask(Fusion): class FusionEmbedLayerNormalization(FusionEmbedLayerNoMask): def __init__(self, model: OnnxModel, mask_indice: Dict, mask_casted: Dict): - super().__init__(model, "EmbedLayerNormalization(with mask)", "SkipLayerNormalization") + super().__init__(model, "with mask") self.mask_indice: Dict = mask_indice self.mask_casted: Dict = mask_casted self.mask_input_name = None diff --git a/onnxruntime/python/tools/transformers/fusion_gelu_approximation.py b/onnxruntime/python/tools/transformers/fusion_gelu_approximation.py index e336938920..32f2e0af5c 100644 --- a/onnxruntime/python/tools/transformers/fusion_gelu_approximation.py +++ b/onnxruntime/python/tools/transformers/fusion_gelu_approximation.py @@ -11,7 +11,7 @@ from fusion_base import Fusion class FusionGeluApproximation(Fusion): def __init__(self, model: OnnxModel): - super().__init__(model, 'FastGelu(GeluApproximation)', ['Gelu', 'BiasGelu']) + super().__init__(model, 'FastGelu', ['Gelu', 'BiasGelu'], 'GeluApproximation') def fuse(self, node, input_name_to_nodes, output_name_to_node): new_node = helper.make_node("FastGelu", diff --git a/onnxruntime/python/tools/transformers/fusion_gpt_attention.py b/onnxruntime/python/tools/transformers/fusion_gpt_attention.py new file mode 100644 index 0000000000..2d060fba82 --- /dev/null +++ b/onnxruntime/python/tools/transformers/fusion_gpt_attention.py @@ -0,0 +1,196 @@ +#------------------------------------------------------------------------- +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. +#-------------------------------------------------------------------------- +import numpy as np +from logging import getLogger +from onnx import helper, numpy_helper, TensorProto +from OnnxModel import OnnxModel +from fusion_base import Fusion +from fusion_utils import FusionUtils + +logger = getLogger(__name__) + + +class FusionGptAttention(Fusion): + """ + Fuse GPT-2 Attention with past state subgraph into one Attention node. + This does not support attention_mask graph input right now. + """ + def __init__(self, model: OnnxModel, num_heads: int): + super().__init__(model, "Attention", "LayerNormalization", "with past") + self.num_heads = num_heads + + def create_attention_node(self, gemm, gemm_qkv, past, present, input, output): + attention_node_name = self.model.create_node_name('GptAttention') + mask_index = '' + attention_node = helper.make_node('Attention', + inputs=[input, gemm.input[1], gemm.input[2], mask_index, past], + outputs=[attention_node_name + "_output", present], + name=attention_node_name) + attention_node.domain = "com.microsoft" + attention_node.attribute.extend( + [helper.make_attribute("num_heads", self.num_heads), + helper.make_attribute("unidirectional", 1)]) + + matmul_node = helper.make_node('MatMul', + inputs=[attention_node_name + "_output", gemm_qkv.input[1]], + outputs=[attention_node_name + "_matmul_output"], + name=attention_node_name + "_matmul") + + add_node = helper.make_node('Add', + inputs=[attention_node_name + "_matmul_output", gemm_qkv.input[2]], + outputs=[output], + name=attention_node_name + "_add") + self.nodes_to_add.extend([attention_node, matmul_node, add_node]) + + def fuse(self, normalize_node, input_name_to_nodes, output_name_to_node): + past = None + present = None + return_indice = [] + qkv_nodes = self.model.match_parent_path( + normalize_node, + ['Add', 'Reshape', 'Gemm', 'Reshape', 'Reshape', 'Transpose', 'MatMul'], + [0, None, 0, 0, 0, 0, 0], + output_name_to_node=output_name_to_node, + return_indice=return_indice + ) # yapf: disable + if qkv_nodes is None: + return + (add_qkv, reshape_qkv, gemm_qkv, reshape_1, reshape_2, transpose_qkv, matmul_qkv) = qkv_nodes + + another_input = add_qkv.input[1 - return_indice[0]] + + v_nodes = self.model.match_parent_path( + matmul_qkv, + ['Concat', 'Transpose', 'Reshape', 'Split', 'Reshape', 'Gemm', 'Reshape'], + [1, 1, 0, 0, 0, 0, 0]) # yapf: disable + if v_nodes is None: + logger.debug("fuse_attention: failed to match v path") + return + (concat_v, transpose_v, reshape_v, split_v, reshape_after_gemm, gemm, reshape_before_gemm) = v_nodes + + # concat <-- Gather(indices=1) <-- past + # | + # unsqueeze + # | + # concat --> present + gather_v = self.model.get_parent(concat_v, 0, output_name_to_node) + if gather_v.op_type != 'Gather': + logger.info("expect Gather for past") + return + if not self.model.find_constant_input(gather_v, 1) == 1: + logger.info("expect indices=1 for Gather of past") + return + past = gather_v.input[0] + if not self.model.find_graph_input(past): + logger.info("expect past to be graph input") + return + unsqueeze_present_v = self.model.find_first_child_by_type(concat_v, + 'Unsqueeze', + input_name_to_nodes, + recursive=False) + if not unsqueeze_present_v: + logger.info("expect unsqueeze for present") + return + concat_present = self.model.find_first_child_by_type(unsqueeze_present_v, + 'Concat', + input_name_to_nodes, + recursive=False) + if not concat_present: + logger.info("expect concat for present") + return + present = concat_present.output[0] + if not self.model.find_graph_output(present): + logger.info("expect present to be graph input") + return + + layernorm_before_attention = self.model.get_parent(reshape_before_gemm, 0, output_name_to_node) + if layernorm_before_attention is None or layernorm_before_attention.op_type != 'LayerNormalization': + logger.debug(f"failed to get layernorm before gemm. Got {layernorm_before_attention.op_type}") + return + + if not another_input in layernorm_before_attention.input: + logger.debug("Add and LayerNormalization shall have one same input") + return + + qk_nodes = self.model.match_parent_path(matmul_qkv, ['Softmax', 'Sub', 'Mul', 'Div', 'MatMul'], [0, 0, 0, 0, 0]) + if qk_nodes is not None: + (softmax_qk, sub_qk, mul_qk, div_qk, matmul_qk) = qk_nodes + mask_nodes = self.model.match_parent_path( + sub_qk, + ['Mul', 'Sub', 'Slice', 'Slice', 'Unsqueeze', 'Sub', 'Squeeze', 'Slice', 'Shape', 'Div'], + [1, 0, 1, 0, 1, 0, 0, 0, 0, 0]) # yapf: disable + if mask_nodes is None: + logger.debug("fuse_attention: failed to match mask path") + return + div_mask = mask_nodes[-1] + + if div_qk != div_mask: + logger.debug("fuse_attention: skip since div_qk != div_mask") + return + else: + # New pattern for gpt2 from PyTorch 1.5.0 and Transformers 2.9.0. + qk_nodes = self.model.match_parent_path(matmul_qkv, ['Softmax', 'Where', 'Div', 'MatMul'], [0, 0, 1, 0]) + if qk_nodes is None: + logger.debug("fuse_attention: failed to match qk path") + return + (softmax_qk, where_qk, div_qk, matmul_qk) = qk_nodes + mask_nodes = self.model.match_parent_path( + where_qk, + ['Cast', 'Slice', 'Slice', 'Unsqueeze', 'Sub', 'Squeeze', 'Slice', 'Shape', 'Div'], + [ 0, 0, 0, 1, 0, 0, 0, 0, 0]) # yapf: disable + if mask_nodes is None: + logger.debug("fuse_attention: failed to match mask path") + return + div_mask = mask_nodes[-1] + + if div_qk != div_mask: + logger.debug("fuse_attention: skip since div_qk != div_mask") + return + + q_nodes = self.model.match_parent_path(matmul_qk, ['Transpose', 'Reshape', 'Split'], [0, 0, 0]) + if q_nodes is None: + logger.debug("fuse_attention: failed to match q path") + return + (transpose_q, reshape_q, split_q) = q_nodes + if split_v != split_q: + logger.debug("fuse_attention: skip since split_v != split_q") + return + + k_nodes = self.model.match_parent_path(matmul_qk, ['Concat', 'Transpose', 'Reshape', 'Split'], [1, 1, 0, 0]) + if k_nodes is None: + logger.debug("fuse_attention: failed to match k path") + return + (concat_k, transpose_k, reshape_k, split_k) = k_nodes + if split_v != split_k: + logger.debug("fuse_attention: skip since split_v != split_k") + return + + # concat_k <-- Transpose (perm=0,1,3,2) <-- Gather(axes=0, indices=0) <-- past + # | + # Transpose (perm=0,1,3,2) + # | + # unsqueeze + # | + # concat --> present + past_k_nodes = self.model.match_parent_path(concat_k, ['Transpose', 'Gather'], [0, 0]) + if past_k_nodes is None: + logger.debug("fuse_attention: failed to match past_k_nodes path") + return + + gather_past_k = past_k_nodes[-1] + if not self.model.find_constant_input(gather_past_k, 0) == 1: + logger.info("expect indices=0 for Gather k of past") + return + past_k = gather_past_k.input[0] + if past != past_k: + logger.info("expect past to be same") + return + + self.create_attention_node(gemm, gemm_qkv, past, present, layernorm_before_attention.output[0], + reshape_qkv.output[0]) + + # we rely on prune_graph() to clean old subgraph nodes: + # qk_nodes + q_nodes + k_nodes + v_nodes + mask_nodes + [reshape_qkv, transpose_qkv, matmul_qkv] + self.prune_graph = True diff --git a/onnxruntime/python/tools/transformers/fusion_gpt_attention_no_past.py b/onnxruntime/python/tools/transformers/fusion_gpt_attention_no_past.py new file mode 100644 index 0000000000..9585e3739a --- /dev/null +++ b/onnxruntime/python/tools/transformers/fusion_gpt_attention_no_past.py @@ -0,0 +1,137 @@ +#------------------------------------------------------------------------- +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. +#-------------------------------------------------------------------------- +import numpy as np +from logging import getLogger +from onnx import helper, numpy_helper, TensorProto +from OnnxModel import OnnxModel +from fusion_base import Fusion +from fusion_utils import FusionUtils + +logger = getLogger(__name__) + + +class FusionGptAttentionNoPast(Fusion): + """ + Fuse GPT-2 Attention without past state into one Attention node. + This does not support attention_mask graph input right now. + """ + def __init__(self, model: OnnxModel, num_heads: int): + super().__init__(model, "Attention", "LayerNormalization", "without past") + self.num_heads = num_heads + + def create_attention_node(self, gemm, gemm_qkv, input, output): + attention_node_name = self.model.create_node_name('Attention') + attention_node = helper.make_node('Attention', + inputs=[input, gemm.input[1], gemm.input[2]], + outputs=[attention_node_name + "_output"], + name=attention_node_name) + attention_node.domain = "com.microsoft" + attention_node.attribute.extend( + [helper.make_attribute("num_heads", self.num_heads), + helper.make_attribute("unidirectional", 1)]) + + matmul_node = helper.make_node('MatMul', + inputs=[attention_node_name + "_output", gemm_qkv.input[1]], + outputs=[attention_node_name + "_matmul_output"], + name=attention_node_name + "_matmul") + + add_node = helper.make_node('Add', + inputs=[attention_node_name + "_matmul_output", gemm_qkv.input[2]], + outputs=[output], + name=attention_node_name + "_add") + + self.nodes_to_add.extend([attention_node, matmul_node, add_node]) + + def fuse(self, normalize_node, input_name_to_nodes, output_name_to_node): + return_indice = [] + qkv_nodes = self.model.match_parent_path( + normalize_node, + ['Add', 'Reshape', 'Gemm', 'Reshape', 'Reshape', 'Transpose', 'MatMul'], + [0, None, 0, 0, 0, 0, 0], + output_name_to_node=output_name_to_node, + return_indice=return_indice + ) # yapf: disable + if qkv_nodes is None: + return + (add_qkv, reshape_qkv, gemm_qkv, reshape_1, reshape_2, transpose_qkv, matmul_qkv) = qkv_nodes + + another_input = add_qkv.input[1 - return_indice[0]] + + v_nodes = self.model.match_parent_path( + matmul_qkv, + ['Transpose', 'Reshape', 'Split', 'Reshape', 'Gemm', 'Reshape'], + [1, 0, 0, 0, 0, 0]) # yapf: disable + if v_nodes is None: + logger.debug("fuse_attention: failed to match v path") + return + (transpose_v, reshape_v, split_v, reshape_after_gemm, gemm, reshape_before_gemm) = v_nodes + + layernorm_before_attention = self.model.get_parent(reshape_before_gemm, 0, output_name_to_node) + if layernorm_before_attention is None or layernorm_before_attention.op_type != 'LayerNormalization': + logger.debug(f"failed to get layernorm before gemm. Got {layernorm_before_attention.op_type}") + return + + if not another_input in layernorm_before_attention.input: + logger.debug("Add and LayerNormalization shall have one same input") + return + + qk_nodes = self.model.match_parent_path(matmul_qkv, ['Softmax', 'Sub', 'Mul', 'Div', 'MatMul'], [0, 0, 0, 0, 0]) + if qk_nodes is not None: + (softmax_qk, sub_qk, mul_qk, div_qk, matmul_qk) = qk_nodes + mask_nodes = self.model.match_parent_path( + sub_qk, + ['Mul', 'Sub', 'Slice', 'Slice', 'Unsqueeze', 'Sub', 'Squeeze', 'Slice', 'Shape', 'Div'], + [1, 0, 1, 0, 1, 0, 0, 0, 0, 0]) # yapf: disable + if mask_nodes is None: + logger.debug("fuse_attention: failed to match mask path") + return + div_mask = mask_nodes[-1] + + if div_qk != div_mask: + logger.debug("fuse_attention: skip since div_qk != div_mask") + return + else: + # New pattern for gpt2 from PyTorch 1.5.0 and Transformers 2.9.0. + qk_nodes = self.model.match_parent_path(matmul_qkv, ['Softmax', 'Where', 'Div', 'MatMul'], [0, 0, 1, 0]) + if qk_nodes is None: + logger.debug("fuse_attention: failed to match qk path") + return + (softmax_qk, where_qk, div_qk, matmul_qk) = qk_nodes + mask_nodes = self.model.match_parent_path( + where_qk, + ['Cast', 'Slice', 'Slice', 'Unsqueeze', 'Sub', 'Squeeze', 'Slice', 'Shape', 'Div'], + [ 0, 0, 0, 1, 0, 0, 0, 0, 0]) # yapf: disable + if mask_nodes is None: + logger.debug("fuse_attention: failed to match mask path") + return + div_mask = mask_nodes[-1] + + if div_qk != div_mask: + logger.debug("fuse_attention: skip since div_qk != div_mask") + return + + q_nodes = self.model.match_parent_path(matmul_qk, ['Transpose', 'Reshape', 'Split'], [0, 0, 0]) + if q_nodes is None: + logger.debug("fuse_attention: failed to match q path") + return + (transpose_q, reshape_q, split_q) = q_nodes + if split_v != split_q: + logger.debug("fuse_attention: skip since split_v != split_q") + return + + k_nodes = self.model.match_parent_path(matmul_qk, ['Transpose', 'Reshape', 'Split'], [1, 0, 0]) + if k_nodes is None: + logger.debug("fuse_attention: failed to match k path") + return + (transpose_k, reshape_k, split_k) = k_nodes + if split_v != split_k: + logger.debug("fuse_attention: skip since split_v != split_k") + return + + self.create_attention_node(gemm, gemm_qkv, layernorm_before_attention.output[0], reshape_qkv.output[0]) + + # we rely on prune_graph() to clean old subgraph nodes: + # qk_nodes + q_nodes + k_nodes + v_nodes + mask_nodes + [reshape_qkv, transpose_qkv, matmul_qkv] + self.prune_graph = True diff --git a/onnxruntime/python/tools/transformers/fusion_layernorm.py b/onnxruntime/python/tools/transformers/fusion_layernorm.py index efabdaea90..3be6747bf2 100644 --- a/onnxruntime/python/tools/transformers/fusion_layernorm.py +++ b/onnxruntime/python/tools/transformers/fusion_layernorm.py @@ -117,7 +117,7 @@ class FusionLayerNormalization(Fusion): class FusionLayerNormalizationTF(Fusion): def __init__(self, model: OnnxModel): - super().__init__(model, "LayerNormalization", "Add") + super().__init__(model, "LayerNormalization", "Add", "TF") def fuse(self, node, input_name_to_nodes: Dict, output_name_to_node: Dict): """ diff --git a/onnxruntime/python/tools/transformers/fusion_skiplayernorm.py b/onnxruntime/python/tools/transformers/fusion_skiplayernorm.py index 2693f96dd0..93e52d37ae 100644 --- a/onnxruntime/python/tools/transformers/fusion_skiplayernorm.py +++ b/onnxruntime/python/tools/transformers/fusion_skiplayernorm.py @@ -47,7 +47,7 @@ class FusionSkipLayerNormalization(Fusion): class FusionBiasSkipLayerNormalization(Fusion): def __init__(self, model: OnnxModel): - super().__init__(model, "SkipLayerNormalization(add bias)", "SkipLayerNormalization") + super().__init__(model, "SkipLayerNormalization", "SkipLayerNormalization", "add bias") def fuse(self, node, input_name_to_nodes, output_name_to_node): if len(node.input) != 4: diff --git a/onnxruntime/python/tools/transformers/optimizer.py b/onnxruntime/python/tools/transformers/optimizer.py index b64a56e841..c4f1e0c07c 100644 --- a/onnxruntime/python/tools/transformers/optimizer.py +++ b/onnxruntime/python/tools/transformers/optimizer.py @@ -6,22 +6,17 @@ # Convert Bert ONNX model converted from TensorFlow or exported from PyTorch to use Attention, Gelu, # SkipLayerNormalization and EmbedLayerNormalization ops to optimize # performance on NVidia GPU and CPU. - +# # For Bert model exported from PyTorch, OnnxRuntime has bert model optimization support internally. -# You can use the option --use_onnxruntime to use model optimization from OnnxRuntime package. +# You can use the option --use_onnxruntime to check optimizations from OnnxRuntime. # For Bert model file like name.onnx, optimized model for GPU or CPU from OnnxRuntime will output as # name_ort_gpu.onnx or name_ort_cpu.onnx in the same directory. +# # This script is retained for experiment purpose. Useful senarios like the following: -# (1) Change model from fp32 to fp16. +# (1) Change model from fp32 to fp16 for mixed precision inference in GPU with Tensor Core. # (2) Change input data type from int64 to int32. # (3) Some model cannot be handled by OnnxRuntime, and you can modify this script to get optimized model. -# This script has been tested using the following models: -# (1) BertForSequenceClassification as in https://github.com/huggingface/transformers/blob/master/examples/run_glue.py -# PyTorch 1.2 or above, and exported to Onnx using opset version 10 or 11. -# (2) BertForQuestionAnswering as in https://github.com/huggingface/transformers/blob/master/examples/run_squad.py -# PyTorch 1.2 or above, and exported to Onnx using opset version 10 or 11. - import logging import coloredlogs import onnx @@ -29,6 +24,7 @@ import os import sys import argparse import numpy as np +from typing import Dict from collections import deque from onnx import ModelProto, TensorProto, numpy_helper, load_model from BertOnnxModel import BertOnnxModel, BertOptimizationOptions @@ -47,17 +43,21 @@ MODEL_CLASSES = { } -def optimize_by_onnxruntime(onnx_model_path, use_gpu=False, optimized_model_path=None, opt_level=99): +def optimize_by_onnxruntime(onnx_model_path: str, + use_gpu: bool = False, + optimized_model_path: str = None, + opt_level: int = 99) -> str: """ - Use onnxruntime package to optimize model. It could support models exported by PyTorch. + Use onnxruntime to optimize model. Args: - onnx_model_path (str): th path of input onnx model. + onnx_model_path (str): the path of input onnx model. use_gpu (bool): whether the optimized model is targeted to run in GPU. optimized_model_path (str or None): the path of optimized model. + opt_level (int): graph optimization level. Returns: - optimized_model_path: the path of optimized model + optimized_model_path (str): the path of optimized model """ import onnxruntime @@ -91,7 +91,22 @@ def optimize_by_onnxruntime(onnx_model_path, use_gpu=False, optimized_model_path return optimized_model_path -def parse_arguments(): +def get_fusion_statistics(optimized_model_path: str) -> Dict[str, int]: + """ + Get counter of fused operators in optimized model. + + Args: + optimized_model_path (str): the path of onnx model. + + Returns: + A dictionary with operator type as key, and count as value + """ + model = load_model(optimized_model_path, format=None, load_external_data=True) + optimizer = BertOnnxModel(model, num_heads=12, hidden_size=768) + return optimizer.get_fused_operator_statistics() + + +def _parse_arguments(): parser = argparse.ArgumentParser() parser.add_argument('--input', required=True, type=str, help="input onnx model path") @@ -192,7 +207,7 @@ def parse_arguments(): return args -def get_optimization_options(args): +def _get_optimization_options(args): optimization_options = BertOptimizationOptions(args.model_type) if args.disable_gelu: optimization_options.enable_gelu = False @@ -214,18 +229,38 @@ def get_optimization_options(args): def optimize_model(input, - model_type, - num_heads, - hidden_size, - opt_level=0, + model_type='bert', + num_heads=12, + hidden_size=768, optimization_options=None, + opt_level=0, use_gpu=False, only_onnxruntime=False): + """ Optimize Model by OnnxRuntime and/or offline fusion logic. + + The following optimizes model by OnnxRuntime only, and no offline fusion logic: + optimize_model(input, opt_level=1, use_gpu=False, only_onnxruntime=True) + If you want to optimize model by offline fusion logic. + optimize_model(input, model_type, num_heads=12, hidden_size=768, optimization_options=your_options) + + Args: + input (str): input model path. + model_type (str): model type - like bert, bert_tf, bert_keras or gpt2. + num_heads (int): number of attention heads. + hidden_size (int): hidden size. + optimization_options (OptimizationOptions or None): optimization options that can use to turn on/off some fusions. + opt_level (int): onnxruntime graph optimization level (0, 1, 2 or 99). When the level > 0, onnxruntime will be used to optimize model first. + use_gpu (bool): use gpu or not for onnxruntime. + only_onnxruntime (bool): only use onnxruntime to optimize model, and no offline fusion logic is used. + + Returns: + object of an optimizer class. + """ (optimizer_class, producer, run_onnxruntime) = MODEL_CLASSES[model_type] input_model_path = input - if opt_level > 1: # Optimization specified for an execution provider. + if opt_level > 1: # Optimization specified for an execution provider. input_model_path = optimize_by_onnxruntime(input_model_path, use_gpu=use_gpu, opt_level=opt_level) elif run_onnxruntime: # Use Onnxruntime to do optimizations (like constant folding and cast elimation) that is not specified to exection provider. @@ -242,15 +277,15 @@ def optimize_model(input, if optimization_options is None: optimization_options = BertOptimizationOptions(model_type) - bert_model = optimizer_class(model, num_heads, hidden_size) + optimizer = optimizer_class(model, num_heads, hidden_size) if not only_onnxruntime: - bert_model.optimize(optimization_options) + optimizer.optimize(optimization_options) - return bert_model + return optimizer -def setup_logger(verbose): +def _setup_logger(verbose): if verbose: coloredlogs.install(level='DEBUG', fmt='[%(filename)s:%(lineno)s - %(funcName)20s()] %(message)s') else: @@ -258,33 +293,33 @@ def setup_logger(verbose): def main(): - args = parse_arguments() + args = _parse_arguments() - setup_logger(args.verbose) + _setup_logger(args.verbose) - optimization_options = get_optimization_options(args) + optimization_options = _get_optimization_options(args) - bert_model = optimize_model(args.input, - args.model_type, - args.num_heads, - args.hidden_size, - opt_level=args.opt_level, - optimization_options=optimization_options, - use_gpu=args.use_gpu, - only_onnxruntime=args.only_onnxruntime) + optimizer = optimize_model(args.input, + args.model_type, + args.num_heads, + args.hidden_size, + opt_level=args.opt_level, + optimization_options=optimization_options, + use_gpu=args.use_gpu, + only_onnxruntime=args.only_onnxruntime) if args.float16: - bert_model.convert_model_float32_to_float16() + optimizer.convert_model_float32_to_float16() if args.input_int32: - bert_model.change_input_to_int32() + optimizer.change_input_to_int32() - bert_model.save_model_to_file(args.output) + optimizer.save_model_to_file(args.output) - if bert_model.is_fully_optimized(): - logger.info("The output model is fully optimized.") + if optimizer.is_fully_optimized(): + logger.info("The model has been fully optimized.") else: - logger.warning("The output model is not fully optimized. It might not be usable.") + logger.info("The model has been optimized.") if __name__ == "__main__": diff --git a/onnxruntime/python/tools/transformers/pytest.ini b/onnxruntime/python/tools/transformers/pytest.ini new file mode 100644 index 0000000000..0d98a076eb --- /dev/null +++ b/onnxruntime/python/tools/transformers/pytest.ini @@ -0,0 +1,9 @@ +[pytest] +log_cli = 1 +log_cli_level = INFO +log_cli_format = %(message)s + +log_file = pytest.log +log_file_level = DEBUG +log_file_format = %(asctime)s [%(levelname)8s] %(message)s (%(filename)s:%(lineno)s) +log_file_date_format=%Y-%m-%d %H:%M:%S \ No newline at end of file diff --git a/onnxruntime/python/tools/transformers/test_data/gpt2_pytorch1.5_opset11/generate_tiny_gpt2_model.py b/onnxruntime/python/tools/transformers/test_data/gpt2_pytorch1.5_opset11/generate_tiny_gpt2_model.py new file mode 100644 index 0000000000..3ee3147235 --- /dev/null +++ b/onnxruntime/python/tools/transformers/test_data/gpt2_pytorch1.5_opset11/generate_tiny_gpt2_model.py @@ -0,0 +1,363 @@ +#------------------------------------------------------------------------- +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. +#-------------------------------------------------------------------------- +# This tool generates a tiny GPT2 model for testing fusion script. +# You can use benchmark_gpt2.py to get a gpt2 ONNX model as input of this tool. + +import onnx +import onnx.utils +import sys +import argparse +import numpy as np +from onnx import ModelProto, TensorProto, numpy_helper +from OnnxModel import OnnxModel +import os +import onnxruntime +import random +from pathlib import Path +import timeit + +DICT_SIZE = 20 +SEQ_LEN = 5 +""" This class creates a tiny bert model for test purpose. """ + +# parameters of input base model. +old_parameters = { + "seq_len": 5, + "hidden_size": 768, + "num_heads": 12, + "size_per_head": 64, + "word_dict_size": [50257], # list of supported dictionary size. + "max_word_position": 1024 +} + +# parameters of output tiny model. +new_parameters = { + "seq_len": SEQ_LEN, + "hidden_size": 4, + "num_heads": 2, + "size_per_head": 2, + "word_dict_size": DICT_SIZE, + "max_word_position": 8 +} + + +class TinyGpt2Model(OnnxModel): + def __init__(self, model): + super(TinyGpt2Model, self).__init__(model) + self.resize_model() + + def resize_weight(self, initializer_name, target_shape): + weight = self.get_initializer(initializer_name) + w = numpy_helper.to_array(weight) + + target_w = w + if len(target_shape) == 1: + target_w = w[:target_shape[0]] + elif len(target_shape) == 2: + target_w = w[:target_shape[0], :target_shape[1]] + elif len(target_shape) == 3: + target_w = w[:target_shape[0], :target_shape[1], :target_shape[2]] + elif len(target_shape) == 4: + target_w = w[:target_shape[0], :target_shape[1], :target_shape[2], :target_shape[3]] + else: + print("at most 3 dimensions") + + tensor = onnx.helper.make_tensor(name=initializer_name + '_resize', + data_type=TensorProto.FLOAT, + dims=target_shape, + vals=target_w.flatten().tolist()) + + return tensor + + def resize_model(self): + graph = self.model.graph + initializers = graph.initializer + + for input in graph.input: + if (input.type.tensor_type.shape.dim[1].dim_value == old_parameters["seq_len"]): + print("input", input.name, input.type.tensor_type.shape) + input.type.tensor_type.shape.dim[1].dim_value = new_parameters["seq_len"] + print("=>", input.type.tensor_type.shape) + + reshapes = {} + for initializer in initializers: + tensor = numpy_helper.to_array(initializer) + if initializer.data_type == TensorProto.FLOAT: + dtype = np.float32 + elif initializer.data_type == TensorProto.INT32: + dtype = np.int32 + elif initializer.data_type == TensorProto.INT64: + dtype = np.int64 + else: + print("data type not supported by this tool:", dtype) + + if len(tensor.shape) == 1 and tensor.shape[0] == 1: + if tensor == old_parameters["num_heads"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["num_heads"], "=>[", new_parameters["num_heads"], "]") + initializer.CopyFrom( + numpy_helper.from_array(np.asarray([new_parameters["num_heads"]], dtype=dtype), + initializer.name)) + elif tensor == old_parameters["seq_len"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["seq_len"], "=>[", new_parameters["seq_len"], "]") + initializer.CopyFrom( + numpy_helper.from_array(np.asarray([new_parameters["seq_len"]], dtype=dtype), initializer.name)) + elif tensor == old_parameters["size_per_head"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["size_per_head"], "=>[", new_parameters["size_per_head"], "]") + initializer.CopyFrom( + numpy_helper.from_array(np.asarray([new_parameters["size_per_head"]], dtype=dtype), + initializer.name)) + elif tensor == old_parameters["hidden_size"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["hidden_size"], "=>[", new_parameters["hidden_size"], "]") + initializer.CopyFrom( + numpy_helper.from_array(np.asarray([new_parameters["hidden_size"]], dtype=dtype), + initializer.name)) + elif tensor == 4 * old_parameters["hidden_size"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + 4 * old_parameters["hidden_size"], "=>[", 4 * new_parameters["hidden_size"], "]") + initializer.CopyFrom( + numpy_helper.from_array(np.asarray([4 * new_parameters["hidden_size"]], dtype=dtype), + initializer.name)) + elif tensor == 3 * old_parameters["hidden_size"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + 3 * old_parameters["hidden_size"], "=>[", 3 * new_parameters["hidden_size"], "]") + initializer.CopyFrom( + numpy_helper.from_array(np.asarray([3 * new_parameters["hidden_size"]], dtype=dtype), + initializer.name)) + elif len(tensor.shape) == 0: + if tensor == old_parameters["num_heads"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["num_heads"], "=>", new_parameters["num_heads"]) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(new_parameters["num_heads"], dtype=dtype), initializer.name)) + elif tensor == old_parameters["seq_len"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["seq_len"], "=>", new_parameters["seq_len"]) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(new_parameters["seq_len"], dtype=dtype), initializer.name)) + elif tensor == old_parameters["size_per_head"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["size_per_head"], "=>", new_parameters["size_per_head"]) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(new_parameters["size_per_head"], dtype=dtype), + initializer.name)) + elif tensor == old_parameters["hidden_size"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + old_parameters["hidden_size"], "=>", new_parameters["hidden_size"]) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(new_parameters["hidden_size"], dtype=dtype), + initializer.name)) + elif tensor == 4 * old_parameters["hidden_size"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + 4 * old_parameters["hidden_size"], "=>", 4 * new_parameters["hidden_size"]) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(4 * new_parameters["hidden_size"], dtype=dtype), + initializer.name)) + elif tensor == 3 * old_parameters["hidden_size"]: + print("initializer type={}".format(initializer.data_type), initializer.name, + 3 * old_parameters["hidden_size"], "=>", 3 * new_parameters["hidden_size"]) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(3 * new_parameters["hidden_size"], dtype=dtype), + initializer.name)) + elif tensor == 1.0 / np.sqrt(old_parameters["size_per_head"]): + print("initializer type={}".format(initializer.data_type), initializer.name, + 1.0 / np.sqrt(old_parameters["size_per_head"]), "=>", + 1.0 / np.sqrt(new_parameters["size_per_head"])) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(1.0 / np.sqrt(new_parameters["size_per_head"]), dtype=dtype), + initializer.name)) + elif tensor == np.sqrt(old_parameters["size_per_head"]): + print("initializer type={}".format(initializer.data_type), initializer.name, + np.sqrt(old_parameters["size_per_head"]), "=>", np.sqrt(new_parameters["size_per_head"])) + initializer.CopyFrom( + numpy_helper.from_array(np.asarray(np.sqrt(new_parameters["size_per_head"]), dtype=dtype), + initializer.name)) + + new_shape = [] + shape_changed = False + for dim in tensor.shape: + if (dim == old_parameters["hidden_size"]): + new_shape.append(new_parameters["hidden_size"]) + shape_changed = True + elif (dim == 4 * old_parameters["hidden_size"]): + new_shape.append(4 * new_parameters["hidden_size"]) + shape_changed = True + elif (dim == 3 * old_parameters["hidden_size"]): + new_shape.append(3 * new_parameters["hidden_size"]) + shape_changed = True + elif (dim in old_parameters["word_dict_size"]): + new_shape.append(new_parameters["word_dict_size"]) + shape_changed = True + elif (dim == old_parameters["max_word_position"]): + new_shape.append(new_parameters["max_word_position"]) + shape_changed = True + else: + new_shape.append(dim) + if shape_changed: + reshapes[initializer.name] = new_shape + print("initializer", initializer.name, tensor.shape, "=>", new_shape) + + for initializer_name in reshapes: + self.replace_input_of_all_nodes(initializer_name, initializer_name + '_resize') + tensor = self.resize_weight(initializer_name, reshapes[initializer_name]) + self.model.graph.initializer.extend([tensor]) + + # Add node name, replace split node attribute. + nodes_to_add = [] + nodes_to_remove = [] + for i, node in enumerate(graph.node): + if node.op_type == "Split": + nodes_to_add.append( + onnx.helper.make_node('Split', + node.input, + node.output, + name="Split_{}".format(i), + axis=2, + split=[ + new_parameters["hidden_size"], new_parameters["hidden_size"], + new_parameters["hidden_size"] + ])) + nodes_to_remove.append(node) + print("update split", + [new_parameters["hidden_size"], new_parameters["hidden_size"], new_parameters["hidden_size"]]) + if node.op_type == "Constant": + for att in node.attribute: + if att.name == 'value': + if numpy_helper.to_array(att.t) == old_parameters["num_heads"]: + nodes_to_add.append( + onnx.helper.make_node('Constant', + inputs=node.input, + outputs=node.output, + value=onnx.helper.make_tensor(name=att.t.name, + data_type=TensorProto.INT64, + dims=[], + vals=[new_parameters["num_heads"] + ]))) + print("constant", att.t.name, old_parameters["num_heads"], "=>", + new_parameters["num_heads"]) + if numpy_helper.to_array(att.t) == np.sqrt(old_parameters["size_per_head"]): + nodes_to_add.append( + onnx.helper.make_node('Constant', + inputs=node.input, + outputs=node.output, + value=onnx.helper.make_tensor( + name=att.t.name, + data_type=TensorProto.FLOAT, + dims=[], + vals=[np.sqrt(new_parameters["size_per_head"])]))) + print("constant", att.t.name, np.sqrt(old_parameters["size_per_head"]), "=>", + np.sqrt(new_parameters["size_per_head"])) + else: + node.name = node.op_type + "_" + str(i) + for node in nodes_to_remove: + graph.node.remove(node) + graph.node.extend(nodes_to_add) + + for i, input in enumerate(self.model.graph.input): + if i > 0: + dim_proto = input.type.tensor_type.shape.dim[2] + dim_proto.dim_value = new_parameters["num_heads"] + dim_proto = input.type.tensor_type.shape.dim[4] + dim_proto.dim_value = new_parameters["size_per_head"] + + for i, output in enumerate(self.model.graph.output): + if i == 0: + dim_proto = output.type.tensor_type.shape.dim[2] + dim_proto.dim_value = new_parameters["hidden_size"] + if i > 0: + dim_proto = output.type.tensor_type.shape.dim[2] + dim_proto.dim_value = new_parameters["num_heads"] + dim_proto = output.type.tensor_type.shape.dim[4] + dim_proto.dim_value = new_parameters["size_per_head"] + + +def generate_test_data(onnx_file, + output_path, + batch_size=1, + use_cpu=True, + input_tensor_only=False, + dictionary_size=DICT_SIZE, + test_cases=1, + output_optimized_model=False): + + for test_case in range(test_cases): + sequence_length = 3 + input_1 = np.random.randint(dictionary_size, size=(batch_size, 1), dtype=np.int64) + tensor_1 = numpy_helper.from_array(input_1, 'input_ids') + + path = os.path.join(output_path, 'test_data_set_' + str(test_case)) + try: + os.mkdir(path) + except OSError: + print("Creation of the directory %s failed" % path) + else: + print("Successfully created the directory %s " % path) + + sess_options = onnxruntime.SessionOptions() + sess_options.graph_optimization_level = onnxruntime.GraphOptimizationLevel.ORT_DISABLE_ALL + sess = onnxruntime.InferenceSession(onnx_file, sess_options, providers=['CPUExecutionProvider']) + + input1_name = sess.get_inputs()[0].name + output_names = [output.name for output in sess.get_outputs()] + inputs = {input1_name: input_1} + + with open(os.path.join(path, 'input_{}.pb'.format(0)), 'wb') as f: + f.write(tensor_1.SerializeToString()) + + for i in range(12): + input_name = f"past_{i}" + input = np.random.rand(2, batch_size, new_parameters["num_heads"], sequence_length, + new_parameters["size_per_head"]).astype(np.float32) + tensor = numpy_helper.from_array(input, input_name) + inputs.update({input_name: input}) + + with open(os.path.join(path, 'input_{}.pb'.format(1 + i)), 'wb') as f: + f.write(tensor.SerializeToString()) + + if input_tensor_only: + return + + result = sess.run(output_names, inputs) + print("result 0 shape:", result[0].shape) + print("result 1 shape:", result[1].shape) + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument('--input', required=True, type=str) + parser.add_argument('--output', required=True, type=str) + parser.add_argument('--output_optimized_model', required=False, action='store_true') + parser.set_defaults(output_optimized_model=False) + args = parser.parse_args() + + model = ModelProto() + with open(args.input, "rb") as f: + model.ParseFromString(f.read()) + + bert_model = TinyGpt2Model(model) + + bert_model.update_graph() + bert_model.remove_unused_constant() + + print("opset verion", bert_model.model.opset_import[0].version) + + with open(args.output, "wb") as out: + out.write(bert_model.model.SerializeToString()) + + p = Path(args.output) + data_path = p.parent + + generate_test_data(args.output, + data_path, + batch_size=1, + use_cpu=True, + output_optimized_model=args.output_optimized_model) + + +if __name__ == "__main__": + main() diff --git a/onnxruntime/python/tools/transformers/test_data/gpt2_pytorch1.5_opset11/gpt2_past.onnx b/onnxruntime/python/tools/transformers/test_data/gpt2_pytorch1.5_opset11/gpt2_past.onnx new file mode 100644 index 0000000000000000000000000000000000000000..350ec55f89f3d7ec24e11ca0f6a8f543609f8b31 GIT binary patch literal 182687 zcmb5X?b0R5b)AQg!OY+#35qo2pj)I!j%-mJN>HCuRp*1EG$YYrMkve(2O)=lb~M43 z01FrlXBq@Bf5d;j0N(&Fz_;PsGW@KSm6?0jIa%lJ`BC%ssmiKd`LTB9sHqbA|Ju{}x4(I~EH9q^{HLFN`S$f^Z(qIq=)?N=$1k3~ z^>?p-{IfS7f9J!m|DzB7z;Ht>#*o>Q59TYeDM7r|L6yAUd^Sv`4CJw1yg>f*4nSD&ekgXn%>}h-17QThl6f#eDUF5yng$` zH($IufAoI&>*HU0dhhiwe)Q$z(J zN5-Z9qU-5*om^sPJ)U+ttKLB+QujWdcS@*xs~Qh}yWYKzmtptDAQjMy9z*v&UVHa` z`Se>qfBWWpKY#P1|MJ7PufBNm<&VC4Lz$N^Rp!uo?9|37Z~I$6|NOSZ(It8xpANN; zKha{`O9%Ds<5XH~T3gJ&#eO|Sg#B)XwGaQ%*Z<>t>eYfY-%-8msA@kh>mBFP2Q-Y1 zs@>zM4x_1osA}(c4qbYRD7x*^+hKIP#FBdhFb)Pk^BPO;4ZyXey^n^`QLTQw=m1(_ zJZ~_)tix!!#$hz66()Q{9Y*uZ!7%zax|J$vnqK_IU%mO8pZ(y?U%q+$*{k{R5ja@? z`SEw<8Nk^3I=}k5z}EGzzBNs{!~Ej?@Bi$tU(K{A|M%#c)G*Vsj;=Y6LGJO5r{7cY z#RJQ4bbj&v|Mt_LznW=8{{JNVCZD)u`<7?er#>~s3ZG*m=RWoS{J(#6u6=4!El(OR z^UKnw=5y&&{{tsxRqhO~eE)C1c>8L;eDt3FaAnBxa7yQoFWfAn{+=Ick_%Z+C@;*6>%c!Thq>!N<*gPq z$#RX}>NG#_J!*(W4RN9&)~VHriOgVXMR%aUNkirY9@VMU>C}DC-|2$G)9P2mgyVKl#b4)A^(K^vB1)`Sh*VKmPoOumADazV`Q@{KMD3 z{_%I7-uvq}Z`;4^uV+EaX-+k8*{K2E-LH~7{OLL@p}})<4HSrT#W9Z7iaX;^M-_KE zs;uo}K2$62%zUWMNq_(817=BI{PKGOoHG;pGFLoLaonF{k3Bzc(yDQP*5l8@_4C|D z+<7o#&6YF-I=5M?I@&oR&^bEI{jq;mC7ykUjzDLoJNH&C*^-7rXSqkOo-`KE5rr;$ zqj)?1FE7LRzdU0U|EiAvOI$o$p1bUn<3DT2T{K=VLV~mU&_%87@qbZkxAA`wI9#Sa z{x9=j{Hs1M5tA-USPkSB`SrCLs(-u)j4tvvtlsVBE~nz=G^4%5b2RNJ6N~c(l2MMhvUwuw1;p9F9-ds!JSGiSYvSPxXeHfU%ZlEV0Bm2FlZ~fOd1rV-n)_YAZ_!?)u*RczZU0;B3)%^47 zy1Fz#xDH%!wBuTV>DrtyZ9lQn?yKo*bi%pw#i@WRo64uK?tnvS1VUSxs&nRAij5E)eD;3B77WEvJ>t%6ND%!UTvLbZP=jH6j7c%!>6|Bu#= zs^V$~D1l$A$Z$9;-RrBGwfLV7*B8Ihz<}3>XK4RS*znCyy;}Q&rw>f1YPv8%1k}NB zIM>0T=d#)l?SF_X)I2WZP+(kgM;(0cV;`4?>wQ0+jR88>k!d=}j@rhg;RP4e>EUe5 z1|+7VnnA3<1S`-$0(R)`q#+xlfmHZEIHPsepw0_j5NhVQH(+DkxM&(wP)7Yd24mgx z1{cNob(nxmdRVlWv_1|23~l*TYI;V*+(WBA^} z#H3L%;V>r{hXNC@N?s#t{EJDLU=r5x4}Ekq{(YI67G9xVcD%hz&k@r!P3JU#;gv?@ zv5PQ!nI_1FfGpE^QRKH#5txEhWEnn1FfnNqLo;XwAut!GFu^J4Q;1;#I_fS={Baw^ zDy7i|)6jupI%lAmRRK^75hx}WIQM~KokXagz%PXMgc7kM+u#2Di#K1s(MP=SexQHp(Ro(( z&*Z{2brGdAL1sMt}7OG>{*L%Q-#G-pBUt3uwyJ+&di5S~H zwe~F))nsr{&pOHZLZa|b-%eBAG8tSF%AbS>Cn$k78C*sv->M3Fq*Ya0M*Asg2D{6K z-BeX-*XPh-%^{xBkca*nH_%#D%^`}3nNI4oCk?EVQ0JtMP^t>9EI7oI=2$0~f>u>< zWA`1RI^9X5{G>j6)+bwd@B_4%iEuT`Spy4>CJrn}u)VQ(QXNCzq$1#K8%#L3&X z8%03;4~1qn>lSD=nXSQnp6sVIr8=w1AaYujwa*TENm~p>5-o<}8oi_}hR6tG26k)Q zn$cpLb+TD%Q`ixCQjL2tAQKEoweZ~dQOz4JY9MHjb>09~7N-~A7Fz$2F4Ikyvu`RK zfz9~#|G~F%K0QOBeanA}q349oYH{_`|NQnRuYbXLrt0fI8s@Py|Hc!T$CbLTy21sQ zQ!x*ktYV(dOA|1m35a+PU3S+QwuQe&s{{iR)JGkyaAYg63Gt8&#F~DdM;EL}G(c(H zB24U|&F=ciZb+ZhLzG2B($peRVR>p18{9Ui2Z|M$U`3k4h5&nDNT1Ndh6SjH2?;fl z_^<6~+2YqVP3U32#@%)c^+*!pQ{w`A`_wJei-MW1!uzYxMr6zkjZ(Y%MeSYJ{hZL* zIH89*yN_SB&E`OnK$`WcnlnXm-5aKxs+tiig!hFS8&_3)U;6U=GQk=A*u$a~>UonV z$}^0)i4KhWIZz}JX|bvp#dD%apv>-96;My5>fp-1y!gP=7I-v&wG#@84V>XXA_~Hj zjZn}Z7&xjuCM3fKx_a|z`Oc&EjLAbP~ z25OeR)Sv@#QdhLO0t%MefmBn{M6_uYl914pdCx>Nc7QT_ShNi(5j^uz<4Wub_Uk3o-kv?7j)_A`+{+C2(q>Obsj2R= zj_S>lAk3&LLT^S(vomJo?$sbiHOR5A#?ZP9don#kB~*jZAiXEUoV9OiKj(_+W%?NT z`4>O^KlA7Mr}#0_r^BpCcZa|9$H*hx+B3gKsum8yAqg=$_khl4rxw%~RV1;RD1C-S za3DNDheA=R_NF4F6zzaMXk2?@Zn#(_nDA6?5=l?5J-R>_H8OmkOCsyJ)bQfo*@ zq-A@dGpf&7pEIR7LHDzYPFz&&GP_z=;vMNvS?+4Y&ex_oIp%K6Eg>W9z>bn(Np^xt zlpmYyx)yE7FN$Fn#W0Ivzm4wE5!Rg3JPwV?&iktusz~Az zj!Z8aFNeJ!?D^4Gz%>SzkZF2DitoHfFo3x!O@RQ@Ryjb{K$|0!~fc<<|615 zS24p?lt)5z`H}Jll3vV6FFyQFuRr_Yt0TPCd-_X^=?qm7UJ=Tw5A*!IB8xI(%Dr?~ zH3eM8`VKUOypLZ|Gp&ny9kj)2&br_c+KI53AuKBFBEbrw~Ux+F``i83S9 zNYZ3&k$8*$iLyAtHTBlr7Tw~1qAbu;JucetUsU~kRVNVnfc<(2wRf5MLrw+!R~W@2 zaYyA%5ipJMpDT5Fp7WyC8i85Ce;R}`Dj0A~jMTl44Z{Bmc5|^xL>z;dYB=^YU9mj< zZpZjL6aU~HZs+|A1jFhBV71=C2Z&r-5MIQ2KeX1k4yb;=OClLiWDSep19*TASNj0f z9ppjt0gzv#y5p@VC&plbG3Y)3(yP^-KtfYA%=_W4BCT8J{aoRfSJknGFYW;@%kIH8s_U{~ zPxF2ltH(C#oHTvH{ z(Te+vn1Y4i3=ZmR3g%c?`>(iR z6Aac4^3xYpBxzbEX}=n0RMy2cAkr1!LjgaH*`QuPY&J%IY4^F967U2K`B9(s+Cw^3 zH0Z)1*6Hii3pH5n`RW#Y;SYUqN5>WWy+)rR9w@C_boLin&;22ojcLGy8Y~%;fGO#? z_VIuZg4m&c%L=~X0VLDid0;X|!4id(m>`-~lvKs0RXGx^oq-+e#zi9wjLvJP=w>ds zuso0LvBy1cxPV18z@pd$ZlZ%mqZI<43!*QJm~aweoB|7s!fb#E>L$)y_%J@-dV9z`~Kwg%n2Mz@lSm=xYqILXwS!NG?8OvyTBA4YA z`vfaB8h|dM_R1@EBw)Km^eug1nI7>`%B{gLM^uJQI+R7ex~kUg3+wiVl*qX5&hikI zt%F#lG};uaQCItj)B1S#eO&<{P?Kl9LW=g_1NADz7EGKQF@|`p4hA(5YxdCoiK5Ew zABTeaqe3Jo%`I5rw~;Af);+2)n8J#_j!RgoFndY51cYD1uDG2vC`MzE2w_ne#=2OX z#7#XirGqE5RDcK;`mt8_25YPr*G3vp&`Nzk4SfHv#Co`B{5X!3>~YaHZIF6mQKExI zdWS{vAPdAnoi;cwv3{oJgw6(;4O+4th)8V0l$fU6MYa$jHenH)u)ro1h!ESP*i4Om zFTs=)p^Sa#rnrJ=5TS3o3y~i6vOB>oZht>#NrM-m)oAd#h(-p?KpgLcChipxC9$xq zk42!KMnwb>TFLViFM~8@q@!P0$~TCBpt=hYt2|whuLb@|3;YHUa8qq5NQ-k)S0u$E z(x?$YY~|^^)u{^+_|8`c0}ydaLIlXn&@FTO$Dsg3a7P_{?@r<%2D?8lnI^hbfeDzcdoaNjx~%dM6MWXqWqxHpCy_Wv*@gqyMKlPY zDB_rh#Tt!y6;U|J33eeB83(=q6Q=`U$x{py{!_IN4bfmiqN{scTa%}&8pTbVPzZFO zP&DF-8czZWsRT}N5pgB~6}2yQ^XsIMm-JU-7^u2$Rf@Qa6Wm1`d~j8FJ+#gABLUhl z;eookYi3hmyG$CBFjjqFWiul{Br`QAU^beWP+ND+bSr5yra{*DL%hO}k41z3>YA$Cre0eu|$tu-Ea@VAqwmvN-;l_IVX37u5 zt@8AfVn0rzKThPPC=dg4R^{pZ$Qc9@>%0^lUv)V?5%Cr!T6p3XZ?`B9XVNhdf!t2f z+RE!>@C11hi#`M8J?f_ZGxNsBwQ@sVYWJP%-dJ*L%yg^{W_>6O2e})W)H*3gLv#*6GCzi5BM# zRi4fjTs!$oj^q<3#V{0kI;y-NPv?ZrYH_v7)6wtINynxhTI$G6g1UiJ8PH~ReRmQQZHdlxOrFDxw zv7cem?)pi=MlqyK>d_D;L*mo|wP0xSY)I^)Hbj$p=^0Mw8LGoZO!&xZq;o*<7v5o8bPM$`BB33i%En_fLZ6^sEDVp{29#&O@ zdQl`NQ6wi(BqvcMr%sWeW30GEkO?PrHc%u87zb4)s8@ClE~gA+iNEsAP-*_k=VlQC z!GS~+gee=%pfehx09!9i$017q4aNa=xQ01llz98JJB*jhjrHRoQM|GkD*4k z<094x-R_(sK2WRe6x|G|bBg$wq0F8=q<9lm0ON|A=s+Jj_lOTfi(yD{LT5u#4CIZ2 zAw~D%N-odw8Au^_#*Ey(Oi?0OnnaD}If>Y~zJI2|GivMN*IN1ZuuY@DBAqo1GF zHc3L3*Q`ON!qfC(s$4V~rPff8RhbGuat4~B$W+jCc9{wqK23Qw)jk+B%Ib@@I+xn(O(od#X2y zJ?h*KZfuD8Vl}u}B|=CNn$lix8vHND)<=)X0FVSRc6p^b4I)gCN8x^7KVHysA9? z3{??cF~uL|R}AuW?xnk`Dc~~7JBYkm<>}`|Q(x&N&Pu)DmV&!-OEL53N1bUO`WHy5 zKJ!0ANJN8TooCpMSSP^Ix`WcXcmL3j5_Jr_#gOKRkfSC;QcEPNDI2oI|Fc>{s*N+M zjmkS9aD8M*b3}+xt0At$)I>?FwFU85CoUSvPsFRCH*Je<@jp=($XXv4t&8cisQU9+ zoj~Nn`K*p&VTUD22w@}sD^lYOE74?%$X<8wUo^!T1xCPsM1?~zvrJJm8f8?l;}}BO zF^-*Ox8g6($|X_fu83eej{QtmoFz}M7=LHtAH2hzc|Qb0nfD_*d>0=ea_tP45bFS1 zYXs;X7kZaWGN7a#7LgB7nhA=?2Y8@B9#kKI=5btilJ=j)7@T1Ydfh>KHQ(8$=AsGC z&;$iTAh+7QU#s#nEI^cgBYl0woLV>c%1tc585Y3veiXBP-tUDIDvR>f7^Up}c|S{d zy*wTC!WzD~2e_;{@8?II$!n7Z#_I9fgxa~J1^trPs*Jp0-p|vXR8qyO_kf$K^L~ER z8Rh$xJRPBJnD;AGgmbz8Cn0${GTS!q=SR)^k9Xzi7gpv4$04Tbyq_O6?_W!#+0BL6 ziPEHqxNID_&yRX`BU+HzR+OWmE=lu%o;Ki3hTDj7bW!feLF@X9*7d8vYk_VW@y*2$ zD;q3NM^T#$iNrTT+o2&F+`qsBV1wEQwL>#xhk9BKX~g%^sVBhu!ZJO0!bCl7!8_K8 zi$*3C(QbH$ZP5+xU$pf_`@(|M_i@pR`xh|<7o}z3COUvQ7vefrz8Rfnzyyq_y|d7Aek{T%{Wfh{$N*7z6J_zk-to<{Rm8bl>Dh?FX!;l8-2I`99vu2{tQ z9QWnv`E(v1+5y7}3DE1l1M<_uWm;POp#9yc1scF_lQb+)bs(s5;A+TsBp1XWtM2*5N7Q%p6(y1mv+3RS$MRXKvnId-tAy8bQ2&z?k%8#63;lEA~;HTnO->N<33SCxHBcn0C zy^N_BJyDlszW!fWrpGyEK1amMoOkE%;Qr zuJn-%-yiEuo~{cR}C#y(OR=#h|mZhs8Ja|_&~jC!Zk8`tOmN|2oot* z11<5e{o_y&Oh}t>QKZC$6@G&W3A1)I;Wnj(k}a@QebstNx&#EtLsPaOp`Z;)%aoo& zZ4o#1$dnGA&{CZ)yrg`LXp4g>Ti+k+MSVV#0nkc)Ky8a|rwtGR0V0y}YEZO6gwzu} zS_dJff%m#7u4n!n20WDs5qkV`Fn{W}EaDh!QM4(vi5TPa5 z1${$-2z1Q1W1nvzT%pTqTI3Vl(hZX6Ibw@0h!n!$MUp%mye^{dNZ>{)vbPSF9TZ_% zpNfp-Hi8JP@Y?ai#@+olc=F&H$k`{RQbk00?x* zBY@b-)BRSB#6i1U2ZQ=iiq$~0JZ%3s6od{E1W;PH=p+tm zAzdCCvjquh5-w65Tu>cUyh7Cc3}c`P)|mvl<7$k?I%Gf}u%|?CZ3k?u8yAff9oppZ z9^0Z@n1DnGPle}th%gN%#7A6|`r)E2E3scMp>{W`bdu1j$e&EBt$`6pdeP5*yO)zGYRn#m&Qk+g}x6HBBZVg<7iI( z6_V*)n2_w`3Pq8M>@AyQhl-@h#|eIFXC{G7PkM@B0x+kvZqWnXkTc(hi6&2{8Y!Ss zAWUQThCFTr6g_+THEp7S(0SOtv8z|zH(ib!=*eO5!ByS$(2P{B<6@~3 zx*`w7^L;a;JpG!2H`rEkXl4Y6Xr|VOSD1~45l~xq&BW1SA$?a#ZLSCq8lppR)#Y*3 znAbtv_*$v$Ke}O+r(YHOK}i)~ztYYVBJpie`~F(E;bnO>}Uzwe=z*=YaB82`{b)FB*r0ZhC<7 zy{ekL^8}^V=gn7y6w&B+SXy6ETEwEUUH9HswVaY+oOWehYgJW;coj{f@TX{y>a9 z9E*OL)FSX)eYk9^H3FjL=_~sdXgArQ!HqmatiFI+U&gQ2U6eTlyx2O2XcohoMD&3E z8li{p6)9WPCh3x59MKWAP ze+0aTM7xXkQJzj273>pxh_k-nzNQQo$jw7d=NVgsHH{U8C{SAW4ikHrw7Y&%uu%-@ zlX@gt(U3H$PiefeA+d|v5cNQDBiE3C!I0Z{(K|d+!*1Q&p+CQ=sFvmgD>SsR_c?8tOo2S8<;!z@vphpF%;9 zn}?>%gFbd2b$8>t0x_;3F$O~3@u0UMJuSv5CH13?EhlXs?M-d^EmLXj&gFN_= zAq|hwch?&6>%=J^V((zcc3i|dana~BQH<{t-3}?#8GYh~&irxFHhU7mb5+F0RqV=D zRL`|Xd>}}VLrUDsHAZ}(KpzY#*Q;JD$7h^j%0@GX;+a)Y=`5~6XmN%rtE;iJRBx zD)!qXRudihu!tH?X)v#78#K8&{`Px^mqcvT73mrN?QZq4rwp(4<#$J5xnvI1%du*%?OhOIB^T}viS>hE`_p5lNly_^9bwXFJjQkls%~TuAlYVUi zmaTAIXIaKbiHKOg_%^27`n5?|vUFYNR#_5S`(oukbqO~FG{icyK!jgm(De%25D;l6 zS}~#^`#JGET8`t{Z`i+HBh|l;$mKiG2t`Ly zTmyuf^^E1dX^6q4^scs** z5}pcKvc3)NBmH~1f*GC8`HFQ>MSw!F%eW0IqAQx5fT)2Al(rb&BUqNGuZ-I_ZiMxX z?p)rlO9}`^*um}9obJ_nlCNS;YM#VvWpi%a$l@kKBqF1NR~sezinY_dT9=0z&Qq)h zfe4|WK6E3ACLu~;r|gXzN%0`WF+48o<%Zn|9|))dfGSox$Mw#s(eqmV1 zZK@2PD&!@I>oMf`Bo1Py9*Hwh(PC=&@=4%CuSY|QzOHhmXBOI?!}Dk8HiOQ;^SScq z&Ac8K?eVOCV(;}oynON%UHL4$iJXQH_ka!SemukVjrUNz2Y0r25)ZM;=?SRV@gA!= zof_(7SurQgzIWKN8}9*kA#>bdXf(Px1HvrUP9yBkoielcK#$RD=aM+Hhu#B`MlAbW zx?9g>J)X+O+12*Bk^glcf#py|2yeKL|5nv9p|9S9sF3JAIOFpnu!`uA9q++B+658^ zOT5Ss`wHF_i_}11eVafNO|pY}D0zIj3k4FEh?0(azw|X{aaYs6&_18e1BfxV`VX}I z@8mxq)-W>Xm>||QZ`(k9gA7YMxOhU zh_k`?|CQE+rEHvl6Vnn9eNeW|JP)QW4g_bYMuZ~=5=-i+p5U1$AukC#NGzfAA;S?N zF9YoGrAYvB*+~`_mH2Y0JVZy*QRukpTP|YUB%#x+gI|5Frh+CqCCxf4qAM4=wb3{f zwH*;Gw-;(^NEh`TGt%t@P?ka-WLfh-u*-00yAYxWy+IS`ou3RYT{6+3HxJi z;M*W9S_z7aSh#z$^FL{0-S_Z_goSA2_60(&0fVdFJi|d7fAkOZmyueneW3sSr*Hk& zH^2Oz5v2#K5Pwrgi!DykD#Yc4UYp174WejA(8Gaf5gmIp9l;d$$8WNy1=L-j-NL-+ zkahwkG-qhf4W1CZB1;VK$T1q8h-|wj)KB(w#D!&B?eVCFmk9UMh&K{a1 z@`E^?@M^i++@OickAumN0}2sO&M}+q<%F38@C=`1G%@eMlRUfds3O`xx9OVC z3w+yMH2LFol&h`HMjXUO)DL9@3~jK)1jGTAM#K^*w}IGaRq`_%_`BCX{@I&?CS9~a z;7}bcMw1lSg5Yx4E#g>l6Ct0_8qQ*lJGn6cXoCxD+ho!s06KD}E zLaCu+|GIJ+v2@uYSdBVk>O`GU0qP=Jv^TP2EtRs7bG$&%jnU`^QD8jcb0gAkk?aOh z%%2=ka;PqhvqabE@AiR(=%9osra=zUZ&YTC(4jpPv<3?j ze-seKjL8AUWSw?wvWWHhiRqI=qfZDd8XyXpRd|{LVgjPT$L%5t3LX_oaz#RZ_6w1G zLxs|(lM6)K=t)Kt@>%*xMwIfmu@W&xo1aL=HkhBRW8L@gh=fV$n4*BsN+5cSm_<6k zoJ354R2yK*D*AT~LY{TxJ6Pw#*J`pdVk{9wP^^xo@V{OHTysiVD!_)BSzR-#K(I;{Ic$Y?EH6@xW;z|V8J6v%Ob#M`f@(jx9M zP&jdHD4Mhlm77ncve3`VE*ffYRLD~tX-X)1?V$j{V?z4f?ce`gWoA!{%vv3l{Bw;4 zAY`d2V?|;fiOUK76Naz1c_uOGpAHAmiitb=C?|d-he(;@KmdBhBm#_IeXcKJVF1Il z^O%x)5+xO$y9MhwGUOzVkUY`mJtt8czAj>qMDE10&u|i60BfpwB$23}?xPm5yoK%Q zi&#{2t6)@L;g5vu=olBV5P=`89wYaXP;D?it8cwR4aRCq(K*CCy-T$blha?FP^q1= zb+ol=gAIyhBQ$KAYwKH&=9Z3x>rfR4WONip9g)MVYkz{ZQnjJy*7=~cB`|AHZSIej z6j0AD2B|Lb5|LWvyFq!RhTjq5DU4a7hlazi6ms*zNu@&Lf$rX!wtOWHYp0qcj{Yx& zfPek=?PuQ${GqE$C~zbON93IQpydWfYg|m85Jy!jFc^+HBpMEjHv*etc$}bIQ=ICy zjice3*J6r|)6D58mIp)2?;ZfKIJKow(yaqTE1V+%BOH|_)SKs^#Bd*23ST!c0m80byA+g21DWI1`=MVsDKhAf$@!b1#NKV9BK%3ql`>HK|}p3|0e6 z`Bu*P#%XPFMI(3DXEdI|1Dm9vxaMOP%zE z1;G~gJ#RFpttyClO6DZh2oX?2b7D^@=D4QJ=o~}N5#p&Hn$t-f0=BU zv+bI(eoEQ6q#&MXr|gZlNJ&QscG@W$A=7=;^wCjqG_ao@7^aerJ@W|RQgtx6-Rt;b;v?8jL>MgqRuDp8Wb_l=PG5@C861PL-)K1y>>V9tGb_mpZT(&XAM{_$z zb2~?KJ4bUn$Ik6sVFsgtnxXubYhRmyphQj z-fx7TE=0x@-7teL?B8x7^F&7|o`}>tvOAA1?QWtc`RIW!Stl50NWr=$Sk^#yB5h+Xh3=wQPpqnB#JYgxeb<_evKGRE?k|8O52ADHXj#iK?p%9SinqGD|D~@vr4PJ!^g#qCv=6*@1VTqxpdt_= zmhN3X@E%Qr9^rK2=!aF@g$%Ekk2#6Tbtm$*RD_Qbk5Hc;*@W#h$beJ4yOeZIntnPu zu;Bb-Pu3>dM;O zi!h>2szB&AO5q5VV)-lqcU>UFuT3DdjZ!!&?T7?IM8b7}5Wm`jJ~`&XvUCrif{Rj7?!YR+|PLLtZ z9LT=0J6p|Z{KyqjFD&$@#E%dp7i#B`|Cka#LOjz>*{%PW5=07I$ zBRt*Jpw zz0-v<39=IvE1xluOjXB%mhR+|*VM@-q~FkUOr~WfOC2U?k?uJl1Dod%P)$i?Ls(qs zb{Q&Jr+iYQ6u75Mv%+T5<|Ym~wt+EUHiDeE_8&LDe{(ygNWa!Gk5VdsQ6=0o8sm#iXN4=31^-=O-- z{)3#0*pMGy0~N}hHH(@|xJ)8k)`zbk|9d)V3Cx59=G)56ujfL?I8DDF4C9g>yw^&9 z?^g1I$qb@+6^nK$u#x(#pC`c9) zt0lHcM)0`S)$)AuqG)1U3;oR(z15+hUmQSqw1|~gwW#NPD;DL3N{m{zDEgFQksFKL zAXp^rOt)wdLJ=$y>(VT0toy(5t#S({eccw`V=a}!d<+Z~45ijP{kCxQE2S_W1!Lhj z5(kCtSloB~F=KL!&tyVnn%$eAXP%0BjHYVLyH=MId$j1tFs+|wxnxAWyUNYJOvqur z(|idd048qaTPTHZanChHF+rL98t?dn*KbvDFs1(E(wTL1_|7DKn_v@Z<)*jUtKp7NNk(vPRL0YVxvgQq+xEZv{i%E>e6m zaT9bhcV@VW59NutoY_2oJ=tI{2iag~V=#=}YgXZ<2NuC8ELGLI&OJ~Zz#V&I}h zC_tM6hOg-3b|l%`7!}~HYxas!8Qz{5U0aM2 zyv=snGpt20Y7lv`F)AQj53$v#o|R=aK*FyFtA9>J5Hy7-#DS69p1pZzuU&6e!Da z^B5oH@n-0d)$NMAQ@n#FDE#b`uv6;<1@11gs1_8KE9KT^q#EtY%(wc=(of0*Hwd^B z4O>bHBxZCQX}Ir%!@3lCS}6rjkxPw<^-ulQVmf%1K|%{K?%?}CT{Nf=t5^P!W0L=6 z8~-loF)CK0qXi6KLkoe@Y~o^uxcE+&>Jd&g(CRfk!5em8m!b(HCYcGhsrp zha@U*f5tL9T8A6>v$P>hMC=?KyD@9`*@`i&Ht%E1$Lq*ftDE4MMQ}8P zVePa*7;_pkijs&hW1qRP_AWs2h~A;AFO-WR{?$88spTH%4_dCdjoK zWvGyrr2A{)|3Xb|M|-TRQaDkY#O3W2zQGx3OX8EC`$@;M1Z~Z8n5{IK*)>*w+bc+K zvf>OeLA+_l8EZ>s*p-emXuS7uMj=Si!pIHg2D`x-lP$Ben9OiC1!ssR2smRVW=3bS zq6b310nV6~nL90m-thLyQX!rfK**%A59POuGgch2)%~W_I84jWa?lf7)^AFUg{VI%dPW(O%lD_{&8fS)*uO2lKGn%{-zIESiQS*m13?Zxad%_xj2e875n(Np5FWX z%@;rU^lP8CFo|Mr9?CF#-Roy_BfeyFJ%}VrX$S)ac-P#2t0r(@5;(32nL>_AqrCzN zx~KT=-AIX2Tr>iWFypOq^yq;XX_ZwP##IJ0ltzkT zdmGE1XC9@L7Wp$_O1adS8n7pIY7|QKXbHaC_Np3$14)O>B-}OR7nR0k4VU08=|$xq zPi`{>HwCPMMC^^9$x`~N117xYa=g-P5@;rG%!*TRjGhmPUgKxKb6*a#6iF@w@+dwU z_OL(tp$maMN`?1aNbLe`IjqR1@Y0ED+&($_AG;9p+v@GSp1PEFay3I7H5UT!7+r{k zH?wvJYExCPY~Vsd<-sMjNdfHRx*D|6pwQ+)ES6mZK; z&gQLVi~eUPi=8e0nHBF;{V&90syuzhoYCZj%7bUTKhJmPwA6v>b~%`Qf3YA~x_q(l zfMj=|_XxP&4C9@+-%?5qy`Z^C!+1yTr`lf7(4-2Blge#xiwOG{h4sG^L4o1)K@71s zRIlVAepFvB5!qa;G1|UyyXgByz-8N&t0jz-6d=53yQWr)^fP#CE$SB0NBk{VTFeeD zutTlMQTwe;E)s@QdKoC8I>mu|0t0RG!6 zMo#$xtr*uh3NBH``+I>Dg3-l-kJ*u`?0O$y6uPm(s9Xvsee2D_1lsYg|83NV#`@y) zvfF^w-v;`qyJyJ9yY4wyqztty@o_-x z7r2-OA}Dr?#3yEpVu!Xx*0^v3y%C>e6}y*6}_b zhUN^U5wX{)@-k^eO!m(E-&9^!(Gaw~Q#iiAYd(`k1adzv++-AsNu$N2(PGkQF=?~} z(&zyD7%R`D5qsW8X9CZ7zw+cv7AMo8gen~ZseGR$CVUlT`4X+3*zrEQ1-Qi#w3dNj z#D*_Q&IGyOZYfhx3*1pIwg0&Y+hpftFfh-vNQ6U|x3=Rj5)SHf5%8p=1E3a|sIX{? z%0YKDM~o~GPYpNVA0uuM_Y?MEA&qHFvK_!jf7EfX=xiHLAdY!## z;mu%#gGz<>U|>75Tg-ngj&oW}J1yqI7TBwbWJd@8HfDuW8aMD#eP%}__Fyo&6PR&e z)BDu>iTU0eGfWJ`9ZiwR$=}s1WxAE&9)QIa;f*h-whzIYg zrp1KL0-=*+Sj2=s!(jU!#TfND$}y=G^FlqUns}n>(qf$ju}-Z^G)WSxO4Y>Aeouvd zS=B^Tc~CV`4Zg?~uD~Fxnuy*mOZ{^Db5k|l24Jw?{OUE2@Q+Q^#LrFDw2i=6h#*=m zg@bIWCVrNskT2E05T~8|3k*xFq@6h5h@6HNlP!yrx3KMeNAu$-(AQ|bTGWLI#~$fh zjLprKLVX|j7jrD^QSWySLWCoP;@IP@7IoU?2yqx^%9(Z{Vtr@I6>G`T*r`b3D#k321+^aL1c3CVVVvG9jUtka8WY4tAl4ut~GiX-@KVY+s)2Ko9;kU1h{e0az8$sc z93F5$`N;Js}TTJqPrm^|0m6iRmQVet72?o59^|vlfeN4&Q8ELlEM_UEL5pE8-F+ zY{d1$m3B9Hox~L8*NU+j>@=^3o5o>-`1)|DLRXl)4nNB>^g3&72Ep+mHiKQ>(}Pps zZ+>vyN5ZdUf{DVucXM~5{%ULn4d}<0*Q*OU;^vEOLUne5clj>-IKuITbqrg(XX|2{ zWK@p`s`0^&@jkj;Aa>fSftdL*s;?F)P-$gk2zgHyLDs~EN<>|@XuFO%9o<+?!NQL2 z!J<(rzfecY16(cYxjJ;I~%`&B>l3tu!THlzBIJIk2hh3>3kP7p8E*FXur zgReOq3_|cL$=5&$?o1`Er5}!ehxK80SR$l}e*kUu-2;Dc^F1z6T9vf`H&E&>0 z@qHQA<0|H4a^Yk`;RMC8dMda**QNc5gOAwDyscuO^?Gau)u0OZ_CY5oiFFkt8czZ) zk9p#5h|O5;$7Zmxd39qDdus~fHvFx1y~PKs0k*`o9Xf{^~9MK|7M(d3wo z&h-R@G2Nm)-fp!BR-@HX=k4rWl(9F|QQD7Vi}A6RO5sF9l2stQ&bIJ&6?1~}NKcNw ze^9tl*6hx%+Hi&zZsfl`-jt_h{r*X!f#hAxgPh<&f_EWycgMS!9ATUB45s_4K?c4D zfk}@Oq(=ipLbG=O(Sj3~QYfewgY=}ie^A~s{$kIl%ZHv>vJdP{5u=?UJ+ zznI25NgvY|a<3FWKJ*l&ZOUt9SrlmjhNbLZEN(d|>n@e*krgE_ z8fgc&R&{`zx_Hnypyas0sw5uFR^7mfKJ^|>PMC6n$p}{AwhP+jIndUs4ss+qsC8vj z$n*rk{-Lh61N#P2pdr#LJp&c`gcXV+6^7|EfcPOKJ*8q;Ujbc)vCy6)=@RWP)## z?u`rz4@3%o`1HQ^r7+eI!Pl_{_Ag9Zh(k+#B++|$RL9YyYA}ai)AD;=4%1 z8F6yM5Dlo*WMLxca(iuYhR4`45i4xvU2#SvT@BITXPtyasptVs7uP}!XSl{!vYJ?N z21@Sc%F=cHKm|R;nFTD}9?sCAeE-dS1dZQrDX&)9ZF}>qv4%@HQzmxIM2s_9lyV0@ zJ$wo!PsDwk`I9>MUC!u;HYH5+AoViJnA!)e|N|9)~Zf-O_>s8SMJnv$>le)>IP}7 zUEv0{Dtzc62gLKN3Ygq!LBj>yNc7u@7Tk{bSX-rV9*od|{gt+bx634U=T1+C9DG=~ zV~r^qwg6YB^VZF0M^~|bcUvZrbRuDmi552Mj#y)HUkO0XqHN;;I+a*ANcy^kxgXtN zjd>UL_g0HR->h>Qj{QN!9L zcFNAE0r|I!HIBz(AM9KTT5nX!SmRW*Jj+;PKsGOPhL`Dc8Y%q-bD9bPYIu`1XULj> zHVPuHl%?l1){vYbXA;`L;q41A-2vcdVh2dFzuW=h(5QhAp0SxZIm6|2D1qU-3nj%| z&2-sLdt)x%qJ<)i;xyQ2JDVHl4cXj?FxgxKIBJ$)?m!5tKnn~t?3daKB z!#ItzxgQ!w?=0@8&1yr?aV=hTl|4@5Y|@BbymMv28SYn>hl-=5&KrRX=`e=fx0JHm zCt@l&bC|?Yr}q@zd0*HJyvJ#T;777*w5m!XBMFR|jJuj$!Eit=PGeGJa){G7!~ev- zicGMBRVGg9t1V9B?64^cy9KIvgR~t5s6}LCX8*Dfcebcw2e7(Gd#vMt|slA)QjqN4OXi?he=-*4cS= z&d!2+wn~V7yF=wAQh`&hyr@f_hMR-v>jm3*`)GJT)Kk7*kQAiHX|QLvZVJH_2jmR3 zI2w>1#@4``;l6UI@!)q;5yuoNY=#^h*LFl$$7!s*Wgsl|ZsyQ1#O~1j(NR%s%H67w zOe3`q2HT;Nvh8CA!nQ;9B-{6TO0?ZBTGAS*O1Fl1B=SayO*)$aHM;)fZ*FD!t%#pBPJ^r>F;3%blIQG*pBkq@vJr;iC{BZG zeS3En_AzQON^Xz_qxQd!`(g9yPiMX^&9i$vRT16vg4 zH{=s}Y*Am3%ml2mz_IhvE0@L}a;MGXl^LSNVx zj_>c9&&&~dOCA?)&V-G#Rbgh1E@qA{W{%j(I#iyzNt1?i%9R(HqZzvKK64~z!k*qf zmR6~XQ?iJZZ?&aM%e}K*b%s_b#yw}kcI_&)$-P68lKoN#c|$g{-aKneKMS6)O-JiI zDC;`}od?NI?oTD9^8lV)*6#=%)fY{qE$NFU;RbtM`__HdLQa#~(TrqSenV1|(A*D@ z%n?m&ZS5z7fN=7ijB%q{)D>5A&R56|)oJwZEGRIYM16vdnS7a@DXi=b{=DuHu zheW2KNSqStXgi*6kSqsHk^E$*@XcV%K~v;Rd0e;;M!RnKLKp`s**d@AF9PuKCQxZfs8&a4cE zl`Qy57tJrz7w9_zC*m~P7w9_*rwjbj`URS-D(~_Inho4l%aBEBUI%%(QzeYl@>-cE zo_FGu4I|Dh%>LD$%6Y64=G0c-q&(TFPE%WKH*IUd-==UPYf7)qu0CSlX{lv$T< z<~fbTIgZmH&q`M~vGucWwrmWBYG2drM?FVfi+d(bQ`C^&dsaBY1XbGq*H(L z9ogNeDx9`)8tmQNEL~N@_Ae%1E|4#ERZNP|;&NptemEkhp~YkidpfHw6n#`F zeA0pJp!sT1mn6(Re%M1S^DAUt3C)&0-i?1T$HLyuwkSvv$a3=7qE5TaeLoTVmrT3N zX+x2R{w32cvn_n%Uo!1N`1|9+-M^T2xtMlgPiJ*WONuUmc9~%h<7CgY3tKlw z{{lm}<6pd)HJd7DFcmY{KEh?fzd)hKm^ZqIE`>FPzCBaggjNx0%Q0N+F3?{%5Rd(kbY&aGI%Yc$(C8!2|u=H zaaQ4rXbo`GJzmnF%InS$i{pwk-E*-ngB_7mMFypUK{N}5X2gqq-I_rZeE6isX>eDy zg2h2>eO%ED&auU5Tyg*tNkVq++~YLroei$&gjUZ1ec0kOE?EIYUX$USW1Pmtu}lXq<@TJ3EKOI1M&P_PyPMAbQFxk=CTVQh94}Iv4LNbb-?; zwvqFsRE+(M*_;b(PW3YjvcSN9u-kC7jwC6eS~9BN_!%#!xB+j^zQo6>X37&5mWf<&r#CF5MqD1ert#GcRIM%kVBG!_e z6gv*5!mxCGpl%Cq{meDfIb?l#T)6oe-pI9~PQs0RRBvK~V&1gK-pC*m{S2pc0hCRr z_ZKLPzk{V}cjDvx;H&;OcA0?(il$%qMrKod5jkY~Du#`Cbu-9|wCp0o z%rgwOk(t~e=S&!y*#XB@3{TX1nowYS-U{m|4BIM(pKY49eb^ml0^=7}F`QC&L1q)Y3%uTrcQHA_ez<8BLqeG$EP54F8p@;x z`{4#agtZ$3q7_wKFB_F{rLAH}+|mK@E4?ovO~h%O_Tx12>CJ$GTUpmKB!PJ+|Ki9x zHp0zS0{^$Ou%pt;Yh_&&S!jl(>|ZQyVf$R6Y3cXcK)SAF7Wln!bvo?jYJmz{)yC?C z(QE#t#~v8Zpex!$HLC}&Xg^|4-%$_7K_Wtg&1QWxY{O!7DA`dDuF)gKssM0ZJ!l7l zC?r`Gr-l#UM4x)G%>W0J(XAQ<6>sVa0;APc4lHJS*ksi^1s1b$Y%cr^RQOfZqp*W) zF1$ziue03_DW*htzUswUS$GkRMvF(*32P24=9Djp#5-Z714wa;L5ai`+esC;)ihoe zu!J>sdM(cj8FIR&kqv2WPfKx6ZkRMjX_`hhqxF?7#XV_sZm%pI*ALvV7t4dlUp61s z{Wt6dNp^-K-p{fwgiT;Gx=1WIppB$yWG7o6&3QNr8F%p0gUX8xJNMC|$7y&9X!!?w z!bZe^_iKn@U^0Pmg}{iK#sva!U!+iDB0Q?$IokMv{fmp@9OpzHp5PessJ_7(#~QI? ztX?0!qmOFo0rz4eY!mcGvk87L+LZMxq~RIbbTe^r%+u9cl`Awzw@IwtW19jd!{-I5 z0=u@Ru>>nOo9NQNi za`@b$2MQ!z<15?FgA0`0-8c=ptFO$loC_@7XzUv?M2+#3&3puQZl@I2mr|QmGKAfz z6gWdE09eoRd<3-MZ%~wDoRz@r-JO-xIE_-yZ6ujwnN-Jy!^kO$lXS_Xa3) z6Q|*q3~Cip1=aCWo}auW$$i2Q$bBS^9;zMN09%gb34!W5LI2$y(>+e3l<9OAiGi|J z$wQLbXK2+1X|Tr)BxwqNhYvlc5hb{_s$())XsT7f1-DD$q*I~=wMoJY~?B8homr~5GNJ*KMpQLZuFQw!&l*$TOvJ4GN8EYi=)U%8=e!6%WehIOs zWI4KzHOy(&I1N{?wz;Ncj6Efv%LZ-q(VdHO>?x^-MrgyX3;*!-50w9!QhCV2;B{Y` zJ3t&7HShsE)h&FZIA~$`?m`JuFs?7P3WB$|j928F!CSO{58^Z^o5wON-^kpEFxgzc zH$+dH&5)7kzPZF{@JNqCMb43&G?1iZFltTs-BFwdMLd8d+Trg&DUjLdrg+s=wm1zv z(FI}2Tji*-2A(ADr$JeFp3SaPC&eTWrsVFdTJTdgJM$Av;;59ny96;R)#EgRDu!%K z?OoGHV8v-L8F$h0qV%G=p9XoAn!vbJUs^Jga49r?CQj+AHBJMg%Kh!{$&phgMdVW| zKm{GiCGHbQY3dyM1mW1{tK^N%NDh{J^U7)*$|3O)qh@&0@@R$kX7laeV{UYc^B zQsD9S6>cjpH~Mp^#!4#iG8ry*b%dMAK81~LinG#wNwvy@gp8G^&85f>9CgVj zILG^y$JYy8>l?d&BkCy{EOJ4$I1MiGfP~Ne6g4F{x$qEtDQ-3;R$nV`GsDK{15 z;JB+sgms+8%3B7)(r;MIVSWb{=vkiR5t!E3OyrKL)fjAtPD*bKX*hx)+n%%dl0E8~ zwi_NKt-(Gwxp*Z0_V1#JjxWjX+nOAK;M(M#xhQsr3;G-f_A#!KYMh3yQFYApByMgr zQpC?1r$Jt+5U0V#Tgy)h-jAyw4?4?}aJmua53iwQl- zGjnuGV#3Fjm-|ABO5PTZ?_Zvru@P!6-tFR^M`mOc^3F&XOrArIX4x4S3PEd>^Xf36L*oRFLo)bLdUMWm) zsjTog?I;NaPuS+8z<^CkeTSg)IPK^O>5!#tdB>>pAlFu3G^wMpJE?CSW-XAA0XNuN zG|k!)%eaqZ>}~2Nlu8Q$;pEmD<3_cpE3Qbwl`P5$cTPLG3%5viP+S!m{M{fi;STw% z9$FM=C{kfni!>zI@-Ed;Qv7Yl({^Te$U#%&)OtLg+RToPMDvsPoc(q%9!!DpeAkIE zo)_`!^q{GeV~uOj6c@+|98$U+QfmzcW4W=Le-qn8lC8Zl!^FT1r4%`Am=IT+aEy1! zNp!U8%sL+BoDWHob5b_Ls-ykN7ah|V=sN-@;xyV9=v##o4}cLanF2|(^$uU4F;j?k zfGmphIND@c_{m`O7 zyO1HdT4eu{X&16>1&bc~mrT2mq-&?}cCyEYUeUjtVF;f;*<(jWbwkrG+{k}>yan2Y zJYl?P_b;Yh*hneRE@v3R`^T0_OuhH5zYmwz_)G@faDQ-C(ojoI-Z@Ungl)WRR&|UyijIV z&O719_AJUU%(DuPdX!5V)I1Bs;@DA}?jcFBt&I_YX9&mUS=_Q&AVfq+=z$qTnOS|Z z#%W-eK~&-({su$~n!!2NIE^EjaZ=w-l;ZZzEl#7}S?P$m1>jmd2lQc$)8OaCbdpkR zXmg0u(A`v{3RU6=oEDp55pA=5V0kUYCUHs4{Y{pzg1D|QNSZ*ik~?hp@DQh=yZPbe z7In`Gl8cpx6Q{u?ZZCv8Yn%o=R*%!*bjHmTGrXWjUq2YS_0b$Um4q?4bH;QSf{z z$so7N4V8$xY|)LMxwx^U^x-zq^No@W#ySd~uS*2&XL95cskuVedFW>{okJ?Iox(SM zCet}&^?F?RziF&{#Q4hzaXjiI=4A-bQq5QlRAg@?`bvIAkxpDd;Q;wQG@m9LQx{mpUFFGquqxzCN z%b4JW?yX`VUe;C28A|Z^F@YKhR0z_Myo!Mm+^G=tR+q|xWU_}1OUc*p6d!4AT(^2X z-y>sX?kx%x^rGusDtM8~{@SDV)KZ#4{XP<1!Mv{izU&TOjxkg;4@ENNwf()$tX&hNM>SmA^c@+ca z_t;>+YbG~H(iMhghV{6L;fZ<`6AH{Q>M90quf}O`(iIJbQCBh0dM(;q)!;i6Bk)L8 zaghV8uVU~FTE(32#%XZ>8lOoTu;Ru97yPVM%vvZV4rsf_H(}8?R!x)+II*_kHOgP-vq}dK?kQ z^vO4@-O#(xb@UQe%dmZ5^e$rZs(1Oy7cFAR&iip1`SfN$!L3yP0_XP*{)MsND%mXT z;Q#jg3tMTrRMthEB{>`{2kl=hZaFFqFqP`D9wjd1J=&>Gg}v}tTMKW}oi!QPdUf{L z1LGOVQYw^?j4eS4@jqeUUUt-jH&Ej5wns$?VNf94$+32S;BKwz07p)$WEJd0vMNpu zGr--ziMiaP$wzkJ>u^G%wr5!dfzj(K1{O0k6UdP^Ex_(WNP&ikuaxMyZQ<>NmEHPs z;rM}fP`FK4k7gg(4^SOPSy}WVYBPe8Dx{FnEJ2FWws1krRz`}m0i+myphVK^Mcya3 z*jCfXV{DTL0_;ffyyR8un#Q9^kRv3>nwODuYH5yA<1|2hLmOg{c4g*deWk0yycrwu zr5WXAKk}D4@yOfgLP&Eq9PxgZbs=n_m!BoE-~cD2X~Z@kxdS+1zikfVh+d?DFIm;@ zqeYL?aFbj9afIFohym}{0NB7}0^g$)-FkQ(SPi=>};O^Mp-y0ni}9rordE zY*WBw_`KC7PatxlThf_@dfJZmSW~5N9*D4UeQvlde1kLOY>A^diTwtJJI+{Y!)|=8 z6?a#KS22s}(k+s*C7dy9!KQg(>Oxw!J)9vzDp@%83^zDqV&!OJRRP=zRi@KtYGm4nv`j+GV0ws5IrN_P*wxfcd+E{<3%x%N5UrKo}L6PKAy99vj zkyX$exzw1EHO`pPI7(?UK})Fl99zP!cmZc%_U=ydYMh2q#zeqh97qNz2?tDS~ z*h2=o9jsx8x@=*j33ET%VhwxZ*(Ob+CNwQ8)+ozg#TtJ0$rgJ`y1H8g5ktckc)qeR zG-U9)y|Q#)KTsi15o=)oM)QnPN~30iw8^D*BgwF;l*-?6$lW!3mavA6^L(V?z?$;6 z+XY^x&uOIeyGK{{I1N|tK|^jL_e($@t@+K$89C z4iJY%4SdkFmRtDoanQo>-PJJ`b6u+vu!PV9vd`T&w}l>*Qh6-H^7Y}E&5a0?%=I#u z-Ri<@20HMrxr)=694Q5s>m}iUdK1D5)+YS!FiwL9xsE&%#QoIod6y7b0( z(uW*=RxJ~d)z-xjj$yva7Ipd2@8m;Zl9JtfKBRVm8mBQqD5b*B@FBmdA98pjzn$Ba zY&v>JS2HA1^C57Mfohv9(wPwHbgGTKd;=d6E3fO4N!XVwFHgBs;PKWCMwGDwMw1iY z|2b5n-@g>jaIvc+ym7$hu-G0?)us%)RGu*iPO4QNBxI~S$2d(WcS7Y2HzU#43$FEz z4G%TaI>=BiA;8LTa`1xeWhW%BNHqYmYWK4un)0qQ$>VzoW{yq z2Ex*BXk?I(2Y2ZHXnDkBdCoeIflgejG1v~Blx-g~5Vk#M@g)u5Gi|pMle7l=+#2GM z_}iwKCSp^!*5ss9tG0UPq71*VzHN>}%D*&oX>l4{-nE`^kW_x*2*)1Z{zHYCL% zn_wO59vnF-4q5&}BV-31N4l%JRQOsmVPQfJ!29mm=(AeX7bPSB%oe2?&UGggIdh=*p5VD*a>7Qk4?eEE#vC!?U15d*M^sH_j;6pIk#R6po|&U5F-KSE z#`_<=>`asn=tv2TQCOrR>$yTgwm6N+^y;+PW{C`b!`WRK2p-#;Bs?TF+%IJwie1qX zj9+(@goY>WyoC9FuF0NBf}A>SCQah~sibrsz>~{*Jp>)q7fob2?29Ik(u;9Bq}(9| zVjZRBZF6%cmhl0SIihKD6w?$_?p$FT$GA}}3JG_}nwTug33o`37&>6K=myCgSH;$8 z>KQTN4#@x?ThwVNav)ZVJR~v=bxn%D?RdIDvYc+`3cEO-YupURoNniu6oK0(-^}bx zflUsYnoK-Rro*NXG(}d$G#JfHu{T;oGE{>12c!G0!=>g;Y||BvakP%26fhy2pcD0Y-sF0hijo}CA@+X%m2W+#FVJ@cPQ+=nFVJ@sPBVFl*DuglSjl(z z0?k;kfq70!62`qMQGa!*%oCHeaac+)!nS8V7RFa&o$bhCb?(3jZ>X}4YGOCC2qRFM zjeF_pBQEj-YW7s}Lw1D|JC*t5oCQPVb$o_kNYCH~K57}j-eT(-fs~J3;l$5A*`kH7 z5S5XuZ*dyzHs)?7&W3HlR)y0Q-C|QX@v~1}w+5e{VX*DU;vLd8hLN2Y>QUjuwqtrv z3wLIOJ0(uzigd9locLL$LEfwiC(=0vg%j1_JGjE(YRXcdWRGR3FXu25gHshww{aS5 z_Eo(m1f;#Drat`K6i(YX4fgx0mO?@{g%dy9(lC+9@xyH96&JRogbbCG(L5rjpoLAu zd~~m7lcn}qj+B&RG_MwQNy4?q56k1RSy#3wHe0r6>tAMbEZo4i36cb|Nj|hF&@SW% zjQ)ikyOKq4gku*GYe~*&&a?}e8+QtC{R`W6Qird@auThwV6(n&`D z!hT=AAkr?RlHBnxZ2zSa&?9k?Oh(8dEEfObG3H37IZAR!B02UibRE5lB{}TC6#-F! z!x+(x(>U+MX@qluD)7s+2tblizKdsJb32#cOf?^bAG6b(OJ!!|x)XkE&*Eeg z?BFFFb<_eIYX#2&u{d

__H0D%QpbK=NYtiZpa&fzTE4V(5r%>TrZBEMkk(s1GkSQ*(cl)v_RN`8fk>VhZ!aW1I$CWT{C^#~0)u>_?TYi*1roJtC;F%gZTlptsjf*EQq$6~1n`j;R&} z&zF-7vQ<^E2q}6*^=#3NpDEL0D}82^KGh;_GWbj!>!^7?E>HCSoFbPv8(VO@nuxVz zKa=Pj3P(r2Eqvo=5}iZgxTYHvZlnS`bg2h8%89+P{I|zFG|$RVbH_yQvq_d36UwGB z)Lb&6_uUE1;jkN3^fQzhPrvPFR-^2rh-0ejcM9%Jyv+}@d6jJp@JQCeF!BC{-}I_D zjp|G8EMtNfy0?mfcqywG?)^?Nfs6#1m-85Nll?Ey^j4QjT0BaPmxJ~-&RsK?&8V%i zVUMU#nNWvgU*n(`uJ?1p7i69q7KLGj;ZPFxgDOzOWovA~C#r9=6YT1g76n}Mob4(6 zjIu^GDKbNf_)7=n0r?uI!OyZ42Bc_{IGYtmeH9}<1Mln344+|NbR?-|SjCiOBN<;} zoCen%p1X=k8# zfu4vK?eKP6@4_}wwn!s}U7n&vJG@=BsC|LJqQ=|FK`d#{B+l+(e5|EXIMR?5elBJ2 z!YQ*0g`@9(eopy<5CzCXNh`X>1m5iL$IaaQ zuo__%L%`UsbdZ7XhBE2FPD=q0VeN*1NJIHmw9;Y^rM`;cd13PXl`lGrCA;j$Y2?$J z0R^{G{R^DmJFsMo1sRwVykO#3byRw-UklnKSLAp(DF4D!d@N|{ejg+*JalJK$PQ92 zTdWPFBgMj^ut!7)&)AiyLrGQ-F2w)DfxDv~jDy54?=T{6peMj$^l9OadT`xO!&V^+ zS-&?CPS~b~<#rk1-XW-%9gtK(PSU$5N$)y+<`EC0M~YPe+4a7Ss9+?c@g%r0~ zQjBnMeZkPK= z96iPw#~Lk;FIwO``lyy3aBuI5Vp9v}N7Hd@&v?{@- z!RNhPPXUwR^MX`iNr`T;X>YX0nljELUWCGNZa9F7jx#VK(HkVO-=J_?7%!IEEEf7N z*2VB5>h|8=+L|Nf^h;A0^C$~EN>~_+lN;kqB38IxBn@(V-QbLgmBqx$0}`{(?x9cTPWe}Pn~I0GejgGi<8ipgOEpgc!F-M8~el6X3ITqu$>vw)r3 zEoF{|-F-Ht#u_f;j2R6!9x~3@L3q`>z-V*|0cP(m&eS-KQqDLhqNB!X5ZoX3k@WP( z+{S`FB}k#TjHP^&Mun8|cj7ck)ikP$vQ6sa4wvGbI2*a4+(+W*absCAvREa;em`MM zL;u~ynI5N6tO|5jrUVzbk2toH$}d}WgEZEza07KDq$wR~MwH;jDrl27)9SRK;UcB@ zC8=Me1-BzU)>bJTSxyQ^r@t+{CE?iVC)URSa~OAlr9AD7Kja?+vhc6y9-sqD-T zFo~m~qP5M*Ibhh|r^jgoa-I~bt*UO~G?OPnu8Sr6kfy@^>r~hNTUy7qaR-o*Tv*u z9n4q7<*tBAt3`#d9q5L?RlaDhI70>*%v3zwXQ;gTip$ud`o;K*4GO+pZA(# z3!UtRpcY31vXpW|!+j;+wIdUBS513;V@I4k8zdM;xxF%O-)bY=HhTpiufsU8YC+vy1;)w`7)C{ zCrA9OaT;8rod`y88eFT}7xX6Lr^6`3V}()sU&sBNv|&)mx&~c=@)twZF8@Oh|!sI|S^4*0=w5Tsi4)BiAqNI1B5+ep= zi(;oJ7P-5cGGfb@Fkyg8d}L7|iO4ZqEpkWAB+(%eiQ7RIYpE2@3lccQUcuYKH_jv_ z!yzeKq!zY?Bba(9n4mg28P3VX5L?!$t$H#X?x^0NL0x*WnIrariFgS8cpnc#bE21N z&ZLV8%ljubpY*Xi?|)NynK?S7sN}u(_meYoM6T1FC$E*)or#$vHlNWLHFLxUHmbb9 z9FgZXRi2rnQ(}(DJUgg7I}@h^qJ(d}Do)9w2hFT+wWUkjz2|L~xa7f{(5ZV)ISCJ? z#&beH%RCg@#bg-YQ4$)SPWFfvO2lnPw0-OU86|{6lN`u zkP$bs{N|8Y#*7=WoB72Z(Zt3venJljC!yQU7if03@%2q{#pQ`yK0!LgsnAaD!W$%W zTopI)6T0GxM-0g2SVsZvd+tIHiA+P0D0gQ_#9Atav%iL293xo-qf$7-09@n4!hJBB z0%M;Ub(S(EF&)MRG2P8O*0>{?nJM;yiAaV@@II2c6PR*1<=7^2NbXV&ObqPBog(KH zB{}5Cy}yJ~oCaM-M@w?pvn7%R%4S$~v|st6Q~CmZN8m)9M*9MNN8!XSFZB!b0ao%I zzCfSMxw3^z1_a2<-I;7zdgshMO7J@@Wf*amVfKIVY|H(gFsIJHz+BeFds>mYI!$ez z9MyD!X$ss6lxC~Wu0G--4J_y-#};7R8e8y*YW!NBe#(L&HspSWV92m;l`$z)YHA+YHRCw}(HLpB)2#WM^p z_f2ayzT1bvnHTC&;lvgxR2Ms)_~|UvwVwv>i9SifRVilK&%`ImG9_7)&hky7DV(SV zU!)3WzmgY)ljI9gkJI4irf|BA(_oX5RpG>Uk*;v!=caJl#%Zu~NwaiS4GTN@*+%q< z3?yarFKkZ|(HsISRMNhGVfPb1LH>fxO8UGM%49q*T`jW8#+hhI+?hTqus2DzC^lQR zXzO3t^h6e!V`2M}Xwl%KLa3uayD)*QbtIH$hmvg39?Eb13;UPk!g1pNXtKxlCb{st zO!nA|#kC>Yg&X0=wXaFz{`HT)^Yq?dzj@pK?P)Eez~%E)lN2y}l|kDWO*YI5c?pN5 z9Vj@D(cWk{;Xpj2>YfgcI8J(5NRe8?b5&!=YI zv&fUX`U!4eZN4OFKaIPF9h5t#c_g_VcJAEcH0qrVuIPk{HkRZiJx+s7KeEjcL9)#; zPJ_KfY_lBP868lyc@0O953CO_)xGBa=CxKOae1u_s7aKVA0Fd0&Q8a2HjTpuAnoBR z{M+kTxN}ChQ{psuE*wvb(_mMR^wue6I6t`b;kBil`-Q@K>$)kE@)P$s4YqdCyIWer z4jkTtO%B}y-&dPGnon3q#*GHVNi{Da7q9Y3>L z6g*!}GRTG>*`kQ**`gaiGu!hy(F?nz1wS)NG8pOz2B(q=h{0($|uqozqL_yE`a z%rEq&kSN4yR9|vu856wFomC7egX$`V59{ybYs^rR#%EI@`(L2To_^%#$)#WhN7ZOw zirvg`J64jGK`Z##=Z-q7`A$c)*GDHaxmeTBg`8?N_tYacFKV+%g!nR+R$OIdHh>e*5wr@vD` zN_6Cez6&fFn^`0*hn?#@Ovdtvo5m_r9o#_8B&wszn0fUb@3LCdIVlQfi*mE=@OC%e zrA&?;kpzGRi}rZC@-62KtI@nm=k4qupe~{ZPQ2X~<6|wA!s+8FX_i9;Zwudemz3Fs znQ5Gl4GMSfVg`jBGt?Qc?viD^!H6{P0^=9Gi+PYUJV+Ni9^v)IEE(6T`H{w?Rm{G0 zkbz%}GPFT@^vO4@-PpTWQN|I3s3Q}psCQ8?Ic2V!tV{MlQ(o(VdVwhBsgmr zy;hw)PD6Nhfv#x7B&!E6FcWc@?5GFhAkoWnS06>VUtlphl?c)xq_Tf;gabI@;C>pwgWrXRwmB|RS7l?KJO))U_Dwr!RM_u87kmB0w$CBIn>iboXNZhshoET-{1_4NP2Q4 z3ARsO3wNBc)aJr2*wlx0F&Bqdv7thnx=ao@ZXftKrOJax=~EXnJn!L5CRRuU9bPZD z*A32?SYhu4KN7;rk=SI54mjggCVMB;wGcU%(^{y-X>g5iQ*+!a7xV3vrPKO>3Vv#1 z{S}vw$9^g05d%e%OYIT>)MP1@SIK07-p3g;8W$-|z!^W?*;_6!8f}tC66yOmQ{yy@ zGTwcXjSM2n5LNvDvv-wYQr%n_m&IXmDNx+kciEdP)}qDTT?)nB-J!VqLUH$_m(4`E}Juj|MZuK;* zL^`1^q;yThLun+ZY z`G5r7B#rc4!2n{*!8OGxX}m##uCr2PvJ)1RvIf|QHY;Bz-p!_gZ;IAgRr+4U#TUDe zElL{bomFk%rL$AkNQfp}`m10dA#r4gRX89Q`!xwUsj(}ZvPJ?e;fR+~WQ}z0^A@Lp zDhApNHJC6CStD^53?Nnv@QqE@h%A5hm@s@Te{{s22DxPUbIKZt(1be?njqkSs<)%l zrja0CxCnt#Wew0jZ+3<}?b3G%r374hd4D4#d&ztq;Q(9oXL*k#lQxaaYKvz{StH&+ zU+T^@IlT|A}8U%a-4@22%kq$Vx6wEbcF+p;wOq6-JYlG2tAO(M~{38 z0?Ox2?x_++pRP1-m7>{r+vCL%RK#h7%Sb86a z2ok9T?Qt6TTKourOFNK&7QG!v>x|_>1N+i?`EV1I+}3J+<8`!ZSGqDxm9R2S!-iWb zfTf^KVrVF1Al0G3X;5B35~SRjoMlPWxx46HfQH^EmUGv#_+sw9c8WT8xdOtq23G^? z-&Qz6ph?KrFXoC!Bgc}%LD073B2EK8$p%gv2k4_}xj5o9B4or5PK$hy&}PVv?}*bt zUT-x{13$?I;@W7B(?Gg!97At$8c3Uu>7bO2><9e7R+b2ifT>{%cN6*-$@66s{c;iv z0MfU)AkI2c=*zy&i%BprM+Iy^x_kA+Au+fjqrKY`a@A#091I|ypeDsZo^a3Bgfb%~ zAx3rSh$}O46-I87Ob?nu%5fi^RbeEbtOrdYcX`~Morh6EU`Vqs zHj4zFkd|Lissv3T54v_3B{CJE1x+Cbx|hQ!&y0i|kg{Ii{lFN#88PUE@;cT@NfoWS zvxV(U=avvki$tgyFTzY4x zgtm~*Ud|Ts*T%pWNGY#~5vTVNve?U}_Ux&_04=4jOO)Vvm|RStNScQe60(1V!;lo~ z((?xy=H)ZsMTMJzmB^lR zPWW2>800cxfQ^2&+qLb}rri(iT6HbrFEEdGy~@~+rT{IH%7h9Ly7LaQf-8qpzo@Go zPDmLq*d-lzA^`MQ{f2^pNjf~7@U@&8vP@ayG?0p2@olO#jNJ5##8@IE z#u5RH#eoDOxw;RHXPc1|c0pcv`;4TN*I^fAc2{N;r+(347o@@V(b>LfcIp=iyCCVj*v=Ami4d?0d_0`63zEt!`bGRC8>p}gQgD0MFX#cGhEM%P7*ecz z6^mB=f6pXkvT;ZU@1F!Rjf7oXP1jjie zhg6sArze96)_^*FMlN~={%CTsg5+WW*imNWJ#xe`Bd0SS65ZQ*1U-Xc^%c}y1mn;% z=&XUv!9eMpo8Gmvi=IJe{VW2ji?g#`&!Ds1K_bdY3?x70ZUW3>hbxgpT+=h63KJUR zj$lg^+5%B2-r#+qG0;}@0PLa~?I8g6d>HKzjB-)}<3L|UJR)1T5Ai4pg^BjEmnt5S z>fR|H`^!c_B2L4uOFUd-6fNX_i@`(SX(s{!1^xdeCJ?(xZeqQ@H3sqd1{!R%5~Tw# z$XV}|jZ)Xh493V$4vOImE7_vFS3`n(AIua%FQ8M-R?A)-0?^Q*YaBNWxLisb1h|RS z@SA5F#h@F^ zKG9h`9T10|m$oVuQlAJP2K z?jaKbe9kGsbZM?4KvpWGkP!KR@f-z}3IST@UBB37pv= z#uPXME6!0NK9D;*kAQ19dE`Luq<+z>CL!^^!XsA|p>oOES;L)?eBDQ9H~pgLc0r2& zxH~)bi@;F82dQO52@Cu!r5L<#9Qp;ApA)o318J+k(MHiPK=wSpKDG`immfh(ntt)^ zJcy$&7fQe$aKu*dVN@FMArgg8m1r+}srWE*{5wVBVA&|ZFPA(kt$O$FTB1!0`A@%g zojP>v-2xsCod&iO`UfXa{oUnA0bjA}7~p+rz%23741=R(7FjR!Tps#hdD)5N915SoI;t#A=HXsYfLD4)2;SsOz zbsrvt>M@Z7h)^OifM6UvMI5FfY5{l_Ca$wmM4Z!uQ1U~o{EVe&9F?c>!3j@hD6HG=HYo}1rf{VpZvVofG9;FPt2o>vKnn54pBVC}24>cYp#2wOG z%SYPXg;1-&;n#>Cg^)u^YCGDBl=GUMV)ttlZDdPFu6tzqcW-IkzSal2h(Q2)zRYZF zDKWkaf`LK;K6@cuw(O;QUZe1U*Iw{z1T6yhz_aiv1EcV_To;j-#g2Z22)xgerRtpczi z8pNlcYFAl4CxF3#;49n_hCWl1Lauib8?k50V3`Ru{}Zs z1ON(g%0GeZ`Ka?@FrnE>H#gUO4-g@NR!+;qr>>o8tq(6QZr8qMTs1(u>1F9MhIR#O8#G;jsg z4;1qB^>ULl5?sLmlEnbMWUK}%Cn$18-ZTS%q}bxtLjcvO_aUd836jeWV8z}Nu$)^< z$qJDW2*F%>JSR#ccs=`828zoF@-^#58W1{P(kVdVGy#A(G}8{Cowr_6&PY@P&%%-+ zc2g{zC^n5x)hYqH=TXksVl}K%MuHsy0_}xF4EGf0th~X3uCr2Pva|k(6c>7$Ulcn~conCtfq#nfMlA%6 zdjCPz=*eSHoFVSc4x1)G!ZZOQi46E6<**2|Lg1qFg%8R!fd5WeBf%D=URGs|bnx>M ztAW8ZSAC@ZS7sQ8tdZCY29R^I$Dg9Wf+A}qFcSb^#wI~fPQhJ*;9WE3Ar~HzQ;wD% zBFL=U*ZM#m^Fpl!ujk*|OIlY*)vHPk)I0FqOT=X%+*9|GMVSMyy`-!WnQcLlSU5z> zioYc`gWSnZn+CbBJVFH1HH5( zi=qf#H-IWg!u*nU7a7E*;f68f^OaK@3YgEMeIRWx5p#qmLxd~w2Pl2udA^SL11b3Q z@TX8fea`c!Z&+2txXf3cBeirk>H>J4FQPTXOvR_jx`n?Ls1!vQ9E<4hEm{Lz}+Dj^a1Wsv( z1$d%26)RU&y2T(nb9z5}v<9va4kJDkP@_lXif|2B<_pi!uSARBkBXNOqGo6hNe^Mue{p zx~NxUlyvZ7vcfUN&OEs2*zNsLttJTu)2!}LpLnAxXe8$mSEbVbHF%>FS3r1`)>#!8 zuG-(t-j^^|fC#q&2~w;JXCeG8wuv=dBNRxa#yn{Y!-2^=#;Ro3{z!nCEI%d ziCrq7r65t_XqXB9mYGKIjpU<@jNx*Fa0FfR*pw0)#k1 zQea2CMkt_3XI%~PlWZWaj`nyBWR1r;^b)Uu9+D0bKa0PCAvokHz(*xsL;NV4=ob{! zv+%5f9VsisVMhf&(HFwqaoA;<5#j&kXZ=Z+SAZCuJ4n~j>QR?$mj(s$3(F@cmy=K+ zksb{Sbwz8Oe>R z14XD&q43CcqpDo8cGfN>C~D!Ov#T~i!f#C=0%mlsj_Z?E^hD?m6e_> z(Ib?vh`Y1nMu&8}f?*OpLiSAgw-g)-dK3yIQe&PJ0UCM~3dqqjdL*<75;4n9>#Ep2 zPf-rT`pVY2rKLP`DYu0B6<$TFwU?0CNbRc&WMGZn1hVWUu_&Yya_QwAf)=)9psc`z zXNkWh;DZ!1iY@8GuTY35*-O$0NJ_alseXl0ZUjm8E2(#o_20hHq}dd6896<#LbV)*7)n3-K5`hkNiq>%B?gc)QarChedI7i zd&wgOib5?4JCEd%(4i>QsPJ(lR4!RNYxo0-TKMSfb{O@vJ1BS&cW38elrUJJ2%ZW; z7Gn7+wT==z6)Izn9qs+yo+Q3O!t42jGmp&I}T!$?3?2Qj3eD zXolxI6yi0+Ldjlo-)gQosIuV%48(hAI2P&ep*?syY44#u_&OnVu>CzWifsIs_s}>E zf+_I&(XN+(FnuuvA*ZsJjuxXnhu2spxF0o|Nfd>_D9K3|5lC81^ew5k`0FKSq2~|K8Unj6Z zj$mcD{VfP8c35xw@?atw$KmVb7P1k9iWoMBVO02V`a0ojr>~Ps%m(_{ecVv^!|CgU zuLTt65r)vUACfFsq+z>j* zf(jm9XXK`1B+!C{*K$URYXT~W#62S?_JSG-_8Cbvp<^$o;GxXO)n;|mF*^2w!XG|5 zyXhDmdqEY3xH~&_jKp4$a8ztZiM=3utl*4-y`cUED0S)>iM=3yi2@&>vWPDoBY^`X za?->p3W|6Yj8Yg%4|VjC(@|c6p$3Q#ktm(Fkb6k>QeT0gCWuoc_LcO8`yuELzDz%L z=?+(?_H9b&j!s$i9~Hnr4(^BtMD39Oq+LXW9g^Yc24|S>U}PLZNfMlhRGk#&BN$m~ z7m?Z*&yuoKz8pjBwNtyG-iZ%ZO4JAKK$I?VvQi3C-Bk!hO_V*T3sUU;D0|cf zf*j@SW*u3xCc`wZM%R&@Q5*J=rGSOUkX3u2)`&A|15qjJEi(-8 zq4SohkE}**AhZc{MrUvZkqVAE+kz9M9+YI^N7=-l2^?igtTI=S1;=RZ;TmA5KU|I4 zfETuWxDcoCakf7k&WJKHLU`HYYxyJIQ(2=nQ1wEI+JKLp{$O7#B`UKYhTtQ0!S zvF$1kl}4;2S*aE_AzUG46Z^FeHB78?6So<-%*h{}jv%N{Vwa*q?7-VbpPjlUzBLUT z!^9Irvx$n$=iL0j2FH{4MnB@Rs^7fPfe#?w^sC?Gvg zZe7Co#?}#^GLQ-OyJCE!kcd+}4wQ{TMD)2s(|EYdsEWwLT?PiF{Vo#*Xz)My)Z#2i zvh)ics+{;_3Et<+Ubp zS8=k_dCx#2*%a^fWu=8Mz=YENjuLZ_m_QJ~5$!G$2H-F5FbMGs{3k!9um(6{z006( zi}Nmn?5lDD4lhSkTJe0BL5fw$5y?}y+}N}cKtV^`Mv%NbBHyXYgn|9VJz+UX9pnhp za6et6sCOAWQP3%UZGzxyW8)<-KF3`KUkf?`W!UdBfcG5rHNYqVxZ-Eo2s%l6bcG2x z&mOlCgqR$0mkID3w-E%q!lwFx@JUT34s7 zPKCjlP;oU3(!FKEfYUj(zAn#ID_t0*7*&2s9VS2zyIT^fu=tSY))`5c7;F4F3X*6u z0!*R2{zEMmuQPJfFhEH#FBzj01V0&NMm`Wcuukv@=!Qe*IT1XmVe~3WVH$$xsw7k{ zSvza*QBk=?A;-0|n}z}Oq47f1|6V#fHH<`3kU~}RL~csyvlE0!DTjsu_NQtXiH8IM z4^cD>5I!FoMuH?kaFo%^@HiX^!zGrEUV`y(m2vnSgaN6 zrT1&0)G^Z41-(FZxb$2NQXO8xfb@B1s0@Tbd%;iuC{DZ#^=X_8b;cq{nTA**$3fA( z34;(3q;6Lq-h}Egk;KcuA`Bp>qJZzH4dYpuxXwxuaZYbSD`mQwrg;%6 za;|_us)z$j=kZmd%d7xoIC-T78w1?bdF4wksN9l51#pDM=&z6qU+m5)L4Yc>&c5J( zH?vhjBiL!;kcmt!_>`wudXiYpR%9>Igo;Szf@W$4Gvwm&B2?sx>Fy&aBcsSneZ^e; zsqq*&Cn?h5NMhB#C|u+8aSW1R2$E~sJr|L_L9FQ({QLB%_^>&{Fe)gDZC&4h!}e2`Oukyv2Tg;FzhI z33}WUd{n&d5%u&)w#(QfB*etxBH^XeOn@Va=ToiD38^lXIcc#Z!MZOC>d|bAm^kcnZ&E5ubGomCgY`volinkAk(FK+EEurnUkvtp>oRFS+hS-EyqV^SK|fw zF106Xo`0@Sc9k=R8!z(DM zqmbuY&LHEan3IH8knYkhX9QFytbE~Pgs;VXfTh?xq@kYOrH8epTy{V!j+P!G$*c?! z7kXr@)XzS!I$v+%kUCZN((RFOIQ;qSC8M*DYf^s}U^{QUgq*?LkW( z>n3TW?TrECtdti%DDCrgvr=TT6C9M{LQngPsyGU-Tu;2pE!i12I(s0Fbd|n-LCAe6p$g=DRBRT&d>(Hi*7;f&10r0BLintBrs}!H|#T@cQDvJng(7uTNUgI``CH{lu5j#-e zKMuQAiAT0~w|iMS<2LX_*%4^D$}J#=zPTm0y$qZ}j!At3izsF9RW`&VRN^)?UU>;r5^)=_#|4LNlwu#*{f$vv$l-5n<)3zIM!g|r`E~ph+tm@b zf%KMge5vx0I>&bZ0_f<7+X$9PagcS=X+{HeR6>oBa#GU-fD@gVWSnXQaFVJFj3sFFp&I|8wxm*Lyf?dP*o$+V+T1YRW$XmfS{Mlge%Q3VPq?s*JV9ZP>6&1+4m$@3r0Or{h}%FLF<^+;O1KUPF)KoV8gy*evt0{Rt3kAZ{M%3R zNCG{+Wgt+nK!#MM0ua)pQqkfz&`Zo) zNtA=SA@;*AHLIQ|2c?UYMbM2(^{C5Lmp(ATGEok4FDeQe5y~D73MqbUGg2-la1vDB zCn2Sxwx@lBC#yT!IwN^=l(12-nj{Csh~l1+^HLH4Y*E!FI#7gqJ_?UqWtz$*aYT_f zQ6&eZe|&Uy)h5WOsCQQ31=lC5(o`M`2~y$t*k%|=Y&Hp4g(wMfP}oQni1Kv7Lm?8VX`16qM}VA06BVQp3o-9VJSbwK}3>7+bdg?HL||4 zwQgx4&yXEa6Ri<|8NJze)?PwlBPXIuFYi{Nq*h{4$d%~Q%R2-uY)Q;behTS}zVf+> z$47)O-$IBd*~?0!>hTew`WA|fhJ1^1Xs6i{0U%P58zH(Q4hk+Q%k5i|Atj@Hf|Y4X z8Xgp7@-8>*GjjS~p)!*;BRvUEgpZ7du96H%8pVtxG!>#I;Xys8IA^3lQ7GnQpOHKg zdJq*VI4L}GJx^|jQ4gX*jT9f9Rbga(vMvTR-$2(VyB$UegCUut@!{Dh|ywvPCet~nu^YYxgoIROLl9vVOE@1Z@oI%)5rJ@`6>0A;ehheq+G|MDIh z2}LDg`t=g^ntU<^J6876(c%anq03T`$o#`dpa`mE`Et;%)VV~>DIW}#02kWJ=RyKI z@LrUTax&EEhlFH?qA5TotItaWpi76Z6TX%|2I{Uu5x^?rb{&#}T>?rW(#vwh4*9lJ)TT~SI#wYIT(r+vR$fU#9315pJ0pQqtoluua@pS?l#0WBp z4Uz}iZu=3y9v!|;_}b~~x8eJzD}-j8zFGjvCS1=q@1hNF%p77 zMnI=AqsJ5@|DR^e_=qW1rKc1x5^})+k{gP90?Ir2km$A;U4T0}SSgP-P!a`^L5(V959)#xdq2t^4S-SQI>;r4n~>Fc z6IG|I0!@yrhcrcKUX89JJL5KNBg<69$j)ci9zYQt)iy$OUmz5r@;YZVZUdu@gmKwK#(d5IeL)a@?pl*|u zVIpn=srSV0gbR-?ZUZH&l(-G#&(l^)izL<`tXA8=cyesHQjU^MD0O8e$x5{}-ldZb zvgO$s0|d}fZ3AD6xdGF0T$=&)ska$GHtIgokUd?lt1?tF#d@idm25%)?>JtZpq`Zv zxo(@0qWkppgGhs?&q!S;eMWAYCc-)|xmbXAlo@$>VJx+=ZAOajb2$u)2R4!RNYcHoz{>n#ZS51SAcN#B%*f}??>yzDXGx7xr((Yj&VIZ-aBE)V&qC6+G z0CZ0Qs_~O-fRu2Km&GVLvT;6c+a<{u}(@X6tdMR;!!|)o?d~<+l(k9 zCfdthns^k62<@6?pllQ-<2LNd#KUDqRYV@{GDvafxXS<<{10&(I16Hez*!WO8ewO? z-$o;WqwJ*tFzDs|S}3#)o+3YmgbbIr31pjduoK;VTn<5@>}#E^@-`XWl*^&>o`IBf za*DRnKWJu)HHebq$i)OWqQ-5Y&DWH#5YIryI{7JX$DwfLk#`yNhnN?yJOrkWujP++ zFGu{UdmlOOGRO#r_HgSE?SX5^y9}Dc+{dlPPfl1~9+9Wbb;1!}d6z*BH|205Xizl+ zb-}LQW$;8niUbxy0d+`Sh2FLr7DA)oYdLR30?K+D_!W0H@u#1aQ*z;Cm8DQX6Y|X^ z+SoitOUe;K{nB^o+M>N47%NA$MQfr1mQR#rDlR*O`K6T|bfXKo*&H*pg$xGyB@B>4 zA9>jHE$$;OSay;-LfO$1iXo6wF~zE5hN0oaKx(SEqZI zPQ6>mxDoMJ5xe=E01!=)w0?8l=;#p41|Jq3i%9viRD;DlBP{sI(~w3 zfOZmzR)}HynpOaZ;uxzS#?l--GTEQ#mo|RhM6VLd!tok0*-?Ua6}}`jH4B3)l<&-pcSwA?6ihhX$`T{DvqV3RZRb@ zX$=R~5XbyLMY!e#;Ef~t>obq7K&uW-d!x0cU!nk!x!$(0VOC}NW{Ms$ROXony}#yjzLPB0)Y`$H^YINmOP z1vR=~ut#d-Xj+`E$7X`)a7~c-AR2~`Ly*COmx2uc+5{O)efvt|t;p9CjqDR&vxK64 z##iu0q`r@<8Hh$o2lW*AA{sG|_LauP!Ai74kl`L;Cpru{%f5Q05$UEr-{aa~2@&4# zHKza$-*K`~5Um#u_d;~~w7t9aY}2v$kDYoKZ`Z$jx1PO=_$P~>HfcFw$p=?SW(P`& zq_=$+QYD!k4N_W$DlJo|jW2tLRY?Slri{`mTxsQeN<@`pcK3E?+u%}Fk|nS~Jj%BL z9={Fr_-&BKZ$XnCIa0<2RSE4+v&}y^6!R~W{8z-@H^PT#7q zL6watg{n|c<;vsGAXgrTs(a4Cr%+JkN_(hSA8mhydR(7S(BsbIP|)Mfr!dgt`kb%; z*XM+JT%Rzkk2{aUK#x1$!a$FQZ?Q@qzV*0DVW7&L$6=t#mB-L>?98|f| z9xm2L<8iph^$7<(?mP}xrGjHE!$FVRbHd%8<8gf=us-fQj_@cIB0!HjHzPbsg$R$U z6alK-c^m<%TzMQJWSu4}BCtO0JdOZWF0>=-p$qNEap<8hau|BxF|rhTYzunao)hBs z98c?mG=d&@jP!vXxQTp#p1uw6Bo&bE&tqFq<;G(q=yT&SQtY|#7#a0E(H;P*Txkz* z)kU}`dF%^%-1afj&UtDZ2zp$f6By{q<3Nw=6NvS3+s8=i=7~drpvP?=BTbvfZ#}M3 zAgFTZaS*6-*~ei)Le{DKI4lV3TzL$8!G*`cSReQK z3yZ;(cGwFpJoacSAa#`|9>ZGj^c>LR_8eE;ip)_S`(l0Ec^u+VDtNRN!i25h&Us-g zc=#4nx$k3PF}U&=_JRwKVJ&!|U04jRw8LI-;ju?s0cnvu@fg;Ehvx{3!Sy+?7hIkr zYz6myAZ!J99*2393Spqfoo~Wc@bIn2Rf4tPfycsPaOE-V1s5K}TJSJ`VKKPU4tv3c zc8|6KGSztEaX9F4=TkW7aeWT#1()XtTfx1ruoc{S?9o;T^JpuC30uLP$HG?d@GYov z-^apYaOE-V1)aypqvB!y0GQlpM}`y^+L7_ZV_&R~n> zL_F{qDMCDO2pK^;)$8B#SWZKiW!d7r^3v0mx zkA=nH%466IF8f%+7VLKeI0?Jc9^|Tv;63Mw_8`#X&SO{$9(XJ)23H=#Ua;~wIzh7d z_#gk_z1(n+Jv_?3lRW<`e)rYJ{=2s}>hGS~xc=_i#?$X=8*jhc+xYr@cvLz6jwNih z?3FjxMd~C=3HFpq_~H%SH!?}=w4*n8=!4q)my-#2=WYjB(RDj`%)CYHMIwU@*`1a3 z8r{m=6-IbDc!StRYtzED$W)O(;wMWa$K4md+@imE?71<#?hi-VwtU-p$j)}O->H(! zXmOd|niojpA1=dteu$(kR{6)~>aoYH@3)=Aznh5`?An`-Sv!-~D^eipK}>P_=Hm#~ zdd^weZpKIReUI}ram;e^>_$<#^ZwAth83og`=>1Q){a@MLSR;2W#)S0W%_;W;oW4s z-uL~<-rrl%=f%>p&I1g7;%sZ$_DvanU|xH3&faUZ?CoLn&)8+0n`!8Pyff(d8;429 z{wGNNZ!4Po?kHLg*-UjpwN!IOf|4Zy|I+$B%u4tClan_okckgUc!FvCarKrd(gExd z-O9|*v5eMw9K|dAdX!|G_|B}6Y9KE(x(cuV;5U}}c`6baaFZX(oQtlk(v+nc4p++tRoD^=$dzT5!>i* zvnQb)o7SgqC)c4PCN7I|jPBw=UZiERw1P#+E9E6A2kl^e7N_C#SW`~GD6(UyA76Ck7qjT_vMjV&4w~-v z5wdh_X?D;^ZcI2D!YZGhMfMGj&o_6Oz^cE0Zyb&;Z_X}D*x;4bENyFUCT8e1wl}5@ zOHtqzN!znKnfP)M$ym4|dzor7W6#Tz_CHQFvWHY9srF}K-T(cLZQK%VX z$bSmas*?+tucstnL&u+D@9K0XsZ(c)2^+hDElPcc9m>6qC0#w!Qf(?V?^nx12Mrxb z(_PNY`+PfrUMv&F6D=LZS0p~cCf}~Zk59YG#@E0-3}U#F&Rmzo-f$9PZqL$ z*+-MJHFJ?Eoi?yEuM_Y@rOsGJ-Oopw_zyQTm0x9yDxH$&ZS|P3RIAANQKyZv2VR=h z5C2R`?#syegRAC<-1o@2hwn&*O}k0ABlFC~U4F5cq_J7?=?rpg&3sa6)_PKUeh~S4 z?`!fVcAc5<>LV6B@Ni`9;QFjn>#}Ti)DiM?<11!tyXUMN!hq~M(8C%1zn-$`U90{Q z(!)Le-L9u57@3SuA^)Dz)kL4bf!i_tp8V)wB8I$4NJlxSIccydOWXh^ORKyt!oT!+n49O9Lyfpf8>e$268nfH0b4Rs5F@lZ@ zD^5$7KgF8-+>qy7#aO*;jrkwVn}0h=bdeeEa9 zmJ}x~O=Ewuv}rLt#<~~0>8Uy^>(NX*ThJzBTF^IldP+UzDetmf3r?K8%E)tgi&1m( zP_w|GS*Bxjk9z9w;q54$J%Z*MJdl3rJ<>V47d=&O%`kQ;G?ITTy2@-ksFV?QXP0rO zL?>2zelC{lXfB>Rb}z|OECIU`z1<|=ePok+j^zV8ePOKW-e^V(9 zNqH;AX#Qz29k^yC8$bIPnF@K^X6IOA^5w)vp9=xB)`p&?A7$)F_N@Mn|MgoTnk-%! ze&*m<@@8yDeztosPde?QvE_a=xmF^r(PMffd;57fn;O3=%Ws?|wVOAj=U#1)I!ED> zzQ=8P>eX~+`X4k%#SyJpwdFrZ-DdCOw#5}aH8$oR8J6{mdF4@3_HxUQB+2XgeEGUu zyy@8OY)$XQyiKxb64>^8HtxVWqj7=n*|%dikU#!4$>OVdd0~8`RQ)1$bO~9 zlaUh#u~oYUliAxQk&tIc$c*B9$nb^h$ii()$&t!G81wTtqrbeKz_#AoVfpaBzESM% zNAq6)HvHxde?F;Dcan5{3wpE5TBCpN5LTn$6gD=L7NzQnMm67CpNH244 z+3BRf{#oqd`_tyXv$m3W4gVqy+h=AK!_ShdP2ZAwwOf!YYdg>s<+9P9{y(!-&7;}O zyTLplWmR_nd~;TDx*vI%H3j_`HHJoJ9zpWfy-F%QKFl)Tj4(!3pUldM?A^YuRKZqr zbNu-#=4t=VM(IDMTLPAZTeqj9pW8w1Xune(U-{Wm>TIIeZ|f{0?K|hACjuAnLLCOt z8_##JeodD0DVZ3r5}J+n$y1wlJa&nn_Fu>z^o^x&;$P=)igu>@{5&YD;@0^Uc=_JQ zzW!ij#$PLv9e1cv&blGov{gxK^T~6s(ySA#$;uDSB2Heda;`KWFlKJw=yZ% zH3N@-svFxneKH%F=YA`Xt|nn`WuM zKC^~nTapXacoA#Qv687AvS+aa(>{CKzgEI2e#t)YBsvoXd^9O zNABj##0#3GXy1Q!m;*jHple?x=jSd|BDpp_Cg-{wWkqT|B_%EoqpP+DGuCGcTQFlB zi>S1l^~;~r{3hRKwx&2^sT-fQM6Iu4_$4n!Ydk1S2DeDed)D|&#{D*dq}r5^2A3&g z3A+iAn+K;7UM`fVR`kp7XBX6)Dr~31xH{-L0t1_^&%Z`!fX}6FhOK;PgO-AxuN3S#g z5$WjLEd}_6Gv63ZSXSC|;C;)ucRBdtzq8Yfk3X_I9gETVowCsqMdtIc(b0Uy;w5ZW z^_jHe`8)hc(VuC(H2!q#l~DR>LVxPlr75pFFA2N$c!Jfiv~$deD+o=#!M3y^@E<%l*zMzqmVZ`eYG*c;Fd}nERR>t{cgpX0ed4${|!L zEywl2ZB#{Rr9Rl4AAH=O4C+6U6VdvW? zq5IbO^YHzDMYb!Og3n)`kwvwcO>TTLNt-mQ%yHoXEPcJ!Wc(u3pAu)R?;$de)AEI3b}B|?Ib&t)|ed+$fRl6=!K#`lR<&o2x+*LJ&XOBl};GV zlHBjjPPJ=dp0WIC<|-aSmR-2cM*lL_2#QKWy86Fo72-{EnV*M)U{fj~duo#Wzio`% z*<>c!UM8VcX}Oo3DPmJ9ABQE4+A-`seZBHM-T&a8RjGKqW7)<(IxW6sOO=(z;e17Fff!~CLN9e$>a??q`1otcBJT3e9!=rN6MEqb2aC5Mcj zUDmPgLDyKxqRIKMuk-QZW&}@mp+Dc9I4}SBbS#^^=L8$}s0#mga7(MwvKzTB*DOV8 z_34;_CwtVKw#l`I-JHLMjlQ;w)oMAJHr$qsk9~iDobKC^KWbH+)Qz2NE?l08HgA%M zu#oU6{C#g-(3MZY51*-cFa5FZ$$1*A5#RIcP{mR{Huj zQZ!+DdNKWwNVC&u-lKj+zT#O+^0;{+{_M?e*6>nJ-lK3e@})x{FA(o4t2{d)&;52T z*?;o~visp`R(8^Kx~g?^^5P;TXGXp&V;-8t(wzFWZ$mG% z{>khR@|=`9U!1;5u+*~hLwR%M@3VNf_9y7>fA6rxwfpmG!KI?EMm%J{f85E&m0*1I z`1g_XbNx*Qbo!O`TtAZy49P}IrEWrAC2dCswmU-0>^(!r7G6aIE`MN~Cyt~kDzu~9 zx)$T>EB#^&O_!E8oIA-#oO3uU)3y&e60a$#_~Rzt^L7!QVcKloY;tz`&GYN#v7RBs z$Wed?f0Kt637g1UnBD2e6N&ka&~fxxzw=h5r7=q9@Ba}&U1=#6O4fUL%yx>8&RlC` z&M=NA%<~(&9WqnrkJ~ABAcEicH89h57|;GWG?k^C{n4z|v?pm7lHKUhJ}YAh!byM2 zbED=TjFg<6&T{l_SN=~}6V@SY4axneiTSWuL49@}^py<}?3QVHL@lbD^YwNATj^i+>ak>w4K843FQ969~N zeTfLl_PVmNjXr(YIBFTTh}H|9PM;5tc8>0Ah>Xd8l<_Htfn&^V<_0_=^QUbAbp3#`)$d zYitW+btjYNU45M$UOt#Fh{;aNH!RQQ6xztrMt3sDMYdpjN+f1uCKV*Vw;WG8|1_F^ z9}V8jkaZ~VgRh9~Bcs|am>3~#{t zhmg>W3*7#0NzPNiA8am)CnKEbr^ZRv9WZ?Fhr0RNH!_Pit^+yck zvsd0AYg(4$Yj=z{?k-$zrZ4(5GRKtrEWk{_JKp?4#%B7*?9d`7hm|4GJ%vm9I@=Jz z(8OcerEa%KgP%*10Zpf9ecZOVf(Wkcf6$D`U7tT)@!V|rv4eSN=)dIZzk7^b1zxg} zmZogft5K}ugsQCb?qjS`!v1E(gTeezk?j2Ui8YObbyAYne%a}`s?j7n$->%=&CJ{9 zxK8+oY&=(&NHfi+zuCChlx$+I)NIO%bR_5tjV-_>@&VNp@c; z%U;qs%&2^y^!qRtr?vhN|S65nIl!bTafJChUty{nB5y*H2zr-JCNn51mVn=9toH#Jyv zlLVYi8^wCPOvXQTm?t3@J8wOl(f{iyH$4^PRZp!S+@9ZfJBp<_y~`LnA)3A^pE;`L zwEXn*xWlx>sra$yPFrfj46)>EH)%rq19PsW#KO-`?597g_$ zZ9^Lebm70UqlTqfch=yaaJK(sZu;%4b!M^IKl75CN0Ie^r{j4>HR2n;>%o%DJ!$OQ zIGR`zEhCe<+$JFla`L|aPT{{l*+|objOVou^`uG5tsoVr7NK9-?k58d6y!;#zA-us zIzqOe8%B#w-xxV~{vw_v@g-hxTVvjTWi`@#VPke>QcXT7@f6ncw<{J)i{)mT^rOlA zpFXn2nWuA7Cr#A(WLL=d=^*Byr*uN~Gi=jSTYGde{qy(bo#)IpSDk$-txZQaw?pjE zQ}-gqk@*?+Sw`l~!1lBV;JMC!CTHvCCuP?*BE!b!rXBo`Mh;(Igq8ohJXyH>AZfUv zndMZPP3GzAnRxlZ#q`;EI0H6K;-BzWqZJ($84?*_^eSCI`alGE)l=1nPojC|Ev6TS zuAo;lZI_zF)>^0CHu{Y_+l@t0=}C^^1<3Sn7p0z(?LF$L8^fm2{2x}MWUJL%Senp3__zw) z%$>Ak%;)PR48oJpKU4a%TGeOr+Y1`AbKlNmr^*aruUaMG7f-ZfLu$`q6B}11T^H0g zhAlkKR#zU&HVnyuhNT#*#(lY0lht3EOp!T|Snbbi0<0OsSm2xSi}Y+j_feWR>30toy2ZmLrGm z7z+;#G|qQePud^+#aNuBCp~duT5SIM%~joA4mzcRIt+ZI>!)Ro^C zk-#%6BmZ2!nx#x!m=#`g!uYO1sAWK8QIb9?(#+XB))@BccjguHBJgEg#6p;qKh`z!_- ze1d#8J1sf*d6Fg5_hsq&hOw+5DZp|K%W9^IpM?B-u%P*6@>?@;hnScV&pWW{LzWn2 zQ^zNjzLcgt-W($HX6+#ximW6F0;Z6PuM@K3Z}+k?#}|?*v5Q%jr57yQH*Ao4%FbI4 zXY~Ji%1uuNd(~5OzI|&pH?NR;YxB@i7mpja%3QWosoI|;PqCBS3H0M#ep+Q78=i>1 zNWPkt=($ba7i^>X)>ARxHYbsdFOvB+{$Yhjwxhqs%p|kMUuOPqjpl04ryl0VsOv+!io51z)jwWB1trlOHraWuAtS)^~vpP8tKADfYmC6))N+(o5 z!!|vYZ9_7?mj1))&xkfRC7x;uJ?`k{c8DE%D(}npB&tULm>Qo}8kMIOhLgr3%cFy{ z%nTu~S@p*~+47u^&D4><8K0A6r5B=JnmO|1HeU{TZEkvaz+BtU>i?lXMxUKcljI56 zZM?5>IMUK%gt7HnDXFJ|z3Qp7-_E0VN35fdhi|5{&TWvI#Mav5=nrZxnPM$y7U$I!&Xr_#QE<7N&&Lpiz^J(VPS6f1OZ9%uI-vfDF@^WFDf znD_V2FctWcg)r169FEb5mgbZyOF$cNnt==cg5`K#`E*|r>O z*vZ(rEKlavW`}yEdE!|$*e?l}lbAZW*p!P2>Du~N*!moMBiqbb!1A9dM<1NLW5nd= zmJ-Jc(R1&M85?eWCRNv^rF+XyA{CCGH){P`mzUYHfQ~Obk|r-zlBao=o(I;yWC)$3 za7kZho1Xf8S`RiQYaSYveghlZHBjm)dmpzguIQ;k5679OmL(&_CSEifR2a=#En3Nz z&TVV#%j`!R#AGx+6fYW{d(I-EN^bC>y(Mj9DIX4%J-D*`sXZTnOl** z{>4ex;}art7C6gJkX7V;qY%1|d@=9d{MC4zsx0-dF^IJLu+9=a^r_|bq8G;B$zGFu zd4{vtlZQzECP_(#1t-Y&b@Iflz4nOpZu7f2a!!<)_EB3hw^uk>oIDHpp;8)l+wY1I zd@#SI(2dTl<@;}#|FCHC{!%xiYkW#xMHZBLO5>Bx-T&(;H$4^NRZrC$8OT!n^wM~q zY7@zI^9Z|ob(|1snAo(G1I!?w|U>#6%=DE+=c z6@KkdYSQ58By#xEJsz*=OJn-YA*B45736-xp=@H41m=WYku-JFd+bTP8hrSfM&|1W z(e%LHbjG?96Y2hb?fB$!-AUQacgWc^S=saGmZa16B6QUCnFd>)H1gx_5azf350Y?7 zBmQRLHeMvSE?;x#0_zmHo*e&u28lmCk}odWi_Y6Ho+mq-nEsKcJAHNHAC~0s22!Ta zPIj`-JF=i&IeKZ{M9WV*qUp^c3)r0<+l(c*4S8eH3DtL*O;0`Totk}$_sQHI^un_6 zN@01Qa&&V$#11{xDt`jja&=ldGu~8|H|ubcE4C_&xVo1$pC4mT|ItRZB<)z;HKk~V zG^tI?_T#JxEz55Ais1Wq>?QA-f0o#{{uq6BHa#_<(MqFl*HuQblhceeNBpEF3Gu3@ z`iF0&TLO;K@e|L{a^vA!0Zn3S?Q!%=SBsMK#TOWpw`3&e9+=M2J?g3B!zR#zBPP-^ zFMg)!hgED;fLR*B%kl~V!xC*z&^KL!?O+9%~A}#YKa$B1<>^OJoD&<%pjJ~))a!xVXg_xnN=(ER0@=dWMN5?9p>ig1Ujb#O^xhM~H^?vvPGd$LxPAJpmOi>FaFMQzPGAH9vh$MS9(;#$GR2r1f#z z;)vM_)o_6 z^D``iW>J=N@{i`Qy=hpXm<3g7C zr?+g}@)DN14|B7}6(5<$_!F|?Q)RQ|#RX>dVky~^AFhc zW#_Hrlg{1$>nS%q73x(_W#5pUyzX0)?yZ!Tk1BfCJUVx(IggEr%sMX}ZI$o`^Va$$ z{OOQ^bV{4Fylzl8DTi&N`PNg%Gv6aSX3u9A%$el7eT`|Ykd$;n)$iDharwQh#%>KO9&dIw(bc1Awrd3(O)$lu1f+o@UMDT#U2DQnGp`Nz}W8t*Wx{QM)!xji=< zKfEN}Q@guSEAleQcq}m+l{o`XXAUM;W^QL21Fn;MjhC}53DeN2k@0x{+Ke^*Jqt|{ zza5`+xg&{RWgl&kd!=!E(wvxD8>aIES=;lJAEMZ$+8xZj!$ul?Ba4x%&zjNqB@6I8 zA(N=MPw9l}yUeDi>hSXLc1Xk4jJ;))Oi-E$J?`k{c8DE%>eK|n$}Z2rElbZB4W85| z$1AKS(^~ZQif zb~Zi5f96K{j~$K1@%I}eziS~iNvKynWe!ki*U=l->a9^jj$p0c&}IQrr);pAnb z97fJUlZ-@98ahY!sHY|k8%$|JPd&}>&P?4hnqA7#G$w8O zckF%f&*rqr_gMSJQ`r2LhuFeJgq;iTj1d zsi*1~Go5tZI)nsdt4AAETuU~ly~&>49KdU~oIrm#^*rL2A4iek+>7|BZW~zn zSx1c8H?OlIhyGya>W_`wIp=TEICV?5I4mCh>BlxSaYQ=S{SDzH{v1AL*ejAdYi52Y z!(P_PuN18``6KJL{;s9@)QWWJ-E*X6!ZW1Q*gU0+pKeI@Px+C>lnN%K(7VWzM;p=_ z4b$)+nv9@(Ea^z`x6$lMq9B$h-U7>^$CF6di$=z8Wp)|OmwY7K^L1f2!;@O_(Dm?F zq}0EDKeBb{IyQRydA6`iBc8F$kLH}d5%euEktX#Yli@?Ak*R&_lfRA*AtS1^mG>!y zOZqz7^i-EhNmznku9M{Va3qr@l~|#@kJ}bk^wiBIhe^M-2iU)P$B+Wy$I1Stkz~e< zgfw_b5w@rJa5ihjELP&fa2D_Cn8Gyh7c6uA5-i?Nqe&K;hth`?j7}eK63c@q{^uVB`0j`ctna}& zk+Hjj*|3NsBwn+ukwr5vHyi(zmmU4vB<+6LM3OwK7&))PQ%kzVzmvu3zn*f_Q(<29)W-q8(xX`xldR=l zl|DD}Caax4t+`-UT{dX#slDq)y`n2BEF|-X>|_VeuO|~4PL=m5+i1S^RHBXhN%{QK z=;uKT$;ljDN$Vt=S)wAJ*^`lJ_|2s~Nb_3L`JI5rBx{kpyzs6c*!z-eSl7OsB)yu9 z@5@@5y?h4+>i?3x&QOvq%O6V)FP%%ytr$vl(QN6reb>DH0mB{|2hk4wRNd6DPtwp^=Wcm z_x?y$Z~P&1^uWSoQ~&iWd%wNXk3xTpK0BM9>hRZFxh*+<7yi_|2x)*eS+Keqyj#>j2Yry?t2%M#(GA=`V@Q!9s#qQ{4fpecrq zqrV@q>M2`meRP{2MZ|?*_VIWjGi=*p_HgW9Y}d{qd|ZptY+wE{{AlUMyzt6oti*^R zd_wOe{KZfHd~U*AWZ8;d+;8hr7C&d8)Kg(`>Z!F$2pe>{htYS;I&!7|X14FgRdmks zYW(!fGvw9DMQla$Rjgu_g>2WjV4Al_OS-OhV&3g@dir+i)yPJb+wzE>IauKr$;nJ3 z7fHKl7#mWn7WppYI9~PReWOX6%5=h?6L_nqyGmbqcZj6ykdLmucAA|o-h`(9b0oXJ zEF)`|c!?R9CJC(@)QVTc`~IFquF>xw4KYg1C`xlLD$gI>Mf}*|POMsNAsW``C!_nm zU)j(FS@^X(Ct0SxHR+q}nR#Tdvuw_Z;^z4^$!VtTe!Sn1LF7)aPONK}u55Onl%&PU z5xnZ~U94rg_H^l=!E|q$C*<(T@-*P)Xew|&g-iNM*z{C|{NG0o{id~L-}WEa&gA}D zAGa;8=&7J$F-CC1l&t#MTZE?hnf0r-(d#9;HR8QXMOOK@ zV(piIHXiM$YvycPlNKCaorDK1A*~lx<=Nj|VMft|% z!cW@HTgfM#yZ_fyZh9)*tDZ_Re!N-to1=VZ)pY#h`~$4A|5L;NOFC96;|aQP&?$EE zk9TZG&XM#;+mXEDPFyo^U$BklTTd+NIvBJQuDjL>zMibC7@qEk2C$rw{+n8jpX%~KJ>e#@o9qC zyQKDZ#{=KXu~lNV>1)IGSMiaOdb=?o+F0KVS#V4{UYM zD1LWsEbI7e21`CIh8Yjq@-_u~^8A}Kn{z)bU~~RjYg*p-F&a$&k-Zy{m6h4J%WRz@ zi_}x$aq6i_xo49aKXjqZI=3g~!~eFt8ncgH+~%L>m;H`ma=rp9Ac$L6QETAgC&f;X@uMGMhXEo-t@AL{VT7gn0A)-|$c@-*|WR|#nO z%KJ#Qv4QM!tJJjUpTCo}(>AkPQNQwTfd#qcQ56>4`lL~I<`cH2!zXj?_G|3dtEVDY z=HD&vMhchob++SvySuZEH~X1~G9}=FuS!ZiW$)v*#T7lZWcd!Z9j}Y~r~Q%jnuO#K z!S#6ih6#D~&dXRk^Mm>Kv~V7A>D&LcccoEL9a&f!(@j4JYRjT{5JXYv7NHTvw%@I0 zbHyks3b-*MOe9fc4T1dlj?Lfs`%Zn|t9Rd5=XxRl{*<^tU*{#}B3J%h>Y<1u)7BoHht3smbC)yiiJ z#k5`Tto&JNEWS>$z%0cBuEVFg#i@IA|L;@p_>}B}K2>)i7$&>+1B+f)(I|x>FK*uB z(!RdN*<@I<+t$R$V-werT*)RH)-a3ueqXJb*Xm^Z?>^4zG6Wu z=7I`%g(o5OU#)l}_5mK-7Efz+cXJL?1k}qoff$+ZrV;%eNmla~_CljO1bnHWBXmDO zDenueX89Ia*K9y09cse8wx(2RWJc_-`NFxX=Xh(^*Qm3dN8M!wU{@@HQ<1xHs@-8$ zq2G#pSBB_q<+E*;_B3;f0Tq=B$lN9B#esckG@;ia44-8|eDrUy%g+dKqi+spo$LqO zL$fp)OZNxuuCh*_I&&}!w`Khbf1>xqPoJt;q(fa| z%3;oF2S|1*q^4m*A>_4;j@w^^)%{hta)uIyX2cWErG3cUxLWSSv!i%5+=E+Cmx27g z{m7|lDs0+X1ZV43LDB9C2uhBl7c=+cj|tHv*mfwYU?tR4M_^&AAvP6zQboyiYF0v2D6=S-6rK%&J>Jp{y@43^bo){`HwSPT&rEjD^E7e=_yI|gt zco(1Y(%(#$jB>z0T^W7Z?m{+OJivm|GCb@$o;sYINCvL3#yh7i$+1x_xU*75>rzr~WRm%qgv-CTLh^_kS8_7vp`t6Vk_5 z?lFcfH#XoJtC2Fdcg|}Qi#tmVR1_(qUL8@Cum436zH_U>?L9}qiFbC?66~njmZh$KUla5w+l%>x?uk>yr`yZ(N9A(mFIr)B z1W`@9Wi4-<>7{D;>Ir1`>aTL^7abb(R@HzSo}Q7Wd#8CvSJ$I;r6uAH0h9}cq)BcK z-&jt<2`d2f&Rjv?WFF0)=K-e4EIE5$iM4fzvEDoB^$EU2R3X8&!834rQUkW+{REF* zm0{FlXZC=~j21j}BkNasvqfj~(I$%rn+(;TIzO0Gh4CkJ{Z?PPs*q1Mq!q(=GY(+x zUtVM0vOh!Q{R0paFbrKi=0crKJ8XLyj#m?{kkh@1ak_tlqh=O#*8-Nk=HaTUQ17C= zAqW7o&{Z%<8pTze`Wy_OD)ADdfOtC@m)GYWJF6NXJ%*1D;uzYXKLq;uk0xTn6#VRB zIJf0^BIr&26b}~MMwjK5G|K%djQWRw)W6z+&!oBN6UHYQwj!wNtHJ??pK{+H97$~X zk742IW^9o-Vns6tHcs179p1>d!aL8J5eA2ZMn(Du2S$hnFeZ8ek%%!6F#(Gsmn`*< z2#yV6)b&y3jbZqq9f#mKh$M`yo{q>`eUIQNubTu%vPS`oM~%pXcr4;yWy*p)#mExPgg>i-+)L* past_data = nullptr, + const std::vector* present_data = nullptr) { int min_cuda_architecture = use_float16 ? 530 : 0; bool enable_cuda = HasCudaEnvironment(min_cuda_architecture); @@ -35,6 +40,8 @@ static void RunAttentionTest( std::vector weights_dims = {hidden_size, 3 * hidden_size}; std::vector bias_dims = {3 * hidden_size}; std::vector mask_index_dims = {batch_size}; + std::vector past_dims = {2, batch_size, head_size, past_sequence_length, head_size}; + std::vector present_dims = {2, batch_size, head_size, past_sequence_length + sequence_length, head_size}; std::vector output_dims = input_dims; if (use_float16) { @@ -51,6 +58,23 @@ static void RunAttentionTest( if (mask_index_data.size() > 0) { // mask index is optional. tester.AddInput("mask_index", mask_index_dims, mask_index_data); + } else { + std::vector dims = {static_cast(mask_index_data.size())}; + tester.AddInput("", dims, mask_index_data); + } + + if (use_past_state) { + if (use_float16) { + if (past_sequence_length > 0) { + tester.AddInput("past", past_dims, ToFloat16(*past_data)); + } + tester.AddOutput("present", present_dims, ToFloat16(*present_data)); + } else { + if (past_sequence_length > 0) { + tester.AddInput("past", past_dims, *past_data); + } + tester.AddOutput("present", present_dims, *present_data); + } } if (enable_cuda) { @@ -256,8 +280,7 @@ TEST(AttentionTest, AttentionUnidirectional) { std::vector input_data = { 0.091099896f, -0.018294459f, -0.36594841f, 0.28410032f, - -0.12125026f, -0.0066160089f, 0.38809127f, -0.22455512f - }; + -0.12125026f, -0.0066160089f, 0.38809127f, -0.22455512f}; std::vector weight_data = { -0.2659236192703247f, @@ -310,8 +333,7 @@ TEST(AttentionTest, AttentionUnidirectional) { -0.09368397295475006f, 0.07878211885690689f, 0.2973634898662567f, - 0.11210034042596817f -}; + 0.11210034042596817f}; std::vector bias_data = { -0.0540979839861393f, @@ -325,20 +347,317 @@ TEST(AttentionTest, AttentionUnidirectional) { 0.3670335114002228f, 0.028461361303925514f, -0.08913630992174149f, - 0.28048714995384216f - }; + 0.28048714995384216f}; // No mask_index std::vector mask_index_data = {}; std::vector output_data = { - 0.28109729f, 0.069518551f, 0.0038009658f, 0.29213354f, 0.3692801f, 0.029495837f, -0.084964074f, 0.28169215f - }; + 0.28109729f, 0.069518551f, 0.0038009658f, 0.29213354f, 0.3692801f, 0.029495837f, -0.084964074f, 0.28169215f}; bool is_unidirectional = true; RunAttentionTest(input_data, weight_data, bias_data, mask_index_data, output_data, batch_size, sequence_length, hidden_size, number_of_heads, false, is_unidirectional); } +TEST(AttentionTest, AttentionEmptyPastState) { + int batch_size = 1; + int sequence_length = 2; + int hidden_size = 4; + int number_of_heads = 2; + + std::vector input_data = { + 0.091099896f, -0.018294459f, -0.36594841f, 0.28410032f, + -0.12125026f, -0.0066160089f, 0.38809127f, -0.22455512f}; + + std::vector weight_data = { + -0.2659236192703247f, + 0.02789675071835518f, + 0.07280516624450684f, + 0.050951678305864334f, + 0.020417947322130203f, + -0.04751841351389885f, + 0.043815530836582184f, + 0.006015353370457888f, + -0.11496957391500473f, + -0.1773347705602646f, + 0.30928605794906616f, + 0.005648412741720676f, + + 0.08960387855768204f, + -0.27270448207855225f, + 0.14847396314144135f, + -0.17960812151432037f, + 0.01788954995572567f, + 0.09993876516819f, + 0.03943513706326485f, + -0.02484400011599064f, + -0.12958766520023346f, + 0.220433309674263f, + 0.1720484346151352f, + 0.22024005651474f, + + 0.059368450194597244f, + 0.1710093915462494f, + -0.3967452347278595f, + -0.1591450721025467f, + 0.1446179747581482f, + -0.20505407452583313f, + 0.12749597430229187f, + 0.32139700651168823f, + 0.139958456158638f, + -0.10619817674160004f, + 0.04528557509183884f, + 0.045598603785037994f, + + -0.007152545265853405f, + 0.109454445540905f, + -0.1582530289888382f, + -0.2646341919898987f, + 0.0920850858092308f, + 0.0701494812965393f, + -0.19062495231628418f, + -0.24360455572605133f, + -0.09368397295475006f, + 0.07878211885690689f, + 0.2973634898662567f, + 0.11210034042596817f}; + + std::vector bias_data = { + -0.0540979839861393f, + -0.06444740295410156f, + 0.03112877532839775f, + -0.08288222551345825f, + 0.07840359210968018f, + 0.039143580943346024f, + -0.45591455698013306f, + -0.11876055598258972f, + 0.3670335114002228f, + 0.028461361303925514f, + -0.08913630992174149f, + 0.28048714995384216f}; + + // No mask_index + std::vector mask_index_data = {}; + + std::vector output_data = { + 0.28109729f, 0.069518551f, 0.0038009658f, 0.29213354f, 0.3692801f, 0.029495837f, -0.084964074f, 0.28169215f}; + + std::vector past_data = {}; + + std::vector present_data = { + 0.053175069391727448f, 0.12795503437519073f, 0.11125634610652924f, -0.0510881207883358f, -0.55345797538757324f, -0.3045809268951416f, -0.36920222640037537f, 0.060108467936515808f, 0.28109729290008545f, 0.069518551230430603f, 0.45718482136726379f, -0.010400654748082161f, 0.0038009658455848694f, 0.29213353991508484f, -0.17697516083717346f, 0.27086889743804932f}; + + bool is_unidirectional = true; + bool use_past_state = true; + int past_sequence_length = 0; + int head_size = 2; + RunAttentionTest(input_data, weight_data, bias_data, mask_index_data, output_data, + batch_size, sequence_length, hidden_size, number_of_heads, false, is_unidirectional, + use_past_state, past_sequence_length, head_size, &past_data, &present_data); +} + +TEST(AttentionTest, AttentionPastStateBatch1) { + int batch_size = 1; + int sequence_length = 1; + int hidden_size = 4; + int number_of_heads = 2; + + std::vector input_data = { + -0.019333266f, -0.21813886f, 0.16212955f, -0.015626367f}; + + std::vector weight_data = { + -0.4738484025001526f, + -0.2613658607006073f, + -0.0978037416934967f, + -0.34988933801651f, + 0.2243240624666214f, + -0.0429205559194088f, + 0.418695330619812f, + 0.17441125214099884f, + -0.18825532495975494f, + 0.18357256054878235f, + -0.5806483626365662f, + -0.02251487597823143f, + + 0.08742205798625946f, + 0.14734269678592682f, + 0.2387014478445053f, + 0.2884027063846588f, + 0.6490834355354309f, + 0.16965825855731964f, + -0.06346885114908218f, + 0.4073973298072815f, + -0.03070945478975773f, + 0.4110257923603058f, + 0.07896808534860611f, + 0.16783113777637482f, + + 0.0038893644232302904f, + 0.06946629285812378f, + 0.36680519580841064f, + -0.07261059433221817f, + -0.14960581064224243f, + 0.020944256335496902f, + -0.09378612786531448f, + -0.1336742341518402f, + 0.06061394885182381f, + 0.2205914407968521f, + -0.03519909828901291f, + -0.18405692279338837f, + + 0.22149960696697235f, + -0.1884360909461975f, + -0.014074507169425488f, + 0.4252440333366394f, + 0.24987126886844635f, + -0.31396418809890747f, + 0.14036843180656433f, + 0.2854192554950714f, + 0.09709841012954712f, + 0.09935075044631958f, + -0.012154420837759972f, + 0.2575816512107849f}; + + std::vector bias_data = { + 0.4803391396999359f, + -0.5254325866699219f, + -0.42926454544067383f, + -0.2059524953365326f, + -0.12773379683494568f, + -0.09542735666036606f, + -0.35286077857017517f, + -0.07646317780017853f, + -0.04590314254164696f, + -0.03752850368618965f, + -0.013764488510787487f, + -0.18478283286094666f}; + + // No mask_index + std::vector mask_index_data = {}; + + std::vector output_data = { + 0.20141591f, 0.43005896f, 0.35745093f, 0.19957167f}; + + std::vector past_data = { + 0.55445826f, 0.10127074f, 0.71770734f, 0.15915526f, 0.13913247f, 0.77447522f, 0.66044068f, 0.27559045f, 0.35731629f, 0.62033528f, 0.24354559f, 0.22859341f, + 0.45075402f, 0.85365993f, 0.097346395f, 0.28859729f, 0.26926181f, 0.65922296f, 0.8177433f, 0.4212271f, 0.34352475f, 0.059609573f, 0.46556228f, 0.7226882f}; + + std::vector present_data = { + 0.55445826f, 0.10127074f, 0.71770734f, 0.15915526f, 0.13913247f, 0.77447522f, -0.30182117f, -0.12330482f, 0.66044068f, 0.27559045f, 0.35731629f, 0.62033528f, 0.24354559f, 0.22859341f, -0.36450946f, -0.19483691f, + 0.45075402f, 0.85365993f, 0.097346395f, 0.28859729f, 0.26926181f, 0.65922296f, -0.027254611f, -0.096526355f, 0.8177433f, 0.4212271f, 0.34352475f, 0.059609573f, 0.46556228f, 0.7226882f, -0.025281552f, -0.25482416f}; + + bool is_unidirectional = true; + bool use_past_state = true; + int past_sequence_length = 3; + int head_size = 2; + RunAttentionTest(input_data, weight_data, bias_data, mask_index_data, output_data, + batch_size, sequence_length, hidden_size, number_of_heads, false, is_unidirectional, + use_past_state, past_sequence_length, head_size, &past_data, &present_data); +} + +TEST(AttentionTest, AttentionPastStateBatch2) { + int batch_size = 2; + int sequence_length = 1; + int hidden_size = 4; + int number_of_heads = 2; + + std::vector input_data = { + -0.10902753f, 0.0041178204f, 0.1871525f, -0.20399982f, + 0.027207348f, -0.25321805f, 0.12869114f, 0.023136809f}; + + std::vector weight_data = { + -0.4738484025001526f, + -0.2613658607006073f, + -0.0978037416934967f, + -0.34988933801651f, + 0.2243240624666214f, + -0.0429205559194088f, + 0.418695330619812f, + 0.17441125214099884f, + -0.18825532495975494f, + 0.18357256054878235f, + -0.5806483626365662f, + -0.02251487597823143f, + + 0.08742205798625946f, + 0.14734269678592682f, + 0.2387014478445053f, + 0.2884027063846588f, + 0.6490834355354309f, + 0.16965825855731964f, + -0.06346885114908218f, + 0.4073973298072815f, + -0.03070945478975773f, + 0.4110257923603058f, + 0.07896808534860611f, + 0.16783113777637482f, + + 0.0038893644232302904f, + 0.06946629285812378f, + 0.36680519580841064f, + -0.07261059433221817f, + -0.14960581064224243f, + 0.020944256335496902f, + -0.09378612786531448f, + -0.1336742341518402f, + 0.06061394885182381f, + 0.2205914407968521f, + -0.03519909828901291f, + -0.18405692279338837f, + + 0.22149960696697235f, + -0.1884360909461975f, + -0.014074507169425488f, + 0.4252440333366394f, + 0.24987126886844635f, + -0.31396418809890747f, + 0.14036843180656433f, + 0.2854192554950714f, + 0.09709841012954712f, + 0.09935075044631958f, + -0.012154420837759972f, + 0.2575816512107849f}; + + std::vector bias_data = { + 0.4803391396999359f, + -0.5254325866699219f, + -0.42926454544067383f, + -0.2059524953365326f, + -0.12773379683494568f, + -0.09542735666036606f, + -0.35286077857017517f, + -0.07646317780017853f, + -0.04590314254164696f, + -0.03752850368618965f, + -0.013764488510787487f, + -0.18478283286094666f}; + + // No mask_index + std::vector mask_index_data = {}; + + std::vector output_data = { + 0.14902574f, 0.62273371f, 0.43022552f, 0.12759127f, + 0.26993567f, 0.23553593f, 0.43190649f, 0.086044826f}; + + std::vector past_data = { + 0.42028648f, 0.55855948f, 0.044569403f, 0.76525789f, 0.13962431f, 0.40977913f, 0.36911047f, 0.83399564f, 0.36905321f, 0.91414654f, 0.17300875f, 0.78793788f, + 0.10279467f, 0.80501258f, 0.089550517f, 0.85371113f, 0.61801594f, 0.91222942f, 0.88626182f, 0.069776468f, 0.10591964f, 0.84836882f, 0.83520192f, 0.0098680854f, + 0.3113814f, 0.63999802f, 0.28603253f, 0.98899829f, 0.044405211f, 0.95105386f, 0.81278932f, 0.63969064f, 0.14494057f, 0.11349615f, 0.87086016f, 0.20983537f, + 0.35107401f, 0.90144604f, 0.68950737f, 0.18928574f, 0.18029204f, 0.074517399f, 0.70763874f, 0.48440042f, 0.58114725f, 0.1048766f, 0.73694098f, 0.17766342f}; + + std::vector present_data = { + 0.42028648f, 0.55855948f, 0.044569403f, 0.76525789f, 0.13962431f, 0.40977913f, -0.22849128f, -0.022080801f, 0.36911047f, 0.83399564f, 0.36905321f, 0.91414654f, 0.17300875f, 0.78793788f, -0.4449589f, -0.17704415f, 0.10279467f, 0.80501258f, 0.089550517f, 0.85371113f, 0.61801594f, 0.91222942f, -0.2994619f, -0.14412443f, 0.88626182f, 0.069776468f, 0.10591964f, 0.84836882f, 0.83520192f, 0.0098680854f, -0.33421949f, -0.18547727f, + 0.3113814f, 0.63999802f, 0.28603253f, 0.98899829f, 0.044405211f, 0.95105386f, -0.033968594f, -0.034833729f, 0.81278932f, 0.63969064f, 0.14494057f, 0.11349615f, 0.87086016f, 0.20983537f, 0.045759238f, -0.26863033f, 0.35107401f, 0.90144604f, 0.68950737f, 0.18928574f, 0.18029204f, 0.074517399f, -0.033201858f, -0.10592631f, 0.70763874f, 0.48440042f, 0.58114725f, 0.1048766f, 0.73694098f, 0.17766342f, -0.054369561f, -0.24562015f}; + + bool is_unidirectional = true; + bool use_past_state = true; + int past_sequence_length = 3; + int head_size = 2; + RunAttentionTest(input_data, weight_data, bias_data, mask_index_data, output_data, + batch_size, sequence_length, hidden_size, number_of_heads, false, is_unidirectional, + use_past_state, past_sequence_length, head_size, &past_data, &present_data); +} + } // namespace test } // namespace onnxruntime