קg[ 8ddlZddlZddlmZmZmZmZmZmZm Z ddl Z ddl m cm ZddlmZddlmZddlmZddlmZmZddlmZmZddlmZdd lmZd d gZe j j!j"j#e j j!j"j$e j j!j%j#gZ&e j j!j'j#e j j!j'j$e j j!j(j#gZ)e j*d d d e j*d d d e j*d e j*d e j*d e j*d e j*d fZ+e j*d d d d e j*d d d d e j*d e j*d e j*d e j*d e j*d fZ,de j-jde j-jde.fdZ/de j-jde j-jdeeeffdZ0defdZ1de j-jdee j-jfdZ2dede.fdZ3dZ4dZ5defdZ6defdZ7defdZ8defd Z9d!efd"Z:defd#Z;d$ed%ed&eed'ed(eddf d)Zee>e?fffd,Z@ d=d.ed/eed0fd1e.defd2ZAd3eddfd4ZBd5ZC d>d6e j-jd7e.d8eeefd9ZD d?d6e j-jd8eeed:eee eEeFe.e jGfeFffd;ZHd+efd<ZIdS)@N)AnyCallableDictListOptionalTupleUnion)capture_pre_autograd_graph)quantized_decomposed_lib)QuantizationAnnotation) _assign_attr _AttrKind) GraphModuleNode)fuse_conv_bn_weights)LeafSpecfold_bn_weights_into_conv_node"remove_tensor_overload_for_qdq_opssourcedestreturncTttz}|tjjjj|j|vr]t|j dtj j std|j d|j d}|j|v]||kS)z Assuming dest is one of the ops inserted by quant workflow, this function finds if source and dest are connected. Assumption is that only quant workflow inserted ops exist between source and dest rz=expected arg[0] of quant workflow ops to be a node but found ) _QUANTIZE_OPS_DEQUANTIZE_OPSappendtorchopsquantized_decomposedchoose_qparamstensortarget isinstanceargsfxr ValueError)rrquant_workflow_opss \/var/www/html/ai-engine/env/lib/python3.11/site-packages/torch/ao/quantization/pt2e/utils.py _is_connectedr*<s '8ei<KRSSS ++ + +$)A, 66 ^PTPYZ[P\^^ y| ++ + + 6>produerusercd}|jD]N}t|tjjr-|jdkr"|jtvrt||r|}nO|V|j D]N}t|tjjr-|jdkr"|jtvrt||r|}nO|dSd}|jdjdkr&|jdjtvr |jd}||fS)z Find q, dq pair corresponding to [producer -> q -> dq -> user] Utils works by finding dq arg of user and ensuring it is connected to producer N call_functionNNr) r%r$rr&ropr#rr*kwargsr)r,r-dq_nodenq_nodes r)_find_q_dq_node_for_userr6MsG Y q%(- ( ( ''O++Wa((   A1ehm,, DO++H// !,,GE| F Qo-- LO "m 3 3a G r+nodec$|jdkr$|jtjjjjkpa|jtjjjjkp=|jtjjjkp|jtjjjkS)Nr/)r1r#rratensym_sizedefault sym_numel)r7s r)_is_sym_size_noder=uss ?" ; K59>2: : 2 ;%).2: : 2 ;%).2 2 2 ;%).1 1 r+cLttd|j}|S)Nc$t|duS)NF)r=)xs r)z(_filter_sym_size_users..s):1)=)=)Fr+)listfilterusers)r7 node_userss r)_filter_sym_size_usersrFs%fGG$*UUVVJ r+ annotationcX|dS|j}|j}t|dkr|dSdS)NFrT)input_qspec_map output_qspeclen)rGrIrJs r)_is_valid_annotationrLs?u 0O*L ?q  \%9u 4r+c $|dS|jdksJ|jd}|}t|D]R\}}t ||s-t dd|d|t||}S|S)Nget_attr.z#Node referenced nonexistent target )r1r#split enumeratehasattr RuntimeErrorjoingetattr)r7m target_atomsattr_itriatoms r)_get_tensor_constant_from_noder[s |t 7j ;$$S))LH\**++4x&& Rchh|BQB?O6P6PRR 8T** Or+c(g}t|D]\}}|j|vr!|||j/|js/|t |kr|||e||j|SN)rQnamer kwarg_onlyrK default_value) orig_args orig_kwargs args_schemaall_argsrYschemas r)_get_all_argumentsrfsH{++22 6 ;+ % % OOK 4 5 5 5 5" 2q3y>>'9'9 OOIaL ) ) ) ) OOF0 1 1 1 1 Or+ctjjjjtjjjjtjjjjtjjjjg}|j|vS)zU Return True if the given node refers to an aten batch norm op QAT supports. ) rrr9 batch_normr;_native_batch_norm_legitcudnn_batch_normmiopen_batch_normr#)r7 supported_opss r)%_is_supported_batch_norm_for_trainingrmsP  !) /7  '/ (0M ;- ''r+r4c|jdko<|jtjjjjtjjjjfvS)z< Return whether the node refers to an aten conv op. r/)r1r#rrr9conv1dr;conv2dr4s r) _is_conv_noderrs@ 4? " qx % %4(r+c|jdkof|jtjjjtjjjjtjjjtjjjjfvS)zF Return whether the node refers to an aten conv_transpose op. r/) r1r#rrr9conv_transpose1dr;conv_transpose2dinputrqs r)_is_conv_transpose_noderwsT 4? " qx ' '/ ' '- 4(r+c>t|pt|S)zN Return whether the node refers to an aten conv or conv transpose op. )rrrwrqs r)_is_conv_or_conv_transpose_noderys     96q999r+conv_fnc6|tjtjfvSr])Frtru)rzs r)_is_conv_transpose_fnr}s q)1+=> >>r+cht|p#|jtjjjjkSr])rmr#rrr9$_native_batch_norm_legit_no_trainingr;rqs r) _is_bn_noders,-a00 S 8uy~JR Rr+ conv_nodeconv_weight_nodeconv_bias_nodebn_noderVc Zt||}t||}t|}|jjj}t |j|j|} t| d|} t| d|} t| d|} t| d|} |jtj j j j krd}n't|rd}ntd|j| |}t||| | || | |\}}t!|j}t#|dkr|d|j}t'|t(sJt+|||t,j|1|j}t+||t)|t,jnr|d z}t+|||t,j|j|5|j|}dddn #1swxYwY||d<t7||_|jtj j jj kr||nQ|jD]I}|jd ks&|jt@j!ks|jdd kr4||J|j"t#|jd kr|j#|dSdS) NrrzBN node target is unexpected ) transpose_biasr/r)$r[rwr#_schema argumentsrfr%r2rrr9rr;rmr'rrBrKrr$strr r PARAMETERgraphinserting_beforerNtuplerhreplace_all_uses_withrDr1operatorgetitemeliminate_dead_code erase_node)rrrrrVconv_wconv_brbn_args_schemabn_argsbn_wbn_bbn_rmbn_rv eps_arg_indexbn_eps fused_weight fused_bias conv_argsweight_attr_namebias_attr_name get_bias_noder-s r)rrsT,, >F' 22I^+5N w~~NNG )'!*a 8 8D )'!*a 8 8D *71:q 9 9E *71:q 9 9E~LTTT .w 7 7J 8'.III ] #F3ufdDI   L* Y^$$I 9~~(. & , ,,,,q"2I4GHHH!'.ZC$7$79LMMMM)G3ZNI4GHHH W % %i 0 0 = =G,,^<z!_fuse_conv_bn_..Es(77AQ777777r+r/rrr)anyrnodesr1r#rrr9rr;rhr%ryrKrr recompile)rVhas_bnr4rrrrs r)_fuse_conv_bn_rDs% 77777 7 7F  W]   4? " "ah IN ? G IN % -7 ' '   LO.q11   $>!,.1)..A.AA.E.E**4& '!    G!!!KKMMMMMr+modelc8i}|jjD]}|jdd}dt df}|rPt |d}|ddd|df}|||j<|S)Nnn_module_stackrrOr) rrmetagettyperBvaluesrPr^)rnode_name_to_scoper4r current_scopebts r)_get_node_name_to_scoper]s68 [ 33&**%6==T$ZZ(  :o,,..//3BU[[--b12a59M%216"" r+Fpatternexample_inputs.is_cudac |rtd|D}t|||}|j||jjD]c}|jdkrV|jtj j j j kr2t|jdkr|j|d|j||S)zF Convert the pattern to an FX graph with decomposed aten ops. cng|]2}t|tjr|n|3S)r$rTensorcuda)rr@s r) z6_get_aten_graph_module_for_pattern..us5 T T TAu|44 ;QVVXXX! T T Tr+r/r)rr rrrrr1r#rrr9copy_r;rKrDr)rrrr2 aten_patternr7s r)"_get_aten_graph_module_for_patternrjs  T T^ T T T  .L **,,,"(00 G & & uy~3;;;DJ1$$   ) )$ / / /**,,, r+ match_patternctjjjjtjjjtjjjjtjjjtjjjjtjjjtjjjjtjjjtjjjjtjjjtjjjjtjjjtjjjjtjjjtjjj jtjjj tjj j j tjj j i }|j jD])}|jdkr|j|vr||j|_*dS)zRemove .tensor overload for quantize/dequantize ops so that we can use the match_pattern that we get from torchdynamo export to match the output of convert_pt2e r/N)rrr quantize_per_tensorr;dequantize_per_tensorr"tensor2quantize_per_channeldequantize_per_channelr9clamprrrr1r#)r_MAPr4s r)rrs-  &:BEIDbDv &<DeiFdFz &:A59CaCu &<CUYEcEy &:BEIDbDv &<DeiFdFz &;CUYEcEx &=EuyGeG| #UY^%9 D &&& 4? " "  8t  AH~AH &&r+ct|ttfrdSt|ttfr"t t t|SdS)NTF)r$intfloatrrBallmap _is_literal)args r)rrsP#U|$$t#t}%%*3{C(())) 5r+gm merge_dupexclude_literalsc d}d}i}|g}|j}|jd}|jjD]&}|jdkr|}|dz }|j|5g} |jD]} t| r| |vr|r | |vr| || 7|j dt|z} | | |jt|dz }|r| || <| | t| } dddn #1swxYwY| |_(| | |S)aReplace the literals in the graph with placeholder nodes that's created on the fly while we traverse the graph, so that the literal arguments in the graph can be matched and replaced To use this, the pattern and replacement graph should have the exact same number of literal args and they should be used in the exact same order in the pattern and replacement graph. If the literal arguments are not used in the same order in pattern and replacement graph, please use `_replace_literals_with_existing_placeholders` instead Args: `gm`: input GraphModule that we'll transform `merge_dup`: boolean flag to indicate that if the same literal appears multiple times in the graph, whether they should correspond to the same placeholder or not `exclude_literals`: a list of literals that will not be replaced with placeholders Example: # 1. Original Graph def pattern(self, x): return x + 3 def replacement(self, x): return x - 3 example_inputs = (torch.randn(1, 3, 3, 3),) pattern_gm = _get_aten_graph_module_for_pattern(pattern, example_inputs) replacement_gm = _get_aten_graph_module_for_pattern(pattern, example_inptus) # 2. Before calling replace literals we'll see the following graph: def pattern(self, x): return x + 3 def replacement(self, x): return x - 3 pattern_gm = _replace_literals_with_new_placeholders(pattern_gm) replacement_gm = _replace_literals_with_new_placeholders(replacement_gm) # 3. After replacing literals with new placeholder nodes def pattern(self, x, new_ph): return x + new_ph def pattern(self, x, new_ph): return x - new_ph Nr placeholderrr)_in_specchildren_specsrrr1inserting_afterr%rrrrrr __post_init__) rrrlast_phcnt literal_to_phin_spec args_specr7new_argsrph_nodes r)'_replace_literals_with_new_placeholdersrshG CFHMkG&q)I 7m # #G 1HC  X % %g . . ' 'Hy ) )s## )3C(C(C 9SM%9%9  c(:;;;;"$("6"6us3xx7G"H"H 000!077 CCCq$918M#.OOC((((XH ' ' ' ' ' ' ' ' ' ' ' ' ' ' '"   IsCD??E E literal_to_ph_idxc`|g}|i}d|jjD}|jjD]}|jdkrg}|jD]T}t |r.||vr*||vr&||}||}||?||Ut |}||_|S)a Replace the literals in the graph with **existing** placeholder nodes, so that the literal arguments in the graph can be matched and replaced To use this, all literal args in the graph should be unique and each of them should correspond to exactly one placeholder node # 1. Original Graph def pattern(self, x_i8, scale, zero_point, quant_min, quant_max): return torch.dequantize_per_tensor(x_i8, scale, zero_point, quant_min, quant_max) def replacement(x_i8, scale, zero_point, quant_min, quant_max): x_i8 = torch.clamp(x_i8, quant_min, quant_max) return ((x_i8.to(torch.float32) - zero_point) * scale).to(dtype=torch.float32) example_inputs = ( torch.randn(1, 3, 3, 3), 1.0, 0, -128, 127, ) pattern_gm = _get_aten_graph_module_for_pattern(pattern, example_inputs) replacement_gm = _get_aten_graph_module_for_pattern(pattern, example_inptus) # 2. Before calling replace literals we'll see the following graph: def pattern(self, x_i8, scale, zero_point, quant_min, quant_max): # scale/zero_point/quant_min/quant_max are burnt in since they are scalar values return torch.dequantize_per_tensor(x_i8, 1.0, 0, -128, 127) def replacement(x_i8, scale, zero_point, quant_min, quant_max): # scale/zero_point/quant_min/quant_max are burnt in since they are scalar values x_i8 = torch.clamp(x_i8, -128, 127) return ((x_i8.to(torch.float32) - 0) * 1.0).to(dtype=torch.float32) # Note that literal args appear in different order in pattern and replacement graph, so # we can't use _replace_literals_with_new_placeholders literal_to_ph_idx = {1.0: 1, 0: 2, -128: 3, 127: 4} pattern_gm = _replace_literals_with_existing_placeholders(pattern_gm, literal_to_ph_idx) replacement_gm = _replace_literals_with_existing_placeholders(replacement_gm, literal_to_ph_idx) # 3. After replacing literals with existing placeholder nodes def pattern(self, x_i8, scale, zero_point, quant_min, quant_max): # scale/zero_point/quant_min/quant_max are burnt in since they are scalar values return torch.dequantize_per_tensor(x_i8, scale, zero_point, quant_min, quant_max) def replacement(x_i8, scale, zero_point, quant_min, quant_max): # scale/zero_point/quant_min/quant_max are burnt in since they are scalar values x_i8 = torch.clamp(x_i8, quant_min, quant_max) return ((x_i8.to(torch.float32) - zero_point) * scale).to(dtype=torch.float32) Nc(g|]}|jdk |S)r)r1)rr7s r)rz@_replace_literals_with_existing_placeholders..Fs$ G G GDdg.F.F4.F.F.Fr+r/)rrr1r%rrr) rrrphsr7rrph_idxrs r),_replace_literals_with_existing_placeholdersrsr  G GBHN G G GC 7o % % 9 % %CC   %///,,,*3/f+(((($$$$?? Ir+cdddtffd }ddtffd }tj|||_tj|||_|S)z Disallow calling `model.train()` or `model.eval()` on the given GraphModule. This is useful for exported models, where these methods don't actually behave as expected. ah Calling train() or eval() is not supported for exported models. Please call `torch.ao.quantization.move_exported_model_to_train(model)` (or eval) instead. If you cannot replace the calls to `model.train()` and `model.eval()`, you may override the behavior for these methods by calling `torch.ao.quantization.allow_exported_model_train_eval(model)`, which does the above automatically for you. Note that this has limited effect on switching behavior between train and eval modes, and should be used only for special ops such as dropout and batchnorm. Tmodec"tr]NotImplementedErrorselfr error_messages r)_trainz$_disallow_eval_train.._trainn!-000r+c"tr]rrs r)_evalz#_disallow_eval_train.._evalqrr+)T)booltypes MethodTypetraineval)rrrrs @r)_disallow_eval_trainr^s  M11411111111$111111"6511EK!%//EJ Lr+)F)FNr0)Jrrtypingrrrrrrr rtorch.nn.functionalnn functionalr| torch._exportr $torch.ao.quantization.fx._decomposedr torch.ao.quantization.quantizerr torch.export.unflattenr rtorch.fxrrtorch.nn.utils.fusionrtorch.utils._pytreer__all__rr rr;r"rrrrrrandn_conv1d_bn_example_inputs_conv2d_bn_example_inputsr&rr*r6r=rFrLr[rfrmrrrwryr}rrrrrrrrrrrrdtyperrrr+r)rs  DDDDDDDDDDDDDDDDDD 444444JIIIIIBBBBBB::::::::&&&&&&&&666666((((((%(   I"6> I"6= I"7?  I"8@ I"8? I"9AEK1aEK1aEKNNEKNNEKNNEKNNEKNNEK1aEK1aEKNNEKNNEKNNEKNNEKNN%(-ux}"% X]%"'(-% 38_%%%%PD4 3F %;       ( ( ( ( ("T t    :t::::?8????4^$^$^$TN^$ ^$  ^$  ^$^$^$^$Dkd2 ; 4U39=M8M3N     "" "#s(O""  """"J&k&d&&&&,,0WW WWtCy)WWWWx-1SWRR RtCy)R U5#tU[+H%I3%N OPRRRRn r+