
    h
K                        S SK JrJrJr  S SKJrJrJrJr  S SK	J
r
  S SKJr  S SKJr  S SKrS SKJrJrJrJr  S SKJr  S SKrS SKrS SKrS SKrS SKrS SKrS S	KJr  S
\4S jrS\
R@                  4S jr!\RD                  " 5       S 5       r#\RD                  " 5       S\$4S j5       r%S\$4S jr&\RD                  " 5       S\$4S j5       r'\RD                  " S5      S 5       r(S\$4S jr)\" SS9 " S S5      5       r* " S S\5      r+g)    )BaseBackend	GPUTargetLanguage)irpassesllvmnvidia)knobs)
PTXASError)	dataclassN)AnyDictTupleOptional)
ModuleType)Pathtargetc                 @    S[         [        [        [        4   4S jnU$ )Nreturnc                     U R                   R                  nUR                   R                  nX#:X  d   S5       eUS:X  a  gg)Nz%lhs and rhs bitwidth must be the same   )   r       )r   r   r   )scalarprimitive_bitwidth)lhs_typerhs_typelhs_bitwidthrhs_bitwidths       Y/var/www/fran/franai/venv/lib/python3.13/site-packages/triton/backends/nvidia/compiler.pycheck_dot_compatibility-min_dot_size.<locals>.check_dot_compatibility   s@    9999+T-TT+1    )r   int)r   r!   s     r    min_dot_sizer%      s!     uS#s]7K   #"r#   r   c                  6    [         R                  R                  $ N)r
   r	   ptxas r#   r    	get_ptxasr*   !   s    <<r#   c                      [         R                  R                  n U b  U $ [        R                  " [        5       R                  S/5      R                  S5      nU$ )Nz	--versionutf-8)r
   r	   mock_ptx_version
subprocesscheck_outputr*   pathdecode)mock_verversions     r    get_ptxas_versionr4   %   sI    ||,,H%%y{'7'7&EFMMgVGNr#   c                     [        U [        5      (       d   e[        [        U R	                  S5      5      u  pUS:X  a  US:  a  SU-   $ SU-   S-
  $ US:X  a  SU-   $ US:X  a  S	U-   $ [        S
U -   5      e)zC
Get the highest PTX version supported by the current CUDA driver.
.      P         F   
   ?   z?Triton only support CUDA 10.0 or higher, but got CUDA version: )
isinstancestrmapr$   splitRuntimeError)cuda_versionmajorminors      r    ptx_get_versionrG   .   s    
 lC((((sL..s34LE{19::>!{Ez{Ez
X[gg
hhr#   archc                 b    U R                   nUc  [        5       R                  n[        U5      nU$ r'   )ptx_versionr*   r3   rG   )optionsrH   rJ   rD   s       r    get_ptx_version_from_optionsrL   A   s0    %%K {**%l3r#   c                 >    [        X5      n[        SU5      nSU 3nU$ )NV   z+ptx)rL   min)rK   rH   rJ   llvm_ptx_versionfeaturess        r    get_featuresrR   I   s.    .w=K 2{+&'(HOr#   c                     [        U S5       n[        R                  " UR                  5       5      R	                  5       sS S S 5        $ ! , (       d  f       g = f)Nrb)openhashlibsha256read	hexdigest)r0   fs     r    	file_hashr[   W   s5    	dD	Q~~affh'113 
		s   2A		
A
capabilityc                 $    U S:  a  SOSnSU  U 3$ )NZ   a sm_r)   )r\   suffixs     r    sm_arch_from_capabilityrc   ]   s!    "$S"FVH%%r#   T)frozenc                   f   \ rS rSr% Sr\\S'   Sr\\S'   Sr\\S'   Sr	\
\   \S	'   S
r\\S'   Sr\\S'   Sr\\S'   Sr\
\   \S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\   \S'   Sr\\   \S'   Sr\\S'   Sr\\   \S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\S '   Sr\\S!'   Sr \\S"'   S# r!S$ r"Sr#g)%CUDAOptionsc      	num_warpsr:   num_ctas   
num_stagesNmaxnreg)r:   r:   r:   cluster_dimsrJ   ptx_optionsir_overrideTenable_fp_fusionFlaunch_cooperative_grid
launch_pdl)fp8e5fp8e4b15supported_fp8_dtypesr)   !deprecated_fp8_dot_operand_dtypestf32default_dot_input_precision)rx   tf32x3ieeeallowed_dot_input_precisionsmax_num_imprecise_acc_defaultextern_libsdebugcudabackend_namesanitize_overflowrH   c                    [        [        5      R                  S-  nU R                  c  0 O[	        U R                  5      nUR                  SS 5      (       d2  [        R                  R                  =(       d    [        US-  5      US'   [        R                  U S[        UR                  5       5      5        U R                  S:  a   U R                  U R                  S-
  -  S:X  d   S5       eg )Nlib	libdevicezlibdevice.10.bcr~   r   r:   znum_warps must be a power of 2)r   __file__parentr~   dictgetr
   r	   libdevice_pathr@   object__setattr__tupleitemsri   )selfdefault_libdirr~   s      r    __post_init__CUDAOptions.__post_init__}   s    h..6 ,,4b$t?O?O:P{D11',||'B'B'mc.[lJlFmK$4k6G6G6I0JK~~!t~~!9K'LQR&R 	0/	0R&Rr#   c           	      d   [        U R                  5      n[        S [        US   5       5       5      US'   SR	                  [        UR                  5       5       VVs/ s H  u  p#U SU 3PM     snn5      n[        R                  " UR                  S5      5      R                  5       $ s  snnf )Nc              3   @   #    U  H  u  pU[        U5      4v   M     g 7fr'   )r[   ).0kvs      r    	<genexpr>#CUDAOptions.hash.<locals>.<genexpr>   s     (hGgtq!Yq\):Ggs   r~   _-r,   )
r   __dict__r   sortedjoinr   rV   rW   encoderY   )r   	hash_dictnamevalkeys        r    hashCUDAOptions.hash   s    '	#((hviXeNfGg(h#h	- hh	@Q9RS9RID4&#9RST~~cjj12<<>> Ts   B,
)$__name__
__module____qualname____firstlineno__ri   r$   __annotations__rj   rl   rm   r   rn   r   rJ   ro   r@   rp   rq   boolrr   rs   rv   r   rw   ry   r|   r}   r~   r   r   r   r   rH   r   r   __static_attributes__r)   r#   r    rf   rf   c   s    IsHcJ "GXc]!#L%#KK!%K#%!d!$)T)J'<%*<46%uSz6'--/I %*I*.!4.KE4L#"t"D#0?r#   rf   c                     ^  \ rS rSr\S\4S j5       rS rS\4S jr	S\SS4U 4S jjr
S\4S	 jrS
 rS rS\\\4   4S jrS r\S 5       r\S 5       rS rS rS rS rS r\R6                  " 5       S 5       rSrU =r$ )CUDABackend   r   c                      U R                   S:H  $ )Nr   )backend)r   s    r    supports_targetCUDABackend.supports_target   s    ~~''r#   c                     Sn[         R                  " X!5      nU(       d  [        SU 35      e[        UR	                  S5      5      $ )Nz	^sm(\d+)$z(TRITON_OVERRIDE_ARCH must have the form r:   )re	fullmatch
ValueErrorr$   group)r   rH   patternmatchs       r    _parse_archCUDABackend._parse_arch   s>    W+GyQRR5;;q>""r#   r   c                 B    U R                  UR                  5      nSU 3$ )Ncuda:)r   rH   )r   rK   r\   s      r    get_target_nameCUDABackend.get_target_name   s#    %%gll3
zl##r#   Nc                 2   > [         TU ]  U5        SU l        g )Ncubin)super__init__
binary_ext)r   r   	__class__s     r    r   CUDABackend.__init__   s     !r#   c                    S[         R                  R                  =(       d    SU R                  R                   30nUR                  [        R                  R                  5        Vs0 s H  o3U;   d  M
  X   c  M  X1U   _M     sn5        [        U R                  US   5      5      nSU;  aG  [        [        R                  5      nUS:  a  UR                  S5        [        [        U5      5      US'   SU;  a  US:  a  SUS'   S	U;  a  [         R                   R"                  US	'   US:X  a  S
OSUS'   [        S0 UD6$ s  snf )NrH   smrv   Y   fp8e4nvrw   r^   )ru   rq   i   @r   r}   r)   )r
   runtimeoverride_archr   rH   updaterf   __dataclass_fields__keysr$   r   setrv   addr   r   languagedefault_fp_fusion)r   optsargsr   r\   rv   s         r    parse_optionsCUDABackend.parse_options   s-   33NDKK<L<L;M7NO)I)I)N)N)Pu)PAY]T]ZaeahZQQZ)Puv))$v,78
