§
     ’Þj€.  ã            	       ó`  — d Z ddlZddlZddlmZ 	 ddlZddlmZ ddl	m
Z
 ddlmZmZ ddlmZ n# e$ r dZY nw xY wdZd	Zd
ZdZdZe�· eddd¦  «        Zej        dej        fd„¦   «         Zej        d„ ¦   «         Zej        d„ ¦   «         Z ej        g d¢¬¦  «        dej        dej        dej        dej        fd„¦   «         Z ej        g d¢¬¦  «        dej        dej        dej        dej        fd„¦   «         Zej        d„ ¦   «         Zd„ Z d„ Z!d„ Z"e�dd l#m$Z$  e$¦   «          dS dS )!a	  silu(group_norm(x)) of 16-bit NCHW activations, the GroupNorm + SiLU before each convolution of image VAEs, as PyTorch computes it under
autocast (GroupNorm and SiLU in FP32, one rounding to the output dtype), without its FP32 temporaries: PyTorch writes an FP32 copy of x, the
FP32 normalization and the FP32 sigmoid (two full-size FP32 tensors per chunk of groups, about 44 bytes moved per element), where these
kernels read x three times and write the output once (8 bytes per element) and allocate only the output.

Triton (compute capability 7.5 and up; BF16 read and written as bits on every GPU, so Turing and the others give the same values):
- each program sums a span of a group (contiguous in NCHW) in FP32 accumulators that carry their rounding errors (two-sum), reduced in FP64,
  and takes the largest |x|: the group's mean is as close to the exact one as FP64 rounding of exact partial sums allows (PyTorch: FP32
  Welford). x is summed times 2^-64 (exact) so that no 16-bit value can overflow the accumulators;
- a second pass sums the squares of (x - m32) * 2^-e the same way, m32 the mean rounded to FP32 and 2^e the power of two above
  max|x| + |m32| (no overflow nor loss whatever the magnitude); the variance corrects the rounding of the mean exactly
  (E[(x - m32)^2] - (m - m32)^2): no cancellation however large the mean against the spread;
- a third pass gives each program one channel's span, with scale = rstd * weight and the shift as scalars: y = (x - m32) * scale + shift
  (computed as ((x - m32) * 2^-e) * (scale * 2^e), the same values for normal numbers), then y * sigmoid(y) with the precise exponential
  and IEEE division (as PyTorch's FP32 sigmoid), rounded once to nearest even.
Activations below MIN_BYTES keep the PyTorch route: the three launches cost more there (a floor of ~0.1 ms on an RTX 4090, ~0.43 ms on
an RTX 2080 Ti under Windows), measured against the chunked route at the Flux VAE shapes. Partial sums are combined in a fixed order:
deterministic for a given GPU and configuration. The launch configuration (elements per load,
warps, pipelining stages, programs per SM) is tuned once per GPU, channels per group and dtype, on scratch output (x is only read); the
lengths, spans and channel counts are runtime values: a new resolution or tile size compiles nothing.é    Né   )Údispatch)Ú	libdevice)Úbf16_bits_to_fp32Úfp32_to_bf16_bits)ÚLaunchTuneri  € é€   ))é   é   é   é   )é   r   r   r   )é   r   é   r   )r
   r   r   r   )r   r   r   r   )r   r   r   r   )r
   r   r   r   i   i   Úwan2gp_group_norm_siluÚVAEÚBITSc                 ó”   — t          j        | |z   |d¦  «        }|rt          |¦  «        n|                     t           j        ¦  «        S )Nr   )ÚtlÚloadr   ÚtoÚfloat32)ÚXÚindexÚmaskr   Úxs        úFC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\group_norm_silu.pyÚ_loadr   /   s@   € åŒG�A˜‘I˜t QÑ'Ô'ˆØ'+ÐAÕ  Ñ#Ô#Ð#°·²µb´jÑ1AÔ1AÐAó    c                 óš   — t          j        dd¦  «        }t          j        t          j        | ||z  z   |z   ||k     d¦  «        d¬¦  «        S )Nr   r	   ç        ©Úaxis)r   ÚarangeÚsumr   )ÚPARTIALÚgroupÚSPANSÚspans       r   Ú_group_totalr*   4   sI   € åŒy˜˜CÑ Ô ˆÝŒv•b”g˜g¨°©Ñ5¸Ñ<¸dÀUºlÈCÑPÔPÐWXÐYÑYÔYÐYr   c                 ó  — t          j        dd¦  «        }t          | ||¦  «        dz  |                     t           j        ¦  «        z  }|                     t           j        ¦  «        }t          j        t          j        |||z  z   |z   ||k     d¦  «        d¬¦  «                             t           j        ¦  «        t          j        |¦  «                             t           j        ¦  «        z   }t          j	        t          j
        |                     t           j        d¬¦  «        dz	  d	z  d
z
  d¦  «        d¦  «                             t           j        ¦  «        }	||d|	z
  dz                       t           j        d¬¦  «        |	fS )zlThe group's mean (FP64), its rounding to FP32 and the power of two 2^-e above max|x| + |m32| (FP32, with e).r   r	   g      ðCr!   r"   T©Úbitcasté4   iÿ  iþ  iˆÿÿÿéx   é   é   )r   r$   r*   r   Úfloat64r   Úmaxr   ÚabsÚminimumÚmaximumÚint64Úint32)
ÚSUMSÚMAXESr'   r(   ÚLENGTHr)   ÚmeanÚmean32ÚboundÚes
             r   Ú_group_statsr@   9   sM  € õ Œy˜˜CÑ Ô ˆÝ˜D %¨Ñ/Ô/Ð2HÑHÈ6Ï9Ê9ÕUWÔU_ÑK`ÔK`Ñ`ˆØ—’�œÑ$Ô$ˆÝ”•r”w˜u u¨u¡}Ñ4°tÑ;¸TÀEº\È3ÑOÔOÐVWÐXÑXÔX×[Ò[Õ\^Ô\fÑgÔgÕjlÔjpÐqwÑjxÔjx×j{Òj{Õ|~ô  }Gñ  kHô  kHñ  HˆÝŒJ•r”z E§H¢H­R¬X¸t HÑ$DÔ$DÈÑ$JÈeÑ#SÐW[Ñ"[Ð]aÑbÔbÐdgÑhÔh×kÒkÕlnÔltÑuÔuˆØ�V˜s Q™w¨2™o×1Ò1µ"´*ÀdÐ1ÑKÔKÈQÐNÐNr   )r;   ÚSPANr(   )Údo_not_specializeÚCENTEREDÚBLOCKÚSTAGESc           	      óä  — t          j        d¦  «        t          j        d¦  «                             t           j        ¦  «        }}|                     t           j        ¦  «        |z  }t          j        ||z   |¦  «        }| ||z  z   }|rt          |||||¦  «        \  }}}}||z  }nd\  }}t          j        |ft           j        ¬¦  «        }t          j        |ft           j        ¬¦  «        }t          j        |ft           j        ¬¦  «        }t          j        d|||	¬¦  «        D ]¢}||z   t          j	        d|¦  «        z   }||k     }t          ||||
¦  «        }|r"t          j        |||z  |z
  d¦  «        }||z  }n,t          j        |t          j        |¦  «        ¦  «        }||z  }||z   }||z
  }||||z
  z
  ||z
  z   z  }|}Œ£t          j        |||z  z   |z   t          j        |                     t           j        ¦  «        |                     t           j        ¦  «        z   d¬¦  «        ¦  «         |s4t          j        |||z  z   |z   t          j        |d¬¦  «        ¦  «         d S d S )Nr   r   )g      ð;r!   )Údtype©Ú
