o 3/i @sdZddlmZddlZddlmZmZddlmZddlm Z m Z m Z ddl m Z mZmZGdd d eZd Zd Zd Zd ZdZe ddddededefddZe ddddedededejfddZe efddddedededefddZe ddddedefdd ZdS)!aqRounding 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)Float32Uint32)ir)arithllvmvector) dsl_user_opInt32Tc@seZdZdZdZdZdS) RoundingModeuRounding modes for epilogue dtype downconversion. RN — Round to nearest even (default hardware behavior) RS — Stochastic rounding (SM100+ only, BF16 output only) rN)__name__ __module__ __qualname____doc__RNRSrr[/lsinfo/ai/hellotax_ai/llm_service/venv_vllm/lib/python3.10/site-packages/quack/rounding.pyr sr lS$lW =lyn<l.vlocipabreturnc Cstjd}tj|t|j||dt|j||dgddddtjjd}t }t tj ||dg||d}t tj ||dg||d}||fS) 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)>rzL{ .reg .u64 prod; mul.wide.u32 prod, $2, $3; mov.b64 {$1, $0}, prod; }z =r,=r,r,rFhas_side_effectsis_align_stack asm_dialectrr ) rTypeparser inline_asmrir_value AsmDialectAD_ATTr i32cutlass extractvalue) rrrr struct_tyresulti32_tyhilorrr mul_wide_u32&s  r/ rand_bitsc CsVttjtt|j||dt|j||dt|j||dgddddtj j dS)zConvert 2 FP32 values to packed BF16x2 using stochastic rounding. Uses Blackwell PTX instruction: cvt.rs.satfinite.bf16x2.f32 dst, src_hi, src_lo, rand rz+cvt.rs.satfinite.bf16x2.f32 $0, $2, $1, $3;z=r,f,f,rFr) r(r rr#r r'rr$rr%r&)rrr0rrrrrcvt_f32x2_bf16x2_rs?s r1counterkeyn_roundscCst|}td}td}td}t|} td} tt} tt} tt} tt}t|D].}t|| ||d\}}t|| ||d\}}||A| A}||A| A}|}|}| | } | |} q,||||fS)aPhilox 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. rr)rPHILOX_ROUND_APHILOX_ROUND_B PHILOX_KEY_A PHILOX_KEY_Branger/)r2r3r4rrc0c1c2c3k0k1round_around_bkey_akey_b_hi_blo_bhi_alo_arrrphilox]s(     rIseedtidc Cst|j}|jd}|ddksJd||d}|ddks)Jd|tjj}tjj|g||d} tjj|gtj|d} t j | ||d} t |D]{} | d} | dd}t j |tjtj| ||d||d }t j |tjtj|||d||d }| d}| d}|dkrt|d >t|B}t|t|}||}tt|t||||d}t|j||d}t j|| tjtj| ||d||d } qNt j| | ||d}|S) zConvert 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. rz&requires even number of elements, got z>num_pairs must be divisible by 4 for stochastic rounding, got )rrr )positionrr)r VectorTypetypeshaper(BFloat16 mlir_typegetr r mlir_undefr9rextractelementrconstantrrIr1rr$ insertelementbitcast)src_vecrJrKrr src_vec_type num_elems num_pairs dst_mlir_type dst_vec_type i32_vec_typei32_vecpair_idxlo_idxhi_idxsrc_losrc_hi group_idx intra_idxr2 rand_batchentropy packed_i32packed_i32_valdst_vecrrrconvert_f32_to_bf16_srsV    ro)renumrr(rr cutlass._mlirrcutlass._mlir.dialectsrrrcutlass.cutlass_dslr r r r PHILOX_N_ROUNDS_DEFAULTr5r6r7r8tupler/r1intrIrorrrrsb      '