!-#&{'G'G#H R$((3+08L1M+ND'(.d:R<J89T)',~~'G'GD#$9Cr9Iq,-"T""% vs   (	E5E<	Ec                     UR                   UR                  UR                  UR                  S   UR                  S   UR                  S   4$ )Nr   r:      )ri   rj   sharedrn   )r   metadatas     r    pack_metadataCUDABackend.pack_metadata   sO    OO!!!$!!!$!!!$
 	
r#   c                     SS K Js  Js  Jn  [	        U R                  UR                  5      5      nUS:  a  UR                  OUR                  [        U R                  5      S.nU$ )Nr   r9   )convert_custom_typesr%   )triton.language.extra.cudar   extrar   r$   r   rH   convert_custom_float8_sm80convert_custom_float8_sm70r%   r   )r   rK   r   r\   codegen_fnss        r    get_codegen_implementation&CUDABackend.get_codegen_implementation   sV    11))',,78
 0:R/?D++TEdEd%

 r#   c                     SSK Jn  SU0$ )Nr   )r   ztriton.language.extra.libdevice)r   r   )r   r   s     r    get_module_mapCUDABackend.get_module_map   s    819==r#   c                 0    [         R                  " U5        g r'   )r	   load_dialects)r   ctxs     r    r   CUDABackend.load_dialects   s    S!r#   c                    [         R                  " U R                  5      nUR                  5         [        R
                  R                  U5        [        R                  R                  U5        US-  S:  a  [        R                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R
                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        UR!                  U 5        U $ )Nr=   	   )r   pass_managercontextenable_debugr   commonadd_inlinerttiradd_rewrite_tensor_pointer(add_rewrite_tensor_descriptor_to_pointeradd_canonicalizeradd_combineadd_reorder_broadcastadd_cseadd_symbol_dceadd_loop_unrollrun)modr   optr\   pms        r    	make_ttirCUDABackend.make_ttir   s    __S[[)
!!"%..r2aKK@@D''+#))"-b!$$R(##B'
s
r#   c                    UR                   bI  U R                  S[        R                  " U R                  5      R                  UR                   5      5        [        R                  " 5       nUR                  b<  UR                  S   Ul	        UR                  S   Ul
        UR                  S   Ul        [        R                  " U R                  5      nUR                  5       n[        R                  R!                  USU 3UR"                  SUR$                  5        [        R&                  R)                  U5        US-  S:  a  [        R&                  R+                  U5        [        R                  R,                  R/                  XT5        [        R&                  R1                  U5        [        R&                  R3                  U5        [        R&                  R5                  U5        [        R&                  R1                  U5        [        R&                  R7                  XSS	:  5        [        R                  R,                  R9                  U5        [        R                  R;                  U5        US-  S
;   GaC  [        R&                  R=                  U5        [        R>                  RA                  U5        [        R                  RC                  U5        [        R>                  RA                  U5        [        R&                  RE                  U5        [        R                  RF                  RI                  XRRJ                  U5        [        R&                  RM                  XRRJ                  5        [        R&                  RO                  U5        [        R&                  RQ                  XRRJ                  U5        GOUS-  S:  Ga  [        R&                  R=                  U5        [        R>                  RA                  U5        [        R                  RC                  U5        [        R&                  RS                  U5        [        R&                  RU                  U5        [        R                  R,                  RW                  U5        [        R&                  RM                  XRRJ                  5        [        R&                  RO                  U5        [        R&                  RY                  XRRJ                  5        [        R&                  RQ                  XRRJ                  U5        [        R&                  RE                  U5        [        R                  R,                  R[                  U5        O[        R                  RC                  U5        [        R>                  RA                  U5        [        R                  R;                  U5        [        R&                  R]                  U5        [        R&                  R7                  XSS	:  5        [        R&                  R_                  U5        [        R                  R,                  Ra                  U5        [        R&                  R1                  U5        [        R                  R,                  Rc                  U5        [        R&                  Re                  U5        [        R&                  Rg                  U5        [        R                  R;                  U5        [        R>                  Ri                  U5        US-  S:  aR  [        R                  R,                  Rk                  U5        [        R                  R,                  Rm                  U5        [        R>                  Ro                  U5        [        R>                  RA                  U5        URq                  U 5        UR                  UR                  UR                  4US'   U Rs                  5       nXqS'   U $ )Nzttg.maxnregr   r:   r   r   r   r=   r   r9   )r   r   r   rn   tensordesc_meta):rm   set_attrr   builderr   get_int32_attrr	   ClusterInforn   clusterDimXclusterDimYclusterDimZr   r   r   r   add_convert_to_ttgpuirri   rj   ttgpuiradd_coalesceadd_f32_dot_tc	ttnvgpuiradd_plan_ctaadd_remove_layout_conversionsadd_optimize_thread_localityadd_accelerate_matmuladd_optimize_dot_operands add_optimize_descriptor_encodingadd_loop_aware_cseadd_fuse_nested_loopsr   r   add_triton_licm add_combine_tensor_select_and_ifhopperadd_hopper_warpspecrl   add_assign_latenciesadd_schedule_loopsadd_pipelineadd_optimize_accumulator_initadd_hoist_tmem_allocadd_promote_lhs_to_tmemadd_warp_specializeadd_remove_tmem_tokensadd_prefetchadd_coalesce_async_copyadd_optimize_tmem_layoutsadd_interleave_tmemadd_reduce_data_duplicationadd_reorder_instructionsr   add_tma_loweringadd_fence_insertionadd_sccpr   get_tensordesc_metadata)r   r   r  r\   cluster_infor  dump_enabledr  s           r    
make_ttgirCUDABackend.make_ttgir   s    ;;"LL

3;;(?(N(Ns{{([\))+''*'7'7':L$'*'7'7':L$'*'7'7':L$__S[[)(**2zl/CS]]TVX[XdXde##B'q NN))"-,,R>44R833B7,,R044R80025EF@@D&&r*v%NN004MM++B/KK''+MM++B/NN;;B?MM  44RVNN//NNCNN--b1NN''NNLI2#NN004MM++B/KK''+NN88<NN//3MM##;;B?NN//NNCNN--b1NN..r>>BNN''NNLINN;;B?MM##::2>KK''+''+&&r*##B'0025EF..r299"=44R833B72226//3&&r*$$R(q MM##44R8MM##77;r"''+
s$0$<$<l>V>VXdXpXp#q 557&5"#
r#   c                    Un[         R                  " UR                  5      nUR                  5         [        R
                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R
                  R                  U5        [        R
                  R                  U5        UR                  U5        UR                  5       US'   U$ )Nr  )r   r   r   r   r   r  r   r   r/  r   r  r   r  r   r0  )r   srcr   rK   r\   r   r  s          r    	ttgir_optCUDABackend.ttgir_opt0  s    __S[[)
""2&r"&&r*((,77;
s&)&A&A&C"#
r#   c                    [        X0R                  R                  5      nUn[        R                  " UR
                  5      nUR                  5         [        R                  R                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R                  R!                  U5        [        R                  R                  R#                  U5        [        R                  R%                  U5        [        R                  R                  R'                  XtU5        [        R(                  R+                  U5        [        R(                  R-                  U5        [        R                  R                  R/                  U5        [        R                  R                  R1                  U5        [        R(                  R+                  U5        [        R(                  R-                  U5        [        R(                  R3                  U5        [4        R6                  R8                  (       d  [        R:                  R=                  U5        UR?                  U5        [@        RB                  " 5         [@        R
                  " 5       n[4        R6                  RD                  (       a  [G        S5      e[@        RH                  " Xh5      n	[K        U5      n
[M        X0R                  R                  5      nSn[        RN                  " 5         [@        RP                  " XX5        [        RR                  " U	5        URT                  (       a6  URT                   VVs/ s H  u  pUPM	     nnn[@        RV                  " X5        [@        RX                  " U	[@        RZ                  5        UR]                  S5      nUb  UUS'   UR]                  S5      US'   UR]                  S5      US'   UR]                  S	5      US
'   UR]                  S5      US'   [_        U	5      nA	AU$ s  snnf )NzYAddress Sanitizer Error: Address sanitizer is currently only supported on the AMD backendnvptx64-nvidia-cudazttg.total-num-warpsri   z
ttg.sharedr   zttg.tensor_memory_size	tmem_sizezttg.global_scratch_memory_sizeglobal_scratch_sizez#ttg.global_scratch_memory_alignmentglobal_scratch_align)0rL   r   rH   r   r   r   r   r	   r   r  add_lower_mmar  r  add_allocate_warp_groupsconvertadd_scf_to_cfadd_allocate_shared_memoryadd_allocate_tensor_memory"add_allocate_global_scratch_memoryadd_to_llvmirr   r   r   add_nvgpu_to_llvmadd_warp_specialize_to_llvmr   r
   compilationdisable_line_infollvmiradd_di_scoper   r   init_targetsenable_asanrC   	to_modulerc   rR   set_short_ptrattach_datalayoutset_nvvm_reflect_ftzr~   link_extern_libsoptimize_moduleOPTIMIZE_O3get_int_attrr@   )r   r6  r   rK   r\   rJ   r   r  r   llvm_modprocrQ   tripler   r0   pathstotal_num_warpsrets                     r    	make_llirCUDABackend.make_llir?  s   27KK<L<LM__S[[)
--b177;//3$$R(11"5::2>99"=++BKH''+b!11"5;;B?''+b!$$R(  22MM&&r*
s,,.((km m>>#/&z2)9)9:&x@##H-.5.A.AB.AltT.AEB!!(2Xt'7'78 **+@A&$3H[! --l; # 0 01I J*-*:*:;[*\&'+.+;+;<a+b'((m
# Cs   Q!c           	         [        X0R                  R                  5      nSn[        U5      n[	        X0R                  R                  5      n[
        R                  " XXx/ UR                  S5      n	[        R                  " SU	5      n
[        U
5      S:X  d   eU
S   US'   US-   SUS-   3n[        R                  " S	S
U 3U	[        R                  S9n	[        R                  " SSU 3U	[        R                  S9n	[        R                  " SSU	5      n	[        R                  R                  (       a  [!        S5        [!        U	5        U	$ )Nr:  Fz(.visible .entry ([a-zA-Z_][a-zA-Z0-9_]*)r:   r   r   r=   r6   z\.version \d+\.\d+z	.version )flagsz\.target sm_\d+z.target sm_z,\s*debug|debug,\s*r`   z // -----// NVPTX Dump //----- //)rL   r   rH   rc   rR   r   translate_to_asmrq   r   findalllensub	MULTILINEr
   r	   
dump_nvptxprint)r   r6  r   r  r\   rJ   rX  rW  rQ   r[  namess              r    make_ptxCUDABackend.make_ptx{  s   238H8HI&&z2[[%5%56##CSEYEY[`a

FL5zQ 8$b);r>*:;ff*i},EsRTR^R^_ff';zl)CSPRP\P\]ff+R5<<""45#J
r#   c                 :   [        5       R                  n[        R                  " SSSS9 n[        R                  " SSSS9 nUR	                  U5        UR                  5         UR                  S-   n[        R                  R                  (       a  SS	/OS/n	UR                  (       a  / OS
/n
[        U5      n[        R                  R                  (       a  SS/O/ nUR                  (       a  UR                  R                  S5      O/ nU/U	QU
QSPUQUQSU 3PUR                  PSPUPn [         R"                  " USSUS9  [$        R                  R'                  UR                  5      (       a   [$        R(                  " UR                  5        [$        R                  R'                  UR                  5      (       a   [$        R(                  " UR                  5        [-        US5       nUR/                  5       nS S S 5        [$        R                  R'                  U5      (       a  [$        R(                  " U5        S S S 5        S S S 5        W$ ! [         R*                   a  n[-        UR                  5       nUR/                  5       nS S S 5        O! , (       d  f       O= f[$        R                  R'                  UR                  5      (       a   [$        R(                  " UR                  5        UR0                  S:X  a  SnO3UR0                  S[2        R4                  -   :X  a  SnOSUR0                   3n[7        U SW SSR9                  U5       S35      eS nAff = f! , (       d  f       GNn= f! , (       d  f       GN>= f! , (       d  f       W$ = f)NFwz.ptx)deletemoderb   rz.logz.oz	-lineinfoz-suppress-debug-infoz--fmad=falsez--opt-level0 z-vz--gpu-name=z-oT)check	close_fdsstderr   z!Internal Triton PTX codegen error   z`ptxas` raised SIGSEGVz`ptxas` failed with error code z
`ptxas` stderr:
z
Repro command: 
rT   )r*   r0   tempfileNamedTemporaryFilewriteflushr   r
   rH  rI  rq   rc   r	   disable_ptxas_optro   rB   r.   r   osexistsremoveCalledProcessErrorrU   rX   
returncodesignalSIGSEGVr   r   )r   r6  r   r  r\   r(   fsrcflogfbin	line_infofmadrH   disable_optptx_extra_options	ptxas_cmdelog_filelogerrorrZ   r   s                        r    
make_cubinCUDABackend.make_cubin  s     ((COSW''u3vNRVJJsOJJL99t#DAFARARAdAd&<=kvjwI--2N3CD*:6D 38,,2P2P=#.VXK ?Boo 5 5c :SU !$(*.1<?PT_`d_eRfhlhqhqswILydS77>>$)),,IIdii(77>>$)),,IIdii($ dD!Q "ww~~d##		$[ O P^ + 00 L$))_"--/C %__77>>$)),,IIdii(<<3&?E\\S6>>%994E=all^LE E7 +558E :33688I3F2Gr"K L LL" "!U ON PO^ s{   NC-M90B3I#M9/M' AM9NM$(M=J		M
J%!B>MM$$M9'
M61M99
N	N
Nc                   ^ ^^ T R                  TR                  5      mU[        R                  :X  a  UUU 4S jUS'   UUU 4S jUS'   OU[        R                  :X  a  UUU 4S jUS'   UUU 4S jUS'   UU 4S jUS	'   UU 4S
 jUS'   g )Nc                 *   > TR                  XTT5      $ r'   )r  r6  r   r\   rK   r   s     r    <lambda>(CUDABackend.add_stages.<locals>.<lambda>  s    4>>#QXZd3er#   r   c                 *   > TR                  XTT5      $ r'   )r3  r  s     r    r  r    s    DOOCSZ\f4gr#   ttgirc                 *   > TR                  XTT5      $ r'   )r7  r  s     r    r  r    s    DNN3RY[e4fr#   c                 *   > TR                  XTT5      $ r'   )r\  r  s     r    r  r    s    t~~cWV`/ar#   llirc                 R   > TR                  XTTR                  R                  5      $ r'   )rh  r   rH   r6  r   rK   r   s     r    r  r    s    dmmC7TXT_T_TdTd.er#   ptxc                 R   > TR                  XTTR                  R                  5      $ r'   )r  r   rH   r  s     r    r  r    s    wX\XcXcXhXh0ir#   r   )r   rH   r   TRITONGLUON)r   stagesrK   r   r\   s   ` ` @r    
add_stagesCUDABackend.add_stages  se    %%gll3
x&eF6NgF7O'fF7Oaveuiwr#   c                 L    [        5       nU SU R                  R                   3$ )Nr   )r4   r   rH   )r   r3   s     r    r   CUDABackend.hash  s&    #%!DKK,,-..r#   )r   )r   r   r   r   staticmethodr   r   r   r@   r   r   r   r   r   r   r   r   r   r   r  r3  r7  r\  rh  r  r  	functools	lru_cacher   r   __classcell__)r   s   @r    r   r      s    (	 ( (#$# $"y "T "#S #,
>S*_ 5 >"    H HT:x,1f	j / /r#   r   ),triton.backends.compilerr   r   r   triton._C.libtritonr   r   r   r	   tritonr
   triton.runtime.errorsr   dataclassesr   r  typingr   r   r   r   typesr   rV   r   rw  r  r|  r.   pathlibr   r%   
NvidiaToolr*   r  r4   r$   rG   rL   rR   r[   rc   rf   r   r)   r#   r    <module>r     s%   E E 8 8  , !  - -   	   	  # #5##    iS i i$  
 
 
 T4 4
& & $'? '? '?TD/+ D/r#   