num_stagesr!   r"   )r   Ú
program_idr   r7   r5   r@   Úzerosr   Úranger$   r   Úwherer6   r4   Ústorer%   r2   r3   )r   r&   r9   r:   r;   rA   r(   rC   rD   rE   r   r)   r'   ÚstartÚendÚbaseÚ_r=   ÚunitÚshiftÚhiÚloÚtopÚoffsetr   Úinsider   ÚtotalÚbacks                                r   Ú_sums_kernelr\   C   si  € õ ”m AÑ&Ô&­¬°aÑ(8Ô(8×(;Ò(;½B¼HÑ(EÔ(EˆeˆØ—’�œÑ!Ô! DÑ(ˆÝŒj˜ ™ vÑ.Ô.ˆØ�5˜6‘>Ñ!ˆØð 	5Ý!-¨d°E¸5À%ÈÑ!PÔ!PÑˆAˆv�t˜QØ˜T‘MˆEˆEà4‰KˆD�%ÝŒX�u�h¥b¤jÐ1Ñ1Ô1ˆÝŒX�u�h¥b¤jÐ1Ñ1Ô1ˆÝŒh˜�x¥r¤zÐ2Ñ2Ô2ˆÝ”h˜q $¨¸&ÐAÑAÔAð 	ð 	ˆFØ˜F‘N¥R¤Y¨q°%Ñ%8Ô%8Ñ8ˆEØ˜S’[ˆFÝ�d˜E 6¨4Ñ0Ô0ˆAØð Ý”H˜V Q¨¡X°Ñ%5°sÑ;Ô;�Ø˜‘E��å”j ¥b¤f¨Q¡i¤iÑ0Ô0�Ø˜‘H�Ø˜‘FˆEØ˜2‘:ˆDØ�2˜ ™Ñ&¨1¨t©8Ñ4Ñ4ˆBØˆBˆBÝ
Œ�˜5 5™=Ñ(¨4Ñ/µ´¸¿º½b¼jÑ8IÔ8IÈBÏEÊEÕRTÔR\ÑL]ÔL]Ñ8]ÐdeÐ1fÑ1fÔ1fÑgÔgÐgØð 	HÝŒH�U˜U U™]Ñ*¨TÑ1µ2´6¸#ÀAÐ3FÑ3FÔ3FÑGÔGÐGÐGÐGð	Hð 	Hr   )ÚHWÚCHUNKr(   r;   ÚCHANNELSÚGROUP_CHANNELSÚOUT_BITSc           	      ó”  — t          j        d¦  «        t          j        d¦  «                             t           j        ¦  «        }}||z  }||z  ||z  z  ||z  z   }t	          ||||	|
¦  «        \  }}}}d|z                        t           j        ¦  «        dz                       t           j        d¬¦  «        }||                     t           j        ¦  «        z
  }t          j        t          |||	¦  «        |
                     t           j        ¦  «        z  |z  |z  ||z  z
  d¦  «        }t          j        ||z   ¦  «                             t           j        ¦  «        t          j	        ||z   ¦  «        z  |z                       t           j
        ¦  «        }t          j        ||z   ¦  «        ||z  }}|                     t           j        ¦  «        |z  } t          j        | |z   |¦  «        }!||z  }"t          j        d|||¬¦  «        D ]¾}#| |#z   t          j        d|¦  «        z   }$|$|!k     }%t          | |"z   |$|%|¦  «        }&|&|z  |z
  |z  |z   }'|'t          j        d	d	t!          j        |' ¦  «        z   ¦  «        z  }'t          j        ||"z   |$z   |rt'          |'¦  «        n|'                     |j        j        ¦  «        |%¦  «         Œ¿d S )
