+ "ݳj3aRt0t^RIt^RIHt^RIHt^RIHtHtH t ^RI t ]!RR7!RR44t Rs ]^k]!4t R R ltR R ltRR Rllt^] P"3RRllt^] P"3RRlltRRRlltR#)N) dataclass)Lock)CallableOptionalTupleT)frozenc,a]tRt^toV3RltRtVtR#)_BlockwiseFp8Implsc<V^8dQh/S[;R&S[S[,;R&S[S[,;R&S[S[,;R&S[S[,;R&#) availablegemm act_quant weight_quantweight_dequant)boolrr)format __classdict__s"f/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torchao/kernel/blockwise_quantization.py __annotate___BlockwiseFp8Impls.__annotate__sWO 8 !!  8$$  X&& N)__name__ __module__ __qualname____firstlineno____annotate_func____static_attributes____classdictcell__)rs@rr r srr c$V^8dQhR\/#r returnr )rs"rrrs  - rc\pVf0\;_uu_4\pVf \4pVsRRR4V#V# +'giT#;i)N) _triton_impls_triton_impls_lock_build_triton_impls)implss r_get_triton_implsr*sI E }  !E}+- %  L5L   Ls < A c$V^8dQhR\/#r"r$)rs"rrr(sAA/Arc a a a ^RIHpV!4'g\RRRRR4#^RIo ^RIHo ^RIHpRUUUu.uF$pRFpRFpV!RVRV/V^R7NK K K& ppppS PV.ROR7S PV 3R lV 3R ll44o \PPR RR 7RR V V 3Rlll4pVPRRl4pS PV 3RlV 3Rll4pS PV 3RlV 3Rll4p S PV 3RlV 3Rll4p \RVVV V R7#uupppi)r) has_tritonFN)Config BLOCK_SIZE_M BLOCK_SIZE_N) num_stages num_warps)configskeyc <V^8dQhRSPRSPRSPRSPRSPRSP/#)r NKM_BUCKETr/r0 BLOCK_SIZE_K constexpr)rtls"rr)_build_triton_impls..__annotate__Csa*'*' <<*' <<*',,*'ll*'ll*'ll*'rc <SP^R7p SP^R7p SPW{4pW,SP^V 4,V,pW,SP^V 4,V,pSP^V 4pWR,V,,VR,,pVVR,V,,VR,,pW?V,,pVVV ,V,,pSPW3SPR7p\ V4FpSP VVR,VVV ,, 8RR7pSP VVR,VVV ,, 8RR7pSP V4pSP V4pVSPVV4VR,,VR,,, pVV , pVV , pV^, pV^, pK VPVPP4pW,SP^V 4,pW,SP^V 4,pW/R,V,,VR,,pVR,V8VR,V8,pSPVVVR7R#) raxisNdtypeg)maskotherrCNNNNNrG) program_idcdivarangezerosfloat32rangeloaddottorB element_tystore) a_ptrb_ptrc_ptra_s_ptrb_s_ptrMr6r7r8r/r0r9pid_mpid_nkoffs_moffs_noffs_ka_ptrsb_ptrsa_s_ptrsb_s_ptrs accumulatoriaba_sb_scc_ptrsrCr<s &&&&&&&&&&&& rblockwise_fp8_gemm_kernel6_build_triton_impls..blockwise_fp8_gemm_kernel?sD$ 1 % 1 % GGA $&1l)CCqH&1l)CCqH1l+1,,vg>1,,vg>aZ'f 499hh ;2::hN qAVG_q1|;K7K%KSVWAVG_q1|;K7K%KSVWA''(#C''(#C 266!Q<#g,6WE EK l "F l "F MH MH NN5;;11 2% !\(BB% !\(BB1,,vg>w!#w!(;< &rzao::blockwise_fp8_gemm) mutates_argsc V^8dQhR\PR\PR\PR\PR\R\P/#r rfrhrgri block_sizer#torchTensorint)rs"rrr=psX << \\ <<\\     rc :<a a VP4'gQhVP4'gQhVP4'gQhVP4'gQhVPR4pVP4V,o VP^4o \P!\P !S 44pVP !.VP4RROS N5R\P/pV V V 3RlpS V,!WWqVS S WVVR7 V#)NrBch<SPSVR,4SPSVR,43#)r/r0rJ)METArYr6tritons&rE_build_triton_impls.._blockwise_fp8_gemm_op..s. KK4/ 0 KK4/ 0 r)r9) is_contiguoussizenumelmathceillog2 new_emptyrsbfloat16) rfrhrgrirqr7r8rjgridrYr6rlr{s &&&&& @@r_blockwise_fp8_gemm_op3_build_triton_impls.._blockwise_fp8_gemm_opos          """"  """" FF2J GGIN FF1I99TYYq\* KK @#2 @ @ @  "$' !#q!Qz rcVP^4pVP!.VP4RROVN5R\P/pV#)rNrBr~)rrrsr)rfrhrgrirqr6rjs&&&&& r__build_triton_impls.._s? FF1I KK @#2 @ @ @rc2<V^8dQhRSP/#r BLOCK_SIZEr:)rr<s"rrr=s!!)+!rc<S P^R7pWC,S P^V4,pS PW,4PS P4pS P S P V44R, pWg, pVPVPP4pS PW,V4S PW$,V4R#)a Quantizes the input tensor `x_ptr` and stores the result in `y_ptr` and the scaling factor in `s_ptr`. Args: x_ptr (triton.Pointer): Pointer to the input tensor. y_ptr (triton.Pointer): Pointer to the output tensor where quantized values will be stored. s_ptr (triton.Pointer): Pointer to the output tensor where scaling factors will be stored. BLOCK_SIZE (tl.constexpr): The size of the block to be processed by each program instance. Returns: None r?|@N) rIrKrOrQrMmaxabsrBrRrS) x_ptry_ptrs_ptrrpidoffsxsyr<s &&&& r$_fp8_blockwise_act_quant_kernel_implA_build_triton_impls.._fp8_blockwise_act_quant_kernel_impls mmm#"))Az":: GGEL ! $ $RZZ 0 FF266!9  % E DD'' ( q! a rc2<V^8dQhRSP/#rr:)rr<s"rrr=s///1||/rc<SP^R7pSP^R7pSPWE4pWe,SP^V4,p Wu,SP^V4,p V R,V,V R,,p V R,V8V R,V8,p SPW ,V R7P SP 4p SP SPV 44R, pW, pVP VPP4pSPW,WR7SPW&V,,V,V4R#)a" Quantizes the input tensor `x_ptr` and stores the result in `y_ptr` and the scaling factors in `s_ptr`. Args: x_ptr (tl.pointer): Pointer to the input tensor. y_ptr (tl.pointer): Pointer to the output tensor where quantized values will be stored. s_ptr (tl.pointer): Pointer to the output tensor where scaling factors will be stored. M (int): Number of rows in the weight matrix. N (int): Number of columns in the weight matrix. BLOCK_SIZE (tl.constexpr): The size of the block to be processed by each program instance. r?NrErrFrH) rIrJrKrOrQrMrrrBrRrS)rrrrYr6rrZr[nr]r^rrCrrrr<s&&&&&& r'_fp8_blockwise_weight_quant_kernel_implD_build_triton_impls.._fp8_blockwise_weight_quant_kernel_impls 1 % 1 % GGA "#bii:&>>#bii:&>>g"VG_4w!#w!(;< GGELtG , / / ; FF266!9  % E DD'' ( q, "U*A.rc2<V^8dQhRSP/#rr:)rr<s"rrr=s--/1||-rcT<SP^R7pSP^R7pSPWE4pWe,SP^V4,p Wu,SP^V4,p V R,V,V R,,p V R,V8V R,V8,p SPW ,V R7P SP 4p SPWV,,V,4pW,pSP W+,WR7R#)a Dequantizes weights using the provided scaling factors and stores the result. Args: x_ptr (tl.pointer): Pointer to the quantized weights. s_ptr (tl.pointer): Pointer to the scaling factors. y_ptr (tl.pointer): Pointer to the output buffer for dequantized weights. M (int): Number of rows in the weight matrix. N (int): Number of columns in the weight matrix. BLOCK_SIZE (tl.constexpr): Size of the block for tiling. Returns: None r?NrErFrH)rIrJrKrOrQrMrS)rrrrYr6rrZr[rr]r^rrCrrrr<s&&&&&& r)_fp8_blockwise_weight_dequant_kernel_implF_build_triton_impls.._fp8_blockwise_weight_dequant_kernel_impls$ 1 % 1 % GGA "#bii:&>>#bii:&>>g"VG_4w!#w!(;< GGELtG , / / ; GGEAI%- . E q,rT)r rrrr) @)rrr))r6r7r8r9rr) torch.utils._tritonr-r r{triton.languagelanguager.autotunejitrslibrary custom_op register_fake)r-r.block_mblock_nr1fp8_gemm_configsrrrrrrlr<r{s @@@rr(r((s. <<!%tT4@@ )  )G$G&J  Wng >! ' %  )   __ &L ZZ*'*'*'X ]]5BGH2))*  ZZ!!!0 ZZ///8 ZZ---:  #6<@  _ s*Ec V^8dQhR\PR\PR\PR\PR\R\P/#rprr)rs"rrrsX 2 2 || 2  2 || 2  2  2  \\ 2rcv\4pVP'g \R4hVPWW#V4#)unsupported without triton)r*r AssertionErrorr)rfrhrgrirqr)s&&&&& rblockwise_fp8_gemmrs4  E ???9:: ::aaj 11rc V^8dQhR\PR\R\PR\\P\P3,/#)r rrqrBr#)rsrtrurBr)rs"rrrsF## ||#!$#38;;# 5<< %&#rcRaa\4pVP'g \R4h^RIoSP 4'gQR4hSP R 4V,^8Xg QRV R24hV\ P\ P39gQR4h\ P!SVR7pSP!.SP 4RR OSP R 4V,N5R\ P/pVV3R lpVPV,!SWEVR 7WE3#) a Quantizes the input tensor `x` using block-wise quantization with block size being BLOCK_SIZEx1. Args: x (torch.Tensor): The input tensor to be quantized. Must be contiguous and its last dimension size must be divisible by `block_size`. block_size (int, optional): The size of the blocks to be used for quantization. Default is 128. dtype (torch.dtype, optional): The dtype to use for the quantized tensor. Default is `torch.float8_e4m3fn`. Returns: Tuple[torch.Tensor, torch.Tensor]: A tuple containing: - The quantized tensor with dtype `dtype`. - A tensor of scaling factors with dtype `torch.float32`. rNInput tensor must be contiguousz@Last dimension size must be divisible by block_size (block_size=)6dtype must be torch.float8_e4m3fn or torch.float8_e5m2rArBcT<SPSP4VR,43#r)rJr)metar{rs&rr|)fp8_blockwise_act_quant..sQWWY\0BCErrr~) r*r rr{rrrs float8_e4m3fn float8_e5m2 empty_likerrMr)rrqrBr)rrrr{sf&& @rfp8_blockwise_act_quantrs"  E ???9:: ??  ???  66": "a ' J:,VWX '    @@@  %(A RQVVXcr]RAFF2J*$<REMMRA ED OOD!Qj9 4KrcV^8dQhR\PR\R\\P\P3,/#)r rrqr#)rsrtrur)rs"rrrs:'' ||'!$' 5<< %&'rc aaa \4pVP'g \R4h^RIo VP 4'gQR4hVP 4^8XgQR4hVP ^4V,^8XdVP ^4V,^8Xg QRV R24hV\P\P39gQR4hVP 4woo\P!WR7pVPSV,SV,\PR7pVVV 3R lpVPV,!WVSSVR 7WE3#) a9 Quantizes the given weight tensor using block-wise quantization with block size being BLOCK_SIZExBLOCK_SIZE. Args: x (torch.Tensor): The weight tensor to be quantized. block_size (int, optional): The block size to use for quantization. Defaults to 128. dtype (torch.dtype, optional): The dtype to use for the quantized tensor. Defaults to `torch.float8_e4m3fn`. Returns: Tuple[torch.Tensor, torch.Tensor]: A tuple containing: - The quantized weight tensor with dtype `dtype`. - A tensor of scaling factors with dtype `torch.float32`. rNrz#Input tensor must have 2 dimensionszABoth dimensions of x must be divisible by block_size (block_size=rrrAch<SPSVR,4SPSVR,43#rryrrYr6r{s&rr|,fp8_blockwise_weight_quant..A. AtL)* AtL)*rr)r*r rr{rdimrrsrrrrrMr) rrqrBr)rrrrYr6r{s &&& @@@rfp8_blockwise_weight_quantrs=  E ???9:: ??  ???  557a<>>>< 66!9z !Q &166!9z+AQ+F KJ.drrr) r*r rr{rrrrsrget_default_dtyper) rrrqr)rrrYr6r{s &&& @@@rfp8_blockwise_weight_dequantrIs"  E ???9:: ??  !2!2V4VV 2 557ars !,,  $'''/3 2V AH 2(+ATAT#N(+%2E2E'T   r