Ë
    ÄÓ^j¨  ã                   óÒ  — d Z ddlmZ ddlZddlmZ ddlmZmZ ddlm	Z	 ddl
mZmZmZ ddlmZmZmZ  G d„ d	e«      Zd
ZdZdZdZdZdZdZdZdZdZedddœdedej:                  defd„«       Zedddœdedej:                  defd„«       Zedddœdedede fd„«       Z!edddœdedededej                   fd„«       Z"eefdddœdeded e#de fd!„«       Z$edddœd"ed#efd$„«       Z%y)%aq  Rounding mode control and stochastic rounding primitives for GEMM epilogues.

Provides a RoundingMode enum for configuring how epilogues downconvert the
accumulator dtype (typically FP32) to the output dtype before storing to gmem.
Stochastic rounding (RS) uses the hardware cvt.rs.satfinite.bf16x2.f32 PTX
instruction and is only supported on Blackwell (SM100+) GPUs.
é    )ÚIntEnumN)ÚFloat32ÚUint32)Úir)ÚarithÚllvmÚvector)Údsl_user_opÚInt32ÚTc                   ó   — e Zd ZdZdZdZy)ÚRoundingModeu³   Rounding modes for epilogue dtype downconversion.

    RN â€” Round to nearest even (default hardware behavior)
    RS â€” Stochastic rounding (SM100+ only, BF16 output only)
    r   é   N)Ú__name__Ú
__module__Ú__qualname__Ú__doc__ÚRNÚRS© ó    úg/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_vendor/quack/rounding.pyr   r      s   „ ñð 
€BØ	
�Br   r   i  i  é   é   l   ±yn< l   S¢$ l   W= l   ¹yn< l   ….Ïv ©ÚlocÚipÚ	base_seedÚtile_coord_mnklÚreturnc                óh   — | |d   t         z  |d   t        z  z   |d   t        z  z   |t        z  z   z   S )Nr   r   é   )ÚEPILOGUE_SR_SEED_M_STRIDEÚEPILOGUE_SR_SEED_N_STRIDEÚEPILOGUE_SR_SEED_BATCH_STRIDEÚEPILOGUE_SR_SEED_SUBTILE_STRIDE©r   r   Úsubtile_idxr   r   s        r   Úepilogue_sr_seedr)   .   sU   € ð Ø˜ÑÔ6Ñ6Ø
˜!Ñ
Ô8Ñ
8ñ	9à
˜!Ñ
Ô<Ñ
<ñ	=ð Ô7Ñ
7ñ	8ñð r   c                ó0   — t        | t        z   ||||¬«      S )Nr   )r)   ÚEPILOGUE_SR_SEED_AUX_OUT_SALTr'   s        r   Úepilogue_aux_out_sr_seedr,   ?   s&   € ô ØÔ1Ñ1ØØØØôð r   ÚaÚbc          	      óÊ  — t         j                  j                  d«      }t        j                  |t        | «      j                  ||¬«      t        |«      j                  ||¬«      gdddd¬«      }t        j                  «       }t        j
                  t        j                  ||dg||¬«      «      }t        j
                  t        j                  ||dg||¬«      «      }||fS )	zuUnsigned 32b x 32b -> 64 wide multiply via PTX `mul.wide.u32`.

    Returns (hi, lo) as a pair of Uint32 values.
    z!llvm.struct<(i32, i32)>r   zL{
  .reg .u64 prod;
  mul.wide.u32 prod, $2, $3;
  mov.b64 {$1, $0}, prod;
}z	=r,=r,r,rF©Úhas_side_effectsÚis_align_stackr   r   )r   ÚTypeÚparser   Ú
inline_asmr   Úir_valuer   Úi32ÚcutlassÚextractvalue)	r-   r.   r   r   Ú	struct_tyÚresultÚi32_tyÚhiÚlos	            r   Úmul_wide_u32r?   Q   sÆ   € ô —‘—‘Ð8Ó9€IÜ�_‰_Øä�1‹I×Ñ 3¨2ÐÓ.Ü�1‹I×Ñ 3¨2ÐÓ.ð	
ð 	[ØØØô
€Fô �U‰U‹W€FÜ	�‰œ×)Ñ)¨&°&¸1¸#À3È2ÔNÓ	O€BÜ	�‰œ×)Ñ)¨&°&¸1¸#À3È2ÔNÓ	O€BØˆrˆ6€Mr   Ú	rand_bitsc                ó$  — t        j                  t        j                  t	        j
                  «       t        | «      j                  ||¬«      t        |«      j                  ||¬«      t        |«      j                  ||¬«      gdddd¬«      «      S )z¡Convert 2 FP32 values to packed BF16x2 using stochastic rounding.

    Uses Blackwell PTX instruction: cvt.rs.satfinite.bf16x2.f32 dst, src_hi, src_lo, rand
    r   z+cvt.rs.satfinite.bf16x2.f32 $0, $2, $1, $3;z=r,f,f,rFr0   )	r8   r   r   r5   r   r7   r   r6   r   )r-   r.   r@   r   r   s        r   Úcvt_f32x2_bf16x2_rsrB   i   sƒ   € ô �=‰=Ü�‰Ü�E‰E‹Gä˜“