Nr   r   iÿ  r.   Tr,   r!   rH   g      ð?)r   rJ   r   r7   r@   r2   r6   r*   r   Úsqrtr   r5   rL   r$   r   Údiv_rnr   ÚexprN   r   rG   Ú
element_ty)(r   ÚYr9   ÚSQUARESr:   ÚWÚBr]   r^   r(   r;   r_   r`   ÚEPSrD   rE   r   ra   ÚchunkÚrowÚchannelr'   r<   r=   rS   r?   ÚpowerÚerrorÚvarianceÚscalerT   ÚcentrerO   rP   rQ   rX   r   rY   r   Úys(                                           r   Ú_apply_kernelru   e   s–  € õ ”] 1Ñ%Ô%¥r¤}°QÑ'7Ô'7×':Ò':½2¼8Ñ'DÔ'DˆsˆØ˜‘.ˆØ˜‘ X°Ñ%?Ñ@À7ÈnÑC\Ñ\ˆÝ ,¨T°5¸%ÀÈÑ OÔ OÑˆˆf�d˜AØ˜‘(—’�rœxÑ(Ô(¨BÑ.×2Ò2µ2´:ÀtÐ2ÑLÔLˆØ�v—y’y¥¤Ñ,Ô,Ñ,ˆÝ”:�l¨7°E¸5ÑAÔAÀFÇIÂIÍbÌjÑDYÔDYÑYÐ\aÑaÐdiÑiÐlqÐtyÑlyÑyÐ{~ÑÔˆÝ”˜˜W™Ñ%Ô%×(Ò(­¬Ñ4Ô4µr´w¸xÈ#¹~Ñ7NÔ7NÑNÐQVÑV×ZÒZÕ[]Ô[eÑfÔfˆÝœ  G¡Ñ,Ô,¨f°t©mˆvˆØ—’�œÑ"Ô" UÑ*ˆÝŒj˜ ™¨Ñ+Ô+ˆØ�R‰xˆÝ”h˜q %¨¸6ÐBÑBÔBð 	oð 	oˆFØ˜F‘N¥R¤Y¨q°%Ñ%8Ô%8Ñ8ˆEØ˜S’[ˆFÝ�a˜$‘h  v¨tÑ4Ô4ˆAØ�T‘˜FÑ" eÑ+¨eÑ3ˆAØ•B”I˜c 3­¬¸°rÑ):Ô):Ñ#:Ñ;Ô;Ñ;ˆAÝŒH�Q˜‘X Ñ%¸xÐ'eÕ'8¸Ñ';Ô';Ð';ÈQÏTÊTÐRSÔRYÔRdÑMeÔMeÐgmÑnÔnÐnÐnð	oð 	or   c                 óJ   — t           j                             | ¦  «        j        S )N)ÚtorchÚcudaÚget_device_propertiesÚmulti_processor_count)r   s    r   Ú_smsr{   }   s   € åŒ:×+Ò+¨EÑ2Ô2ÔHÐHr   c                 ó®  — |\  }}}}|||z  z  ||z  }}t          dt          t          t          j        ||¦  «        ||z  |z  ¦  «        ¦  «        }t          j        t          j        ||¦  «        |¦  «        |z  }t          j        ||¦  «        }t          j        d||ft
          j        | j        ¬¦  «        }|d          	                    t
          j
        ¦  «        }|d         df|d         dffD ]/\  }}t          ||f         | ||d         ||||||||	|d¬¦  «         Œ0t          dt          t          j        ||¦  «        ||z  ||z  z  ¦  «        ¦  «        }t          j        t          j        ||¦  «        |¦  «        |z  }t          t          j        ||¦  «        ||z  f         | ||d         |d         |||||||||||||	|
|¬	¦  «         d S )
Nr   r   ©rG   Údevicer   r   FT)Ú	num_warpsÚenable_fp_fusion)r   )r3   ÚminÚ	MAX_SPANSÚtritonÚcdivrw   Úemptyr2   r~   Úviewr   r\   ru   )ÚsourceÚtargetÚweightÚbiasÚbatchÚchannelsÚgroup_channelsÚhwÚepsÚbitsÚout_bitsÚconfigÚsmsÚblockÚwarpsÚstagesÚper_smÚgroupsÚlengthÚspansr)   ÚsumsÚmaxesÚpartialÚcenteredÚchunksrl   s                              r   Ú_launchr    ‚   s  € Ø#)Ñ €Eˆ5�&˜&Ø˜h¨.Ñ8Ñ9¸>ÈBÑ;NˆF€FÝ�•3•y¥&¤+¨f°eÑ"<Ô"<¸fÀs¹lÈfÑ>TÑUÔUÑVÔV€EÝŒ;•v”{ 6¨5Ñ1Ô1°5Ñ9Ô9¸EÑA€DÝŒK˜ Ñ%Ô%€EÝŒ;˜˜6 5Ð)µ´ÀvÄ}ÐUÑUÔU€DØ�ŒG�LŠL�œÑ'Ô'€EØ# Aœw¨Ð.°°a´¸$°Ð@ð >ð >Ñˆ�Ý�e˜V�_Ô% f¨g°t¸A´wÀÀvÈtÐUZÐ\dÐfkÐmsÐuyð  FKØ7<ð	>ñ 	>ô 	>ð 	>ð 	>å�•C�œ B¨Ñ.Ô.°¸±ÀÈÑAQÑ0RÑSÔSÑTÔT€FÝŒK�œ B¨Ñ/Ô/°Ñ7Ô7¸%Ñ?€EÝ•6”;˜r 5Ñ)Ô)¨5°8Ñ+;Ð<Ô=¸fÀfÈdÐSTÌgÐW[Ð\]ÔW^Ð`eÐgmÐosÐuwÐy~ð  AFð  HNØ>FÈÐX[Ð]bÐdjÐlpÐrzð  GLðMñ Mô Mð Mð Mð Mr   c                 ó†  — t           duo¸|                      ¦   «         dk    o |                      ¦   «         oŒ|                      ¦   «         |                      ¦   «         z  t
          k    oZ| j        d         |z  dk    oF|t          j        t          j	        fv o,t          j        | t          j        t          j	        fd¬¦  «        S )zƒWhether the kernels run silu(group_norm(x)) to dtype: contiguous NCHW BF16 or FP16 x of MIN_BYTES or more and output, whole groups.Nr   r   r   )é   é   )ÚdtypesÚcapability_floor)rƒ   ÚdimÚis_contiguousÚnumelÚelement_sizeÚ	MIN_BYTESÚshaperw   Úbfloat16Úfloat16r   Úcuda_inference)r   Ú
