iS3 .SSKrSSKJr SSKrSqSqSqSqSqSq Sr SS\RS\RS\RS \RS \ S \R4 S jjr S\R4S \RS \ S\RS \\R\R44SjjrS\R4S \RS \ S \\R\R44SjjrSS \RS\RS \ S \R4Sjjrg)N)TupleFc^ ^ ^ [(agSqSSKJn U"5(dSqgSqSSKm SSKJm SSKJn SVVVs/sH!nSHnSH nU"X#S .US S 9PM M M# nnnnT RU/S QS 9T RST RST RST RST RST RST R4 U 4Sjj55m [RRSSS9S$S[RS[RS[RS[RS[S[R4 U U 4Sjjj5nUR S$Sj5nUqT RS T R4U 4S!jj5nUqT RS T R4U 4S"jj5n U qT RS T R4U 4S#jj5n U qgs snnnf)%NTr) has_tritonF)Config) @)rr r )) BLOCK_SIZE_M BLOCK_SIZE_N) num_stages num_warps)NKM_BUCKET BLOCK_SIZE_K)configskeyrrrrrrc <>TRSS9n TRSS9n TRX{5nX-TRSU 5-U-nX-TRSU 5-U-nTRSU 5nXSS2S4U--USSS24-nUUSSS24U--USS2S4-nX?U--nUUU -U--nTRX4TRS9n[ U5HnTR UUSSS24UUU -- :SS9nTR UUSS2S4UUU -- :SS9nTR U5nTR U5nUTRUU5USS2S4-USSS24-- nUU - nUU - nUS- nUS- nM URURR5nX-TRSU 5-nX-TRSU 5-nX/SS2S4U--USSS24-nUSS2S4U:USSS24U:-nTRUUUS9 g)Nraxisdtypeg)maskotherr ) program_idcdivarangezerosfloat32rangeloaddottor element_tystore) a_ptrb_ptrc_ptra_s_ptrb_s_ptrMrrrrrrpid_mpid_nkoffs_moffs_noffs_ka_ptrsb_ptrsa_s_ptrsb_s_ptrs accumulatoriaba_sb_scc_ptrsr tls d/var/www/html/ai-image-ml/venv/lib/python3.13/site-packages/torchao/kernel/blockwise_quantization.pyblockwise_fp8_gemm_kernel4_lazy_init_triton..blockwise_fp8_gemm_kernel:st$ 1 % 1 % GGA $&1l)CCqH&1l)CCqH1l+41,,vdAg>a1,,vag>aZ'f 499hh ;2::hN qAVD!G_q1|;K7K%KSVWAVAtG_q1|;K7K%KSVWA''(#C''(#C 266!Q<#ag,6T1WE EK l "F l "F MH MH NN5;;11 2% !\(BB% !\(BB41,,vdAg>q$w!#tQw!(;< &zao::blockwise_fp8_gemm) mutates_argsr@rBrArC block_sizereturnc >^ ^ UR5(deUR5(deUR5(deUR5(deURS5nUR5U-m URS5m [R"[R "T 55nUR "/UR5SSQT P7S[R06nU U U 4SjnT U"XXqUT T XVUS9 U$)NrrcX>TRTUS5TRTUS54$)Nrrr$)METAr3rtritons rGC_lazy_init_triton.._blockwise_fp8_gemm_op..{s. KK4/ 0 KK4/ 0 rJ)r) is_contiguoussizenumelmathceillog2 new_emptytorchbfloat16) r@rBrArCrMrrrDgridr3rrHrTs @@rG_blockwise_fp8_gemm_op1_lazy_init_triton.._blockwise_fp8_gemm_opjs          """"  """" FF2J GGIN FF1I99TYYq\* KK @#2 @ @ @  "$' !#q!Qz rJcURS5nUR"/UR5SSQUP7S[R06nU$)NrrPr)rXr]r^r_)r@rBrArCrMrrDs rG__lazy_init_triton.._s? FF1I KK @#2 @ @ @rJ BLOCK_SIZEc>T RSS9nXC-T RSU5-nT RX-5RT R5nT R T R U55S- nXg- nURURR5nT RX-U5 T RX$-U5 g)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 rr|@N) r#r%r)r+r'maxabsrr,r-) x_ptry_ptrs_ptrrfpidoffsxsyrFs rG$_fp8_blockwise_act_quant_kernel_impl?_lazy_init_triton.._fp8_blockwise_act_quant_kernel_impls mmm#"))Az":: GGEL ! $ $RZZ 0 FF266!9  % E DD'' ( q! a rJcd>TRSS9nTRSS9nTRXE5nXe-TRSU5-n Xu-TRSU5-n U SS2S4U-U SSS24-n U SS2S4U:U SSS24U:-n TRX -U S9R TR 5n TR TRU 55S- nX- nUR URR5nTRX-XS9 TRX&U--U-U5 g)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. rrrNr"rh) r#r$r%r)r+r'rirjrr,r-)rkrlrmr3rrfr4r5nr7r8ror rprqrrrFs rG'_fp8_blockwise_weight_quant_kernel_implB_lazy_init_triton.._fp8_blockwise_weight_quant_kernel_impls/ 1 % 1 % GGA "#bii:&>>#bii:&>>ag"VD!G_4q$w!#tQw!(;< GGELtG , / / ; FF266!9  % E DD'' ( q, "U*A.rJc>TRSS9nTRSS9nTRXE5nXe-TRSU5-n Xu-TRSU5-n U SS2S4U-U SSS24-n U SS2S4U:U SSS24U:-n TRX -U S9R TR 5n TRXU--U-5nX-nTR X+-XS9 g)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 rrrNr")r#r$r%r)r+r'r-)rkrmrlr3rrfr4r5rvr7r8ror rprqrrrFs rG)_fp8_blockwise_weight_dequant_kernel_implD_lazy_init_triton.._fp8_blockwise_weight_dequant_kernel_impls$ 1 % 1 % GGA "#bii:&>>#bii:&>>ag"VD!G_4q$w!#tQw!(;< GGELtG , / / ; GGEAI%- . E q,rJr )_triton_initializedtorch.utils._tritonr_triton_availablerTtriton.languagelanguagerautotunejit constexprr^library custom_opTensorint register_fake_blockwise_fp8_gemm_impl_fp8_blockwise_act_quant_kernel"_fp8_blockwise_weight_quant_kernel$_fp8_blockwise_weight_dequant_kernel)rrblock_mblock_nrfp8_gemm_configsrardrsrwrzrHrFrTs @@@rG_lazy_init_tritonrs?. <<! )  )G$G&J $ >! ' %  )   __ &L ZZ*' <<*' <<*',,*'ll*'ll*'ll*'*'X ]]5BG   << \\ <<\\     H2))* 6 ZZ!)+!!0'K# ZZ//1||//8*Q& ZZ-/1||--:,U(k s(G(r r@rBrArCrMrNc\[5 [(d [S5e[XX#U5$)Nunsupported without triton)rrAssertionErrorr)r@rBrArCrMs rGblockwise_fp8_gemmrs+  9:: #AAJ ??rJrprc^^[5 [(d [S5eSSKmTR 5(dS5eTR S5U-S:Xd SUS35eU[ R[ R4;dS5e[ R"TUS 9nTR"/TR 5SSQTR S5U-P7S [ R06nUU4S jn[U"TX4US 9 X44$) 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`. rrNInput tensor must be contiguousrPz@Last dimension size must be divisible by block_size (block_size=)6dtype must be torch.float8_e4m3fn or torch.float8_e5m2rrcL>TRTR5US54$Nrf)r$rY)metarTrps rGrU)fp8_blockwise_act_quant..sQWWY\0BCErJrf) rrrrTrWrXr^ float8_e4m3fn float8_e5m2 empty_liker]r'r)rprMrrrrqr`rTs` @rGfp8_blockwise_act_quantrs "  9:: ??  ???  66": "a ' J:,VWX '    @@@  %(A RQVVXcr]RAFF2J*$<REMMRA ED#D)!QjI 4KrJc Z^^^[5 [(d [S5eSSKmUR 5(dS5eUR 5S:XdS5eUR S5U-S:XaUR S5U-S:Xd SUS 35eU[R[R4;dS 5eUR 5umm[R"XS 9nURTU-TU-[RS 9nUUU4S jn[U"XUTTUS 9 X44$)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`. rrNrz#Input tensor must have 2 dimensionsrzABoth dimensions of x must be divisible by block_size (block_size=rrrcX>TRTUS5TRTUS54$rrRrr3rrTs rGrU,fp8_blockwise_weight_quant..<. AtL)* AtL)*rJr)rrrrTrWdimrXr^rrrr]r'r) rprMrrrrqr`r3rrTs @@@rGfp8_blockwise_weight_quantrs2   9:: ??  ???  557a<>>>< 66!9z !Q &166!9z+AQ+F KJTRTUS5TRTUS54$rrRrs rGrU.fp8_blockwise_weight_dequant.._rrJr) rrrrTrWrrXr^rget_default_dtyper)rprqrMrrr`r3rrTs @@@rGfp8_blockwise_weight_dequantrDs"  9:: ??  !2!2V4VV 2 557ars_  "&%)"'+$OUn @ || @  @ || @  @  @  \\ @(+ATAT# ||#!$#38;;# 5<< %&#N(+%2E2E' ||'!$' 5<< %&'V9<   ||     25   \\  rJ