×#Ñ#¨°Ð#Ó3Ü˜“
×#Ñ#¨°Ð#Ó3Ü�yÓ!×*Ñ*¨s°rÐ*Ó:ðð
 :ØØ"Ø ô	
óð r   ÚcounterÚkeyÚn_roundsc                ó®  — t        | «      }t        d«      }t        d«      }t        d«      }t        |«      }	t        d«      }
t        t        «      }t        t        «      }t        t        «      }t        t        «      }t        |«      D ]D  }t        ||||¬«      \  }}t        ||||¬«      \  }}||z  |	z  }||z  |
z  }|}|}|	|z   }	|
|z   }
ŒF ||||fS )a  Philox 4x32b counter-based random number generator.

    Given a 32b counter and a 32b key, returns four pseudo-random uint32 words
    produced by running n_rounds of the Philox 4x32 bijection. Each round
    performs two wide 32x32->64 multiplies with the Philox constants.
    r   r   )r   ÚPHILOX_ROUND_AÚPHILOX_ROUND_BÚPHILOX_KEY_AÚPHILOX_KEY_BÚranger?   )rC   rD   rE   r   r   Úc0Úc1Úc2Úc3Úk0Úk1Úround_aÚround_bÚkey_aÚkey_bÚ_Úhi_bÚlo_bÚhi_aÚlo_as                       r   Úphiloxr[   †   sã   € ô 
�‹€BÜ	�‹€BÜ	�‹€BÜ	�‹€BÜ	�‹€BÜ	�‹€Bä”^Ó$€GÜ”^Ó$€GÜ”<Ó €EÜ”<Ó €Eä�8‹_ò ˆÜ! " g°3¸2Ô>‰
ˆˆdÜ! " g°3¸2Ô>‰
ˆˆdØ�B‰Y˜‰^ˆØ�B‰Y˜‰^ˆØˆØˆØ�%‰ZˆØ�%‰Z‰ðð ˆr�2�rˆ>Ðr   ÚseedÚtidc                óø  — t        j                  | j                  «      }|j                  d   }|dz  dk(  s
J d|› �«       ‚|dz  }|dz  dk(  s
J d|› �«       ‚t        j
                  j                  }t         j                  j                  |g||¬«      }	t         j                  j                  |gt        j                  |¬«      }
t        j                  |
||¬«      }t        |«      D �]t  }|dz  }|dz  dz   }t        j                  | t        j                  t        j                  |||¬«      ||¬	«      }t        j                  | t        j                  t        j                  |||¬«      ||¬	«      }|dz  }|dz  }|dk(  rMt	        j                   |d
z  «      t	        j                   |«      z  }t#        |t	        j                   |«      «      }|   }t%        t'        |«      t'        |«      |||¬«      }t	        j                  |«      j)                  ||¬«      }t        j*                  ||t        j                  t        j                  |||¬«      ||¬	«      }�Œw t        j,                  |	|||¬«      }|S )z¿Convert an MLIR FP32 vector to BF16 with stochastic rounding.

    Processes elements in pairs using Philox PRNG for entropy and the hardware
    cvt.rs.satfinite.bf16x2.f32 instruction.
    r   é   z&requires even number of elements, got é   z>num_pairs must be divisible by 4 for stochastic rounding, got )r   r   r   )Úpositionr   r   é   )r   Ú
VectorTypeÚtypeÚshaper8   ÚBFloat16Ú	mlir_typeÚgetr   r   Ú
mlir_undefrK   r	   Úextractelementr   Úconstantr   r[   rB   r   r6   ÚinsertelementÚbitcast)Úsrc_vecr\   r]   r   r   Úsrc_vec_typeÚ	num_elemsÚ	num_pairsÚdst_mlir_typeÚdst_vec_typeÚi32_vec_typeÚi32_vecÚpair_idxÚlo_idxÚhi_idxÚsrc_loÚsrc_hiÚ	group_idxÚ	intra_idxrC   Ú
rand_batchÚentropyÚ
packed_i32Úpacked_i32_valÚdst_vecs                            r   Úconvert_f32_to_bf16_srr‚   ®   sA  € ô —=‘= §¡Ó.€LØ×"Ñ" 1Ñ%€IØ�q‰=˜AÒÐSÐ!GÈ	À{ÐSÓSÐØ˜Q‘€IØ�q‰=˜AÒð Ø
HÈÈÐTóÐô ×$Ñ$×.Ñ.€MÜ—=‘=×$Ñ$ i [°-ÀSÐ$ÓI€Lä—=‘=×$Ñ$ i [´%·/±/ÀsÐ$ÓK€LÜ�o‰o˜l°¸Ô;€Gä˜)Ó$ó !
ˆØ˜A‘ˆØ˜A‘ Ñ!ˆä×&Ñ&ØÜ—^‘^¤E§O¡O°VÀÈÔLØØô	
ˆô ×&Ñ&ØÜ—^‘^¤E§O¡O°VÀÈÔLØØô	
ˆð  ‘Mˆ	Ø˜q‘Lˆ	Ø˜Š>Ü—n‘n Y°"¡_Ó5¼¿¹ÀsÓ8KÑKˆGÜ ¬¯©¸Ó)=Ó>ˆJà˜YÑ'ˆÜ(¬°«¼'À&»/È7ÐX[Ð`bÔcˆ
ä Ÿ™ zÓ2×;Ñ;ÀÈÐ;ÓKˆÜ×&Ñ&ØØÜ—^‘^¤E§O¡O°XÀ3È2ÔNØØô
Šð7!
ôF �l‰l˜<¨°c¸bÔA€GØ€Nr   )&r   Úenumr   r8   Úcutlass.cuteÚcuter   r   Úcutlass._mlirr   Úcutlass._mlir.dialectsr   r   r	   Úcutlass.cutlass_dslr
   r   r   r   r#   r$   r%   r&   r+   ÚPHILOX_N_ROUNDS_DEFAULTrG   rH   rI   rJ   ÚCoordr)   r,   Útupler?   rB   Úintr[   r‚   r   r   r   ú<module>r�      sÌ  ðñõ ã Ý ß #Ý ß 6Ñ 6ß 5Ñ 5ô�7ô ð "Ð ØÐ Ø "Ð Ø"#Ð Ø *Ð àÐ à€Ø€Ø€Ø€ð ð 	ØòØðà—Z‘Zðð òó ðð  ð 	ØòØðà—Z‘Zðð òó ðð" Ø.2°tò �Fð ˜vð Àò ó ðð. ð 	ØòØðàðð ðð ‡]�]òó ðð8 ð ,ð$ð
 	Øò$Øð$à	ð$ð ð$ð ò$ó ð$ðN ð 	Øò?à
ð?ð 
ò?ó ñ?r   