num_groupsrG   s      r   Úeligibler°   “   sâ   € å˜$Ðð Y 1§5¢5¡7¤7¨a¢<ð Y°A·O²OÑ4EÔ4Eð YÈ!Ï'Ê'É)Ì)ÐVW×VdÒVdÑVfÔVfÑJfÕjsÒJsð YÐxyÔxð  ABô  yCð  FPñ  yPð  TUò  yUð YØ�%œ.­%¬-Ð8Ð8ðYÝ=EÔ=TÐUVÕ`eÔ`nÕpuÔp}Ð_~ð  RXð  >Yñ  >Yô  >YðZr   c                 óx  ‡ ‡‡‡‡‡‡‡‡‡‡‡‡— ‰ j         \  }}}}	||	z  |‰z  cŠŠt          j        ‰ j         |‰ j        ¬¦  «        }
‰ j        t          j        k    |t          j        k    cŠŠ‰r‰                      t          j        ¦  «        n‰ ‰r|
                     t          j        ¦  «        n|
cŠŠ‰                     ‰ j        t          j	        ¬¦  «        ‰                     ‰ j        t          j	        ¬¦  «        cŠŠt          ‰ j        j        ¦  «        Š‰                      ¦   «         ‰                      ¦   «         z  Šˆˆˆˆˆˆˆˆˆˆˆˆˆ fd„}t                               d‰› d‰ j        › d|› �‰ j        t           |‰t"          k    ¦  «        }t%          ‰‰‰‰||‰‰‰‰‰t'          |¦  «        ‰¦  «         |
S )zsilu(group_norm(x.float(), num_groups, weight, bias, eps)) rounded to dtype, x borrowed (not modified); eligible(x, ...) first.r}   )r~   rG   c            
      ó¤  •‡ ‡— ‰	t           k    r‰n	‰d d…         Š ‰                      ¦   «         ‰                      ¦   «         z  t           k    rQ‰ d d …d t          dt           ‰                      ¦   «         ‰                      ¦   «         z  ‰z  z  ¦  «        ‰z  …f         Š t	          j        ‰ j        ‰j        ‰j        ¬¦  «        Šˆˆˆˆˆˆˆ ˆˆ
ˆf
d„S )Nr   r}   c                 óf   •
— t          ‰‰‰
‰‰j        d         ‰j        d         ‰‰‰‰‰| ‰	¦  «        S )Nr   r   )r    r«   )r’   rŠ   r�   r�   r�   rŽ   r‘   ÚpartÚsampler“   r‰   s    €€€€€€€€€€r   ú<lambda>z2group_norm_silu.<locals>.scratch.<locals>.<lambda>©   sP   ø€ �g d¨F°F¸DÀ$Ä*ÈQÄ-ÐQUÔQ[Ð\]ÔQ^Ð`nÐprÐtwÐy}ð  @Hð  JPð  RUñ  Vô  V€ r   )	Ú_TUNE_MAX_BYTESr¨   r©   r3   rw   r…   r«   rG   r~   )r´   rµ   rŠ   r�   r�   r�   rŽ   r¯   r‘   Úsizer“   r‡   rˆ   r‰   r   s   @@€€€€€€€€€€€€€r   Úscratchz group_norm_silu.<locals>.scratch¤   s  øøø€ Ø¥Ò0Ð0ˆvˆv°f¸R¸a¸R´jˆØ�:Š:‰<Œ<˜$×+Ò+Ñ-Ô-Ñ-µÒ?Ð?Ø˜˜˜Ðz�C ¥?°t·z²z±|´|Àd×FWÒFWÑFYÔFYÑ7YÐ]gÑ7gÑ#hÑiÔiÐlzÑzÐzÐzÔ{ˆDÝ”˜TœZ¨v¬|ÀAÄHÐMÑMÔMˆð Vð  Vð  Vð  Vð  Vð  Vð  Vð  Vð  Vð  Vð  Vð  Vð  Vð  	Vr   Úgroup_norm_siluú|)r«   rw   r…   r~   rG   r¬   r†   Úint16r   r   r{   r   r¨   r©   Ú_TUNERr’   Ú_CANDIDATESÚ_TUNE_MIN_BYTESr    Útuple)r   r¯   r‰   rŠ   r�   rG   r‹   rŒ   ÚheightÚwidthÚoutr¹   r’   r�   r�   rŽ   r‘   r¸   r“   r‡   rˆ   s   `````        @@@@@@@@r   rº   rº   ™   sù  øøøøøøøøøøøøø€ à%&¤WÑ"€Eˆ8�V˜UØ %™¨°ZÑ)?Ð€BˆÝ
Œ+�a”g U°1´8Ð
<Ñ
<Ô
<€CØ”W¥¤Ò.°½¼Ò0G€N€Dˆ(Ø,0Ð7�Q—V’V�EœKÑ(Ô(Ð(°aÐRZÐ9c¸¿ºÅ%Ä+Ñ9NÔ9NÐ9NÐ`c€N€FˆFØ—9’9 A¤HµE´M�9ÑBÔBÀDÇGÂGÐSTÔS[ÕchÔcpÀGÑDqÔDq€L€FˆDÝ
ˆqŒxŒ~Ñ
Ô
€CØ�7Š7‰9Œ9�q—~’~Ñ'Ô'Ñ'€DðVð Vð Vð Vð Vð Vð Vð Vð Vð Vð Vð Vð Vð Vð Vð Vð Võ �]Š]Ð,°Ð.SÐ.SÀ!Ä'Ð.SÐ.SÈEÐ.SÐ.SÐUVÔU]Õ_jÐlsÐuyõ  ~Mò  vMñ  Nô  N€FÝˆF�F˜F D¨%°¸>È2ÈsÐTXÐZbÕdiÐjpÑdqÔdqÐsvÑwÔwÐwØ€Jr   )Ú!install_triton_compilation_logger)%Ú__doc__Ú	functoolsrw   Ú r   rƒ   Útriton.languageÚlanguager   Útriton.language.extrar   Ú	bf16_bitsr   r   Útriton_tunerr   ÚImportErrorrª   r‚   r¾   r¿   r·   r½   ÚjitÚ	constexprr   r*   r@   r\   ru   Úcacher{   r    r°   rº   Ú%shared.kernels.triton_compilation_logrÄ   © r   r   ú<module>rÓ      sÇ  ððið ið( Ð Ð Ð à €€€à Ð Ð Ð Ð Ð ðØ€M€M€MØ Ð Ð Ð Ð Ð Ø/Ð/Ð/Ð/Ð/Ð/Ø?Ð?Ð?Ð?Ð?Ð?Ð?Ð?Ø)Ð)Ð)Ð)Ð)Ð)Ð)øØð ð ð Ø€F€F€Fðøøøð €	Ø€	ð F€Ø€Ø€à	ÐØˆ[Ð1°1°eÑ<Ô<€Fà„ZðB B¤Lð Bð Bð Bñ „ZðBð „ZðZð Zñ „ZðZð „ZðOð Oñ „ZðOð €V„ZÐ"=Ð"=Ð"=Ð>Ñ>Ô>ðHÈRÌ\ð HÐbdÔbnð HÐxzô  yEð HØœ<ðHð Hð Hñ ?Ô>ðHðB €V„ZÐ"bÐ"bÐ"bÐcÑcÔcðoÐxzô  yEð oØ œlðoØ24´,ðoØJLÌ,ðoð oð oñ dÔcðoð. „ðIð Iñ „ðIðMð Mð Mð"Zð Zð Zðð ð ð. 
ÐØWÐWÐWÐWÐWÐWà%Ð%Ñ'Ô'Ð'Ð'Ð'ð Ðs   ’1 ±;º;