o 'i/%@sddlmZmZmZmZddlmmZddl Z ddl m Z ddl m Z ddlmZddlmZddlmmZddlmmZddlZddlZddlmZddlmZmZddlm Z ddl!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*ddl+m,Z,ddl-m.Z.dd l/m0Z0dd l1m2Z2d d lm3Z3m4Z4m5Z5m6Z6dd lm7Z7m8Z8dZ9e(dddde%de"de$de%fddZ:e(dddde%de"de$fddZ;e(dddde#de"de"fddZ Gdd d Z?e j@d!e jAd"e jAfd#d$ZBe j@d"e jAfd%d&ZCd'd(ZDGd)d*d*ZEejFd+eGd,eGd-eGd.eGd/eHd0eHd1eHd2ee jId3ee jId4ee jId5eee jId6eGd7eeGeGfd8eeGeGfd9eGd:eHd;eJde7f$dd?eejAejAfd@eejAejAfdAejAdBejAd2eHd3eHd4eHd6eGdCeejAd9eeGfdDdEZLdS)F)OptionalTupleTypeUnionN)ir)cpasynctcgen05) from_dlpack) Int32Int64Uint8Uint64TInteger dsl_user_opextract_mlir_valuesnew_from_mlir_values)llvm)get_compute_capability)flashinfer_api) WorkTileInfo)get_cutlass_dtypecutlass_to_torch_dtype get_num_smmake_ptr)CallableListlocipobjindexvaluereturncCs>|d|d>M}|||d>O}t|tsJd||S)Nzobj=) isinstancer )r"r#r$r r!r)a/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/flashinfer/cute_dsl/blockscaled_gemm.py with_byteCsr+cCs||d?d@tS)Nr'r&)tor )r"r#r r!r)r)r* read_byteKsr-addrc Cs@ttjt|j||dt|j||dgddddtjjdS)Nrz-atom.add.release.gpu.global.s32 $0, [$1], $2;z=r,l,rTF)has_side_effectsis_align_stack asm_dialect)r r inline_asmri32ir_value AsmDialectAD_ATT)r.r$r r!r)r)r*atomic_add_release_globalPs r7c @s~eZdZddddejdeejdejdeeefdej f dd Z d d Z d d Z e ddddedeeeeffddZdS)MaskedSchedulerParamsNrmasked_m dst_signalscc_tilercluster_shape_mnkc Cs|ddkrtd|dtj||d}|dj} ||_||_||_||_| |_||_ |dd|_ ||_ tj tj |j|dd||d||d|_dS)Nrzunsupported cluster_shape_k )tiler)rNNNr) ValueErrorcute zipped_divideshaper9r:r;r<problem_shape_ntile_mnl_cluster_shape_mnkcluster_shape_mn_loc make_layoutceil_divproblem_layout_ncluster_mnl) selfr9r:r;r<r=r r!gcrEr)r)r*__init__cs&   zMaskedSchedulerParams.__init__cCsNgg}|_|j|j|j|j|jfD]}t|}||7}|jt|q|SN) _values_posr9r:r;r<rFrappendlen)rLvaluesr" obj_valuesr)r)r*__extract_mlir_values__s z-MaskedSchedulerParams.__extract_mlir_values__cCsjg}t|j|j|j|j|jg|jddD]\}}|t||d|||d}qt t |d|j iS)NTstrictr ) zipr9r:r;r<rFrPrQrr8tuplerH)rLrSobj_listr"n_itemsr)r)r*__new_from_mlir_values__s z.MaskedSchedulerParams.__new_from_mlir_values__max_active_clustersr%cCs|}g|j|RSrO)rG)rLr]r r!num_persistent_clustersr)r)r*get_grid_shapesz$MaskedSchedulerParams.get_grid_shape)__name__ __module__ __qualname__rBTensorrPointerrintShaperNrUr\rr rr_r)r)r)r*r8bs.   " r8c@seZdZdedededededejdefdd Zd ee j fd d Z d ee j d dfddZ e eddddedeeeefdeeeeffddZeddddeded eeeeffddZej d)dedeedeedeed eeeeff ddZe   d*ddddeedeedeed eeeeffddZe dddd efd d!Ze d"ddd#d$efd%d&Zed efd'd(ZdS)+MaskedSchedulerparamsr^current_work_linear_idxcurrent_batch_idx accum_tile_mcta_id_in_clusternum_tiles_executedcCs.||_||_||_||_||_||_||_dSrO)rhr^_current_work_linear_idx_current_batch_idx _accum_tile_mrl_num_tiles_executed)rLrhr^rirjrkrlrmr)r)r*rNs  zMaskedScheduler.__init__r%cCs^t|j}|t|j|t|j|t|j|t|j|t|j|SrO)rr^extendrnrorprlrq)rLrSr)r)r*rUs z'MaskedScheduler.__extract_mlir_values__rScCst|dksJt|j|dg}t|j|dg}t|j|dg}t|j|dg}t|j|dd}t|j|dg}t|j ||||||S)Nr'rrr>r) rRrr^rnrorprlrqrgrh)rLrSnew_num_persistent_clustersnew_current_work_linear_idxnew_current_batch_idxnew_accum_tile_mnew_cta_id_in_clusternew_num_tiles_executedr)r)r*r\s4    z(MaskedScheduler.__new_from_mlir_values__Nr block_idxgrid_dimcCs|}tj|||dtj|j||d}|\}}}t|} td} td} t||jdt||jdtdf} td} t||| | | | | S)Nrrr)rBsizerGr rg)rhr{r|r r!r^bidxbidybidzrirjrkrlrmr)r)r*creates,  zMaskedScheduler.creater]cCs|j|||dSNr)r_)rhr]r r!r)r)r*r_szMaskedScheduler.get_grid_shapedsm_pending_packed dsm_counter num_c_stagec Cs|jjd}|j}|j}|t|jj||jjd||krt||jjjdkrtt |duo4|jj durAt ||||dd}|t|jj||jjd7}|t d7}|t|jj||jjd||krt||jjjdks)||_||_|j|jjjdk}|r|jt|jj|j|jjd||k}|||j|||jf} tddt| |jg|jjt dRddD} t| ||fS)Nrr)r#r$css.|]\}}}t|t|t|VqdSrO)r ).0xyzr)r)r* [s  zCMaskedScheduler._get_current_work_for_linear_idx..TrV)rhrErprorBrJr9r<rDcutlass const_exprr:r+r rYrXrlrGr) rLrirrr num_tiles_nrk batch_idxis_validcur_cluster_coordcur_tile_coordr)r)r* _get_current_work_for_linear_idxsv        z0MaskedScheduler._get_current_work_for_linear_idxcCs|j|j|||dS)Nrrr)rrn)rLrrrr r!r)r)r*get_current_workgs  z MaskedScheduler.get_current_workcCs|j||d\}}|Sr)r)rLr r! tile_info_r)r)r*initial_work_tile_infoxsz&MaskedScheduler.initial_work_tile_infor) advance_countr r!rcCs2|jt|t|j7_|jtd7_dSNr)rnr r^rq)rLrr r!r)r)r*advance_to_next_work}s z$MaskedScheduler.advance_to_next_workcCs|jSrO)rqrLr)r)r*rmsz"MaskedScheduler.num_tiles_executedrOr@)r`rarbr8r rBCoordrNlistrValuerUr\r staticmethodrrrr_jitrr r rerrrrrpropertyrmr)r)r)r*rgs    *    I rgc-@s6eZdZdZdedeeefdeeefdefddZdd Ze j d e j d e j d e j d e j de j de j de e j de e j dejdejfddZe jde jde jde jde j de jde j de jde j de jde j d e e jd!e j d"e e j d#e jd$e jd%e jd&e jd'e jd(e jd)ee je jd*fd+e jd,ef,d-d.Zd/e j d0e j d1ee je j e j ffd2d3Zd4ejd5e j d6e j d+e jd7eejefd1ee je j e j ff d8d9Z d:e jd;e j d4ejdZ!d4ejd?ee je jfd6e j d+e jd>> gemm = Sm100BlockScaledPersistentDenseGemmKernel( ... sf_vec_size=16, ... mma_tiler_mn=(256, 128), ... cluster_shape_mn=(2, 1) ... ) >>> gemm(a_tensor, b_tensor, sfa_tensor, sfb_tensor, c_tensor, max_active_clusters, stream) sf_vec_size mma_tiler_mnrG sm_versioncCsddg}||vsJ|d|tj|_||_|ddk|_||_g|dR|_|jr1tjj ntjj |_ d|_ d|_ d|_d |_d t|j|jg|j R|_d|_d|_d |_t||_d }||_d S)a@Initializes the configuration for a Blackwell dense GEMM kernel. This configuration includes several key aspects: 1. MMA Instruction Settings (tcgen05): - acc_dtype: Data types for MMA accumulator, always set to Float32 - sf_vec_size: Scalefactor A/B vector size. - mma_tiler_mn: The (M, N) shape of the MMA instruction tiler. 2. Cluster Shape: - cluster_shape_mn: The (ClusterM, ClusterN) shape of the CTA cluster. :param sf_vec_size: Scalefactor vector size. :type sf_vec_size: int :param mma_tiler_mn: Tuple (M, N) shape of the MMA instruction. :type mma_tiler_mn: Tuple[int, int] :param cluster_shape_mn: Tuple (ClusterM, ClusterN) shape of the cluster. :type cluster_shape_mn: Tuple[int, int] sm_100sm_103zJ are the only supported SM versions for cute-dsl backend, but encountered rr)rrr>rsr r>iN)rFloat32 acc_dtyperuse_2cta_instrsrG mma_tilerrCtaGroupTWOONE cta_group occupancyepilog_warp_id mma_warp_id tma_warp_idrRthreads_per_ctacta_sync_bar_idepilog_sync_bar_idtmem_ptr_sync_bar_idutilsget_smem_capacity_in_bytes smem_capacitynum_tmem_alloc_cols)rLrrrGrsupported_sm_versionsSM100_TMEM_CAPACITY_COLUMNSr)r)r*rNs0    z2Sm100BlockScaledPersistentDenseGemmKernel.__init__cCsd}|jd|jd||jjf|_|jd|jrdndt|jdd|jdf|_t |j|j |j |j |j |j|jdd}t |j|j |j |j |j tjjjj|jdd}d}|jd|jd|jd|f|_|jd|jd|jd|f|_|jdt|jj|jd|jdf|_ttg|jdR|jjf|_ttg|jdR|jjf|_t|jjd|_t|jjd|_t|jjd|_|jdk|_ |jdk|_!|jdk|_"t#|j|j|j$|j%|_&|'||j|j|j |j(|j |j&|j%|j$|j |j |j)|j* \|_+|_,|_-t.||j|j|j,|_/t0||j|j(|j,|_1t23||j|j |j,|_4t25||j|j |j,|_6t7|j%|j$|j&|j-|_8dS)aJSet up configurations that are dependent on GEMM inputs This method configures various attributes based on the input tensor properties (data types, leading dimensions) and kernel settings: - Configuring tiled MMA - Computing MMA/cluster/tile shapes - Computing cluster layout - Computing multicast CTAs for A/B/SFA/SFB - Computing epilogue subtile - Setting up A/B/SFA/SFB/C stage counts in shared memory - Computing A/B/SFA/SFB/C shared memory layout - Computing tensor memory allocation columns rrrr>Nr)9ra_dtypewidthmma_inst_shape_mnkrrBround_upmma_inst_shape_mnk_sfb sm100_utils"make_blockscaled_trivial_tiled_mma a_major_mode b_major_modesf_dtyperrnvgpurrr mma_tiler_sfbr}thr_idrDcta_tile_shape_mnk tiled_dividerIrGcluster_layout_vmnkcluster_layout_sfb_vmnknum_mcast_ctas_anum_mcast_ctas_bnum_mcast_ctas_sfb is_a_mcast is_b_mcast is_sfb_mcastcompute_epilogue_tile_shapec_layoutc_dtypeepi_tile_compute_stagesb_dtyperr num_acc_stage num_ab_stagermake_smem_layout_aa_smem_layout_stagedmake_smem_layout_bb_smem_layout_stagedblockscaled_utilsmake_smem_layout_sfasfa_smem_layout_stagedmake_smem_layout_sfbsfb_smem_layout_stagedmake_smem_layout_epic_smem_layout_staged)rLmma_inst_bits_k tiled_mma tiled_mma_sfbmma_inst_tile_kr)r)r*_setup_attributes:s            z;Sm100BlockScaledPersistentDenseGemmKernel._setup_attributesa_tensorb_tensor sfa_tensor sfb_tensorc_tensormasked_m_tensorr: alpha_tensorr]streamc )s\|j_|j_|j_|j_tj|_ tj|_ tj|_ t jjkr>tdjdjt|jj} t|j| }t|jj} t|j| }tjj j jjjjdd} tjj j jjtjjjjjdd}t | j!j}t"j#| j!}t$j%d}tj&|||j'| j(j\}}t)j#| j!}t$j*d}tj+|||j'| j(j\}}t"j#| j!}t$j,d}tjj&|||j'| j(jt j-d\}}t.j#| j!}t$j/d}tjj+|||j0|j1jt j-d\}}t2j|} t2j|}!t2j|}"t2j|}#| |!|"|#|_3t$j4d}$t56t57||$j8\}%}&9|||j:j#| \_;}'d_?| ||||||||||%|&|j(j1j%j*j,j/j4j8j;j@|'jAd d ggj#d Rj>2| d dS) aExecute the GEMM operation in steps: - Setup static attributes before smem/grid/tma computation - Setup TMA load/store atoms and tensors - Compute grid size with regard to hardware constraints - Define shared storage for kernel - Launch the kernel synchronously :param a_tensor: Input tensor A :type a_tensor: cute.Tensor :param b_tensor: Input tensor B :type b_tensor: cute.Tensor :param sfa_tensor: Scale factor tensor A :type sfa_tensor: cute.Tensor :param sfb_tensor: Scale factor tensor B :type sfb_tensor: cute.Tensor :param c_tensor: Output tensor C :type c_tensor: cute.Tensor :param masked_m_tensor: Masked layout tensor M :type masked_m_tensor: cute.Tensor :param max_active_clusters: Maximum number of active clusters :type max_active_clusters: cutlass.Constexpr :param stream: CUDA stream for asynchronous execution :type stream: cuda.CUstream :param alpha_tensor: Optional 1D tensor of shape (l,) containing per-batch scaling factors. :type alpha_tensor: cute.Tensor :raises TypeError: If input data types are incompatible with the MMA instruction. zType must match: z != Nr>NNNr) internal_typeNNrcsdeZdZUejjejjfe d<ejjejjfe d<ejjejj fe d<ejjejj fe d<eje d<ej e d<ejj ejjj ejjfjfe d<ejj ejjjejjfjfe d<ejj ejjjejjfjfe d <ejj ejjjejfjfe d <ejj ejjjejfjfe d <d S) zISm100BlockScaledPersistentDenseGemmKernel.__call__..SharedStorageab_full_mbar_ptrab_empty_mbar_ptracc_full_mbar_ptracc_empty_mbar_ptrtmem_dealloc_mbar_ptrtmem_holding_bufsCsAsBsSFAsSFBN)r`rarbrBstructMemRangerr r__annotations__rr Alignrcosizerouterbuffer_align_bytesrrrrrrrr)rr)r* SharedStoragesX     rr)gridblockclustersmemr)B element_typerrrrr LayoutEnum from_tensormma_major_moderrrrr TypeErrorrrtile_atom_to_shape_SFrDrrB make_tensoriteratorrrrrrrrrrr}rcluster_shape_to_tma_atom_ArGslice_rmake_tiled_tma_atom_Arrcluster_shape_to_tma_atom_Brmake_tiled_tma_atom_BrInt16cluster_shape_to_tma_atom_SFBrrr size_in_bytesnum_tma_load_bytesrrmake_tiled_tma_atomCopyBulkTensorTileS2GOpr _compute_gridrtile_sched_paramsrrshared_storagekernellaunchr))rLrrrrrrr:rr]r sfa_layout sfb_layoutrr atom_thr_sizea_op a_smem_layout tma_atom_a tma_tensor_ab_op b_smem_layout tma_atom_b tma_tensor_bsfa_opsfa_smem_layout tma_atom_sfatma_tensor_sfasfb_opsfb_smem_layout tma_atom_sfbtma_tensor_sfb a_copy_size b_copy_size sfa_copy_size sfb_copy_sizeepi_smem_layout tma_atom_c tma_tensor_crrr)rr*__call__s**             , z2Sm100BlockScaledPersistentDenseGemmKernel.__call__rrr1mA_mklr5mB_nklr9mSFA_mklr=mSFB_nklrDmC_mnlalpharrrrrrrNrr(c Csbtj}tj|}||jkr)t|t|t|t| t| t|jj dk}tj \}}}|t|jj }|dk}tjtj }| |}| |} tj \}!}"}"t}#|#|j}$|$j}%|$j}&ttjj}'|j|jd}(ttjj|(})tjj|$j|j|'|)|j|d}*ttjj}+t |j!|rdnd},ttjj|,}-tj"j|$j#|j$|+|-|d}.|r||jkrd}/tj%tj&|%|/Wdn1swYtj't|j(dkrtj)|$j*j+|j,|j-d}0|$j.j+|j,|j-d}1|$j/j+|j,|j-d}2|$j0+|}3|$j1+|}4d}5d}6d}7d}8t23|j4p7|j5p7|rZtj6||dd }5tj6||dd }6tj6||dd }7tj6|| dd }8t7|t8|j9d d }9t7|t8|j9d d }:t7|t8|j9d d };t7| t8|j9d d }|:|}?|:|}@|?;|9}A|?<|:}B|?;|;}C|@<|<}D|?=|=}Et>t8|dj }Ft?||d|Ft@|1ddt@|Add\}G}Ht>t8|dj }It?||d|It@|2ddt@|Bdd\}J}K|F}LtjAj?||d|Lt@|3ddt@|Cdd\}M}NtB|M}MtB|N}Nt>t8|dj }OtjAj?| | d|Ot@|4ddt@|Ddd\}P}QtB|P}PtB|Q}Q|C|1}R|D|2}S|E|j9dd}T|FtG|T|j$}Ut|j(dkrtjHn tjjI|jJ|jKd||jkrtL|tj tjM}V|VN}WtOtjPjQ|j}X|WjRr|WjS}Y|Ydt|jj |Yd|Ydf}Z|Hd|Zdd|Zdf}[|Kd|Zdd|Zdf}\|Nd|Zdd|Zdf}]|Qd|Zdd|Zdf}^|XTt2Ud}_|XjV|>kr |*W|X}_t2jXd|>dddD]y}`|*Y|X|_tjZ||[d|XjVf|Gd|Xj[f|*\|X|5dtjZ||\d|XjVf|Jd|Xj[f|*\|X|6dtjZ||]d|XjVf|Md|Xj[f|*\|X|7dtjZ| |^d|XjVf|Pd|Xj[f|*\|X|8d|X]t2Ud}_|XjV|>kr|*W|X}_q|V^|V_\}W}"|WjRs|*`|X||jakrEdt |jag|j!R}atjjI|jb|adtjjc|jdd|&d}bte|b|Ujf}ctjg|bthi|c|jjd}dtkl||j9|jmt8|d}ete|d|e}ftjg|bthi|cthi|f|jjd}gtkn||j9|jmt8|d}hte|g|h}i|o|3|f\}j}k}l|o|4|i\}m}n}otL|tj tjM}V|VN}WtOtjPjp|j}ptOtjPjQ|j$}q|WjRr@|WjS}Y|Ydt|jj |Yd|Ydf}Z|cddd|qj[f}r|pTt2Ud}s|pjV|>kr|r|*q|p}s|r|.Y|q|rthjsjtdt2u|>D]}`|r |*v|p|sdddd|pj[f}t|k|t}u|n|t}vtZ|j|u|ltZ|m|v|otj|Rdgd}wt2jX|wddD]:}xdd|x|pj[f}ydd|xf}z|rthjsjw|f|zjx|rthjsjy|i|zjxtz||r|R|y|S|y|r|rthjsjtdq|*{|p|p]t2Ud}s|pjV|>kr$|r$|*q|p}sq|r.|.||q|q]|V^|V_\}W}"|WjRsK|.`|q||jakr/||j!dkr]tjj}|j~|&|ddt |jag|j!R}atjjI|jb|adtjjc|jdd|&d}bte|b|Ujf}c|!}{||{|c|E||\}|}}}~t|~j |j}|||||{|0\}}}||{| |E||0\} }}tL|tj tjM}V|VN}WtOtjPjp|j$}ttjjdt |j!dt |j!}tjj|j|d}t23|jdur|jdksJd|jj d}|d ksJd!td}td}td}|WjRr|WjS}Y|Ydt|jj |Yd|Ydf}Z|dddg|ZR}|}ddddd|j[f}|.v|t@|dt|}t@|dt|}tj|j dgd}|Vj|}t2X|D]}|ddd|f}tZ||||~||~}t23| dur|| |WjSd}||j}|||||j}tZ|||ddd|ftjjtjjjtjjjd"dt |j!}tjjI|j|d||j!dkrtZ| |d|f|d|f||t23|jdur|dt}t|||k}|rtjj|jddd#n |Yn|YtjjI|j|dt23|jdurl|!d}||j!dkrl|dkrl||krlt|||krlt|jt|dd$|d7}||krlt|||ksNqntj% |.{|Wdn 1swY|]|V^|Vj_|||jd%\}W}|WjRs||j!dkrtjj|ddt |j!}tjjI|j|d||j!dkr|rtj|%|dAtj|%dtjj|b|j~|dt23|jdur)tjjddd#|!d}||j!dkr#|dkr%||kr't|jt|dd$|d7}||ks dSdSdSdS|`dSdS)&zW GPU device kernel performing the Persistent batched GEMM computation. r>rr)barrier_storage num_stagesproducer_groupconsumer_grouptx_countcta_layout_vmnk)rMrNrOrPrRrN)swizzle) mcast_mode)NrNr@)rNNrrs)mode)rrNr)rNrr) barrier_idnumber_of_threads)unroll) tma_bar_ptr mcast_mask) alignmentptr_to_buffer_holding_addr)dtyperFT) unroll_full) is_two_cta)rNrOrzmust be representable in 1 byter'zneed to be packable into a u64)space)read)r$r)rBarchwarp_idxmake_warp_uniformrrprefetch_descriptorr}rrDr{block_idx_in_clusterget_flat_coord thread_idxr SmemAllocatorallocater)rrpipelineCooperativeGroupAgentThreadrrPipelineTmaUmmarrdata_ptrrr$rRrPipelineUmmaAsyncrr elect_one mbarrier_initmbarrier_init_fencerGcluster_arrive_relaxedr get_tensorr innerrrrrrrrrcreate_tma_multicast_mask local_tilerr get_slice partition_A partition_B partition_CrI tma_partition group_modesr filter_zerosmake_fragment_Amake_fragment_Bpartition_shape_Cmake_fragment_CrQ cluster_waitbarrierrrrgr|rmake_pipeline_statePipelineUserTypeProducer is_valid_tiletile_idx reset_countBooleancountproducer_try_acquirerangeproducer_acquirecopyr#producer_get_barrieradvancerr producer_tailrrretrieve_tmem_ptrrrlayout recast_ptrrfind_tmem_tensor_col_offsetrrmake_tmem_layout_sfarmake_tmem_layout_sfbmainloop_s2t_copy_and_partitionConsumerconsumer_try_waitsetField ACCUMULATErange_constexpr consumer_waitSFArSFBgemmconsumer_releaseproducer_commit alloc_tmemrepilog_tmem_copy_and_partition make_fragmentrepilog_smem_copy_and_partitionepilog_gmem_copy_and_partitionPipelineTmaStorerr:r9r r r rankrmretileloadr,store fence_proxy ProxyKind async_shared SharedSpace shared_ctarr-cp_async_bulk_wait_groupr7toint sizeof_i32relinquish_tmem_alloc_permitmbarrier_arrive mbarrier_wait dealloc_tmem)rLrrr1rGr5rHr9rIr=rJrDrKrLrrrrrrrrr(rdrr~rrmma_tile_coord_v is_leader_ctacta_rank_in_clusterblock_in_cluster_coord_vmnkblock_in_cluster_coord_sfb_vmnktidxrrstoragerrab_pipeline_producer_groupnum_tma_producerab_pipeline_consumer_group ab_pipelineacc_pipeline_producer_groupnum_acc_consumer_threadsacc_pipeline_consumer_group acc_pipelinenum_tmem_dealloc_threadsrrrrra_full_mcast_maskb_full_mcast_masksfa_full_mcast_masksfb_full_mcast_maskgA_mklgB_nklgSFA_mklgSFB_nklgC_mnl k_block_cntthr_mma thr_mma_sfbtCgAtCgBtCgSFAtCgSFBtCgC a_cta_layouttAsAtAgA b_cta_layouttBsBtBgBsfa_cta_layouttAsSFAtAgSFAsfb_cta_layouttBsSFBtBgSFBtCrAtCrB acc_shape tCtAcc_fake tile_sched work_tileab_producer_statermma_tile_coord_mnl tAgA_slice tBgB_slice tAgSFA_slice tBgSFB_slicepeek_ab_empty_statusk_blocktmem_ptr_read_threads acc_tmem_ptr tCtAcc_base sfa_tmem_ptr tCtSFA_layouttCtSFA sfb_tmem_ptr tCtSFB_layouttCtSFBtiled_copy_s2t_sfatCsSFA_compact_s2ttCtSFA_compact_s2ttiled_copy_s2t_sfbtCsSFB_compact_s2ttCtSFB_compact_s2tab_consumer_stateacc_producer_statetCtAccpeek_ab_full_statuss2t_stage_coordtCsSFA_compact_s2t_stagedtCsSFB_compact_s2t_staged num_kphases kphase_idx kphase_coordsf_kphase_coordepi_tidxtiled_copy_t2r tTR_tAcc_basetTR_rAcctTR_rCtiled_copy_r2stRS_rCtRS_sCbSG_sCbSG_gC_partitionedacc_consumer_statec_producer_group c_pipeline num_expertsrdsm_pending_idxrbSG_gCtTR_tAcc subtile_cntnum_prev_subtiles subtile_idx tTR_tAcc_mnacc_vecc_bufferepilog_threadswill_write_signalslane_idr)r)r*r*sp                                                             ]                    {                         z0Sm100BlockScaledPersistentDenseGemmKernel.kernelsSFtSFr%c Csjt|}t|}tt|j|j}t||}|d}| |}t ||} | |} || | fS)a Make tiledCopy for smem to tmem load for scale factor tensor, then use it to partition smem memory (source) and tensor memory (destination). :param sSF: The scale factor tensor in smem :type sSF: cute.Tensor :param tSF: The scale factor tensor in tmem :type tSF: cute.Tensor :return: A tuple containing (tiled_copy_s2t, tCsSF_compact_s2t, tCtSF_compact_s2t) where: - tiled_copy_s2t: The tiled copy operation for smem to tmem load for scale factor tensor(s2t) - tCsSF_compact_s2t: The partitioned scale factor tensor in smem - tSF_compact_s2t: The partitioned scale factor tensor in tmem :rtype: Tuple[cute.TiledCopy, cute.Tensor, cute.Tensor] r) rBrmake_copy_atomr Cp4x32x128bOprr make_s2t_copyr{ partition_Sget_s2t_smem_desc_tensor partition_D) rLr/r0 tCsSF_compact tCtSF_compact copy_atom_s2ttiled_copy_s2t thr_copy_s2ttCsSF_compact_s2t_tCsSF_compact_s2ttCtSF_compact_s2tr)r)r*r?s        zISm100BlockScaledPersistentDenseGemmKernel.mainloop_s2t_copy_and_partitionrtAccrrcCst|j|j|j|j||}t|d|}t ||d}| |} | |} t|d|} | | } t | dj|j} || | fS)a Make tiledCopy for tensor memory load, then use it to partition tensor memory (source) and register array (destination). :param tidx: The thread index in epilogue warp groups :type tidx: cutlass.Int32 :param tAcc: The accumulator tensor to be copied and partitioned :type tAcc: cute.Tensor :param gC_mnl: The global tensor C :type gC_mnl: cute.Tensor :param epi_tile: The epilogue tiler :type epi_tile: cute.Tile :param use_2cta_instrs: Whether use_2cta_instrs is enabled :type use_2cta_instrs: bool :return: A tuple containing (tiled_copy_t2r, tTR_tAcc, tTR_rAcc) where: - tiled_copy_t2r: The tiled copy operation for tmem to register copy(t2r) - tTR_tAcc: The partitioned accumulator tensor - tTR_rAcc: The accumulated tensor in register used to hold t2r results :rtype: Tuple[cute.TiledCopy, cute.Tensor, cute.Tensor] )NNrrN)NNrrrr@rrNNN)NNNrrrrr)rget_tmem_load_oprrrrrB flat_dividermake_tmem_copyr{r4r6rrD)rLrr?rrr copy_atom_t2rtAcc_epir thr_copy_t2rr% gC_mnl_epitTR_gCrr)r)r*rjs2      zHSm100BlockScaledPersistentDenseGemmKernel.epilog_tmem_copy_and_partitionrrrc CsJt|j|j|j|}t||}||}||}| |} || |fS)a Make tiledCopy for shared memory store, then use it to partition register array (source) and shared memory (destination). :param tiled_copy_t2r: The tiled copy operation for tmem to register copy(t2r) :type tiled_copy_t2r: cute.TiledCopy :param tTR_rC: The partitioned accumulator tensor :type tTR_rC: cute.Tensor :param tidx: The thread index in epilogue warp groups :type tidx: cutlass.Int32 :param sC: The shared memory tensor to be copied and partitioned :type sC: cute.Tensor :type sepi: cute.Tensor :return: A tuple containing (tiled_copy_r2s, tRS_rC, tRS_sC) where: - tiled_copy_r2s: The tiled copy operation for register to smem copy(r2s) - tRS_rC: The partitioned tensor C (register source) - tRS_sC: The partitioned tensor C (smem destination) :rtype: Tuple[cute.TiledCopy, cute.Tensor, cute.Tensor] ) rget_smem_store_oprrrrBmake_tiled_copy_Dr{r6r) rLrrrr copy_atom_r2sr thr_copy_r2srrr)r)r*rs     zHSm100BlockScaledPersistentDenseGemmKernel.epilog_smem_copy_and_partitionatomc CsVt|d|}|}t|dd}t|dd} t|dtd|| \} } || | fS)aMake tiledCopy for global memory store, then use it to: partition shared memory (source) and global memory (destination) for TMA store version. :param tidx: The thread index in epilogue warp groups :type tidx: cutlass.Int32 :param atom: The copy_atom_c to be used for TMA store version, or tiled_copy_t2r for none TMA store version :type atom: cute.CopyAtom or cute.TiledCopy :param gC_mnl: The global tensor C :type gC_mnl: cute.Tensor :param epi_tile: The epilogue tiler :type epi_tile: cute.Tile :param sC: The shared memory tensor to be copied and partitioned :type sC: cute.Tensor :return: A tuple containing (tma_atom_c, bSG_sC, bSG_gC) where: - tma_atom_c: The TMA copy atom - bSG_sC: The partitioned shared memory tensor C - bSG_gC: The partitioned global tensor C :rtype: Tuple[cute.CopyAtom, cute.Tensor, cute.Tensor] rArr>r)rBrCrrrrI) rLrrNrrrgC_epirDsC_for_tma_partitiongC_for_tma_partitionrr$r)r)r*rs zHSm100BlockScaledPersistentDenseGemmKernel.epilog_gmem_copy_and_partition mma_tiler_mnkrrrrrrrrrc Cs|ddkrdnd} d}t|||d}t|||d}t||| d}t||| d}t|||d}t||t||t| |t| |}d}t||}||}| | |||}|| | ||| ||| |7}| ||fS)aComputes the number of stages for A/B/C operands based on heuristics. :param tiled_mma: The tiled MMA object defining the core computation. :type tiled_mma: cute.TiledMma :param mma_tiler_mnk: The shape (M, N, K) of the MMA tiler. :type mma_tiler_mnk: tuple[int, int, int] :param a_dtype: Data type of operand A. :type a_dtype: type[cutlass.Numeric] :param a_major_mode: Major mode of operand A. :type a_major_mode: tcgen05.OperandMajorMode :param b_dtype: Data type of operand B. :type b_dtype: type[cutlass.Numeric] :param b_major_mode: Major mode of operand B. :type b_major_mode: tcgen05.OperandMajorMode :param epi_tile: The epilogue tile shape. :type epi_tile: cute.Tile :param c_dtype: Data type of operand C (output). :type c_dtype: type[cutlass.Numeric] :param c_layout: Layout enum of operand C. :type c_layout: utils.LayoutEnum :param sf_dtype: Data type of Scale factor. :type sf_dtype: type[cutlass.Numeric] :param sf_vec_size: Scale factor vector size. :type sf_vec_size: int :param smem_capacity: Total available shared memory capacity in bytes. :type smem_capacity: int :param occupancy: Target number of CTAs per SM (occupancy). :type occupancy: int :return: A tuple containing the computed number of stages for: (ACC stages, A/B operand stages, C stages) :rtype: tuple[int, int, int] rrr>r) rrrrrrrrBr#)rrRrrrrrrrrrrrrra_smem_layout_stage_oneb_smem_layout_staged_onesfa_smem_layout_staged_onesfb_smem_layout_staged_onec_smem_layout_staged_oneab_bytes_per_stagembar_helpers_bytesc_bytes_per_stagec_bytesrr)r)r*rsn2        z9Sm100BlockScaledPersistentDenseGemmKernel._compute_stagesr;rc Cs>t|d}g|dR}t|||||}t||} || fS)aUse persistent tile scheduler to compute the grid size for the output tensor C. :param c: The output tensor C :type c: cute.Tensor :param cta_tile_shape_mnk: The shape (M, N, K) of the CTA tile. :type cta_tile_shape_mnk: tuple[int, int, int] :param cluster_shape_mn: Shape of each cluster in M, N dimensions. :type cluster_shape_mn: tuple[int, int] :param max_active_clusters: Maximum number of active clusters. :type max_active_clusters: cutlass.Constexpr :return: A tuple containing: - tile_sched_params: Parameters for the persistent tile scheduler. - grid: Grid shape for kernel launch. :rtype: Tuple[MaskedSchedulerParams, tuple[int, int, int]] rr)rBrr8rgr_) rr:r;rrGr]r<cluster_shape_mnlr(rr)r)r*r'rs   z7Sm100BlockScaledPersistentDenseGemmKernel._compute_gridab_dtypecCsd}|tjtjtjhvrd}|dvrd}|tjtjhvrd}|tjkr)|dkr)d}|tjtjhvr7|dkr7d}|tjtjtjtjtjhvrGd}|S)aO Check if the dtypes and sf_vec_size are valid combinations :param ab_dtype: The data type of the A and B operands :type ab_dtype: Type[cutlass.Numeric] :param sf_dtype: The data type of the scale factor :type sf_dtype: Type[cutlass.Numeric] :param sf_vec_size: The vector size of the scale factor :type sf_vec_size: int :param c_dtype: The data type of the output tensor :type c_dtype: Type[cutlass.Numeric] :return: True if the dtypes and sf_vec_size are valid, False otherwise :rtype: bool TF>r[rrr[)r Float4E2M1FN Float8E5M2 Float8E4M3FN Float8E8M0FNUrFloat16BFloat16)r]rrrrr)r)r*)is_valid_dtypes_and_scale_factor_vec_sizes0zSSm100BlockScaledPersistentDenseGemmKernel.is_valid_dtypes_and_scale_factor_vec_sizea_majorb_majorc_majorcCs&d}|tjur|dkr|dksd}|S)a} Check if the dtypes and sf_vec_size are valid combinations :param ab_dtype: The data type of the A and B operands :type ab_dtype: Type[cutlass.Numeric] :param c_dtype: The data type of the output tensor :type c_dtype: Type[cutlass.Numeric] :param a_major: The major dimension of the A tensor :type a_major: str :param b_major: The major dimension of the B tensor :type b_major: str :param c_major: The major dimension of the C tensor :type c_major: str :return: True if the layouts are valid, False otherwise :rtype: bool TkF)rr^)r]rrerfrgrr)r)r*is_valid_layoutssz:Sm100BlockScaledPersistentDenseGemmKernel.is_valid_layoutscCsd}|ddvr d}|ddvrd}|d|ddkrdnddkr$d}dd }|d|dd ksV|ddksV|ddksV|dd ksV|dd ksV||drV||dsXd}|S) a Check if the mma tiler and cluster shape are valid :param mma_tiler_mn: The (M, N) shape of the MMA instruction tiler :type mma_tiler_mn: Tuple[int, int] :param cluster_shape_mn: The (ClusterM, ClusterN) shape of the CTA cluster :type cluster_shape_mn: Tuple[int, int] :return: True if the mma tiler and cluster shape are valid, False otherwise :rtype: bool Tr)rrFrrr>cSs|dko ||d@dkS)Nrrr))rr)r)r* sz`Sm100BlockScaledPersistentDenseGemmKernel.is_valid_mma_tiler_and_cluster_shape..r[rr))rrGr is_power_of_2r)r)r*$is_valid_mma_tiler_and_cluster_shapes&         zNSm100BlockScaledPersistentDenseGemmKernel.is_valid_mma_tiler_and_cluster_shapemnrhlc CsVd} dd} | ||dk|||fr'| ||dk|||fr'| ||dk|||fs)d} | S)a Check if the tensor alignment is valid :param m: The number of rows in the A tensor :type m: int :param n: The number of columns in the B tensor :type n: int :param k: The number of columns in the A tensor :type k: int :param l: The number of columns in the C tensor :type l: int :param ab_dtype: The data type of the A and B operands :type ab_dtype: Type[cutlass.Numeric] :param c_dtype: The data type of the output tensor :type c_dtype: Type[cutlass.Numeric] :param a_major: The major axis of the A tensor :type a_major: str :param b_major: The major axis of the B tensor :type b_major: str :param c_major: The major axis of the C tensor :type c_major: str :return: True if the problem shape is valid, False otherwise :rtype: bool TcSs*|rdnd}||}d|j}||dkS)Nrrr)r)r^is_mode0_major tensor_shapemajor_mode_idxnum_major_elementsnum_contiguous_elementsr)r)r*check_contigous_16B_alignment> s   zjSm100BlockScaledPersistentDenseGemmKernel.is_valid_tensor_alignment..check_contigous_16B_alignmentrmrnFr)) rmrnrhror]rrerfrgrrur)r)r*is_valid_tensor_alignment s%zCSm100BlockScaledPersistentDenseGemmKernel.is_valid_tensor_alignmentc  Cs`d} t||||s d} t||| | | sd} t||sd} t|||| ||| | | s.d} | S)aD Check if the gemm can be implemented :param ab_dtype: The data type of the A and B operands :type ab_dtype: Type[cutlass.Numeric] :param sf_dtype: The data type of the scale factor tensor :type sf_dtype: Type[cutlass.Numeric] :param sf_vec_size: The vector size :type sf_vec_size: int :param c_dtype: The data type of the output tensor :type c_dtype: Type[cutlass.Numeric] :param mma_tiler_mn: The (M, N) shape of the MMA instruction tiler :type mma_tiler_mn: Tuple[int, int] :param cluster_shape_mn: The (ClusterM, ClusterN) shape of the CTA cluster :type cluster_shape_mn: Tuple[int, int] :param m: The number of rows in the A tensor :type m: int :param n: The number of columns in the B tensor :type n: int :param k: The number of columns in the A tensor :type k: int :param l: The number of columns in the C tensor :type l: int :param a_major: The major axis of the A tensor :type a_major: str :param b_major: The major axis of the B tensor :type b_major: str :param c_major: The major axis of the C tensor :type c_major: str :return: True if the gemm can be implemented, False otherwise :rtype: bool TF)rrdrirlrv)r]rrrrrGrmrnrhrorerfrg can_implementr)r)r*rwL s$1 z7Sm100BlockScaledPersistentDenseGemmKernel.can_implement)1r`rarb__doc__rerstrrNrrBrrcrrdr ConstexprcudaCUstreamrFr*TiledMmaCopyAtomLayoutComposedLayoutrTiler8 TiledCopyrr rboolrrrrrNumericrOperandMajorModerrrr'rdrirlrvrwr)r)r)r*rs)   ?          t +  ? % /        t   " 9    (  4       r sf_ref_tensor sf_mma_tensorcCsNt|dd}t|dd}tt|D]}|j|}||||<qdSzdConvert scale factor tensor from MKL layout to mma specification M(32x4xrest_m)xK(4xrest_k)xL layoutrrsrN)rBrrrr}rget_hier_coord)rri mkl_coordr)r)r*cvt_sf_MKL_to_M32x4xrm_K4xrk_L s  rcCs t|dd}t|dd}dSr)rBr)rr)r)r*'cvt_sf_MKL_to_M32x4xrm_K4xrk_L_mma_spec src CsRdd}|||}|||f}d} d} |||| d| d||| | d| d| f} d} d} tj|tj| tjjtjdd d d }tj| tj| tjjtjddd d }tt|t||j |d d }| ddd d |||| ||||j | }|ddd|ddf}|j |d d }tj||d dd\}}tj|||d d}|||fS)NcS||d|Srr)abr)r)r*rJ z,create_scale_factor_tensor..ceil_divrrrrr)rr>rrsrrrr>rrs)min_valmax_val) permute_order init_type init_configT) non_blockingr>r[)is_dynamic_layout assumed_align)r) cutlass_torchcreate_and_permute_torch_tensortorchfloat32TensorInitTypeRANDOMRandomInitConfigrr r,permute unsqueezeexpandreshapecute_tensor_likeconvert_cute_tensor)romnrhrr^devicerJsf_k ref_shapeatom_matom_k mma_shaperef_permute_ordermma_permute_orderref_f32_torch_tensor_cpucute_f32_torch_tensor_cpucute_f32_torch_tensorref_f32_torch_tensor cute_tensorcute_torch_tensorr)r)r*create_scale_factor_tensor sz         rc @seZdZdedededededededejd ejd ejd ejd ed eeefdeeefdedef ddZe j de j de j de j de j de j de j de e j de j de jfddZdS)MaskedBatchedMatmulCuteDSLrmrnrhrorerfrgr]rr alpha_dtyperrrGsm_countrcCs||_||_||_||_||_||_||_||_| |_| |_ | |_ | |_ | |_ ||_ t|| | | | |||||||| sftd|d| d| d| d| d|d|d|d|d|d|d|d|tj}t||j d|j d||_||_dS)Nz-MaskedBatchedMatmulCuteDSL: Unsupported with z, z, rr)_m_n_k_l_a_major_b_major_c_major _ab_dtype _sf_dtype_c_dtype _alpha_dtype _sf_vec_size _mma_tiler_mn_cluster_shape_mnrrwrrr HardwareInfominget_max_active_clusters_max_active_clusters _sm_version)rLrmrnrhrorerfrgr]rrrrrrGrr hardware_infor)r)r*rN sPP  z#MaskedBatchedMatmulCuteDSL.__init__a_ptrb_ptrsfa_ptrsfb_ptrc_ptr masked_m_ptrdst_signals_ptr alpha_ptrcurrent_streamc  Cstj|tj|j|j|jf|jdkrdnddd} tj|tj|j|j|jf|jdkr,dnddd} tj|tj|j|j|jf|j dkrEdnddd} dd} | |j|j }d }d }|j| |j|d |d | |||d |d |f}|j| |j|d |d | |||d |d |f}d }tj|tj||dd}tj|tj||dd}t |t |tj|tj|jfddd}t |durtj|tj|jfdddnd}t|j |j|j|jd| | ||| ||||j| dS)Nrm)rrr>)rrr>)order)rrncSrrr)rr)r)r*rJf rz5MaskedBatchedMatmulCuteDSL.__call__..ceil_divrrrrr)r)rrrGr)rBrmake_ordered_layoutrrrrrrrrrrrrrrrr)rLrrrrrrrrrrrrrJrrr mma_shape_a mma_shape_brrrrrr)r)r*rFB s     z#MaskedBatchedMatmulCuteDSL.__call__N)r`rarbreryrr^rrNrBrrdrr{r|rFr)r)r)r*r sl        ?  rrmrnrhrorerfrgr]rrrrrrGrrenable_dst_signalscs dtttjdttjf fdd tjtdiddd|dd |d |d |d d ddd| d| d| d|d|gdt R d dtj dtj dtj dtj dtj dtj dttj dttj ffdd }|S)!N input_tensorsr%c s|durddtdD\}}}}}}}}sd}n@|\} } } } } }}}|duks,J| | | | | ||durF|nd|durO|ndf\}}}}}}}}t|tjjdd}t|tjjdd}t|tjjdd}t|tjjdd}t|tjjdd}ttj|tjjdd}|durttj|tjjddnd}|durdurt|tjjddnd}||||||||gS)NcSsg|]}dqS)r[r))rrr)r)r* szWget_cute_dsl_compiled_masked_gemm_kernel..get_cute_pointers..r'r[)r) rrqrrB AddressSpacegmemrr Uint32)r a_data_ptr b_data_ptr sfa_data_ptr sfb_data_ptr c_data_ptrmasked_m_data_ptrdst_signals_data_ptralpha_data_ptr a_tensor_gpu b_tensor_gpusfa_tensor_gpusfb_tensor_gpu c_tensor_gpumasked_m_tensor_gpudst_signals_tensor_gpualpha_tensor_gpurrrrrrrr)r]rrrrr)r*get_cute_pointers s     zCget_cute_dsl_compiled_masked_gemm_kernel..get_cute_pointersrmrnrhrorerfrgr]rrrrrrGrrrrrrrrrrc sT|durtjftdd}t}g||||||||g|R|S)Nr{)r^r)remptyrrr) rrrrrrrrr)rrr*rormrnr)r* tensor_apiW s.  z.tensor_apir)r@) rrrtensorrBrdcompilerrrrc)rmrnrhrorerfrgr]rrrrrrGrrrrr)) r]rrrrr*rormrnrr*(get_cute_dsl_compiled_masked_gemm_kernel s z     (r)r:rlhsrhsoutr9r:c Ks|\} } |\} }|}| j\}}}| j\}}}|dkr|d}| dd}| dd}| dur3t| j} | dd}| d d}t| d ksLJd | t| j\}}|d kr_|d kr_td tdid|d|d|d|dddddddt|dt|dt|d |durdn)t|d|d|d|d| dd||d|du| | | |||||dSd|d|d|d| dd||d|du| | | |||||dS)aL Executes a masked, batched matrix multiplication (GEMM) with scale factors and optional alpha scaling at output. Args: lhs (Tuple[torch.Tensor, torch.Tensor]): Tuple containing the left-hand side input tensor (A) and its scale factor tensor (SFA). - A should be in (m, k, l) order, but physically (l, m, k). For fp4 tensor with 8-bit storage, we expect the shape to be (m, k/2, l). - SFA should be in (m32, m4, rm, k4, rk, l) order, but physically (l, rm, rk, m32, m4, k4) rhs (Tuple[torch.Tensor, torch.Tensor]): Tuple containing the right-hand side input tensor (B) and its scale factor tensor (SFB). - B should be in (n, k, l) order, but physically (l, n, k). For fp4 tensor with 8-bit storage, we expect the shape to be (n, k/2, l). - SFB should be in (n32, n4, rn, k4, rk, l) order, but physically (l, rn, rk, n32, n4, k4) out (torch.Tensor): Output tensor to store the result, with shape (l, m, n). masked_m (torch.Tensor): 1D tensor of shape (l,) specifying the valid row count for each batch (used for masking). ab_dtype (str): Data type for A and B matrices. Supported: "float4_e2m1fn", "float8_e4m3fn", "float8_e5m2". sf_dtype (str): Data type for scale factors. Supported: "float8_e8m0fnu", "float8_e4m3fn". c_dtype (str): Data type for output matrix C. Supported: "float16", "bfloat16", "float32", "float8_e4m3fn", "float8_e5m2". sf_vec_size (int): Vector size for scale factors. Typically 16 or 32. sm_count (int, optional): Number of SMs to use. Default: max available SMs under the CTA configuration. mma_tiler_mn (Tuple[int, int], optional): Shape of the MMA tiler (M, N). Default: (128, 128). cluster_shape_mn (Tuple[int, int], optional): Shape of the CTA cluster (ClusterM, ClusterN). Default: (1, 1). alpha_dtype (str, optional): Data type for alpha scaling factors. alpha (torch.Tensor, optional): Optional 1D tensor of shape (l,) containing per-batch scaling factors. Perform per-batch scaling out = alpha * out. Notes: - Legends of the input tensors: * `l` is the batch size, `m/n` is the number of rows, and `k` is the number of columns. * `m/n32`, `m/n4`, `k4` are constant values 32, 4, 4 respectively. * `m32 * m4 * rm` should be same as `M`, which is `m` padded up to the nearest multiple of 128. * `n32 * n4 * rn` should be same as `N`, which is `n` padded up to the nearest multiple of 128. * `k4 * rk` should be same as `K`, which is `k / sf_vec_size` padded up to the nearest multiple of 4. - The function applies masking per batch using masked_m. - If alpha is provided, each batch output is multiplied by its corresponding alpha value. out = alpha * (A @ B). - The result is written to c_tensor. float4_e2m1fnr>r)rrrG)rrNrLrrzUnsupported kwargs: z,SM110 is not supported for cute-dsl backend.rmrnrhrorerfrgr]rrrrrsm_r)rrrrrrrrr)) rDpoprrrRrrArr)rrrr9r]rrrr:rkwargsa_torch sfa_torchb_torch sfb_torchc_torchrmrhrornrrrGrLrmajorminorr)r)r*grouped_gemm_nt_masked s1               r )Mtypingrrrrcuda.bindings.driverbindingsdriverr{r cutlass.cuterBcutlass.pipelinerl cutlass.torchrr cutlass.utilsrcutlass.utils.blackwell_helpersblackwell_helpersr cutlass.utils.blockscaled_layoutblockscaled_layoutr functools cutlass._mlirrcutlass.cute.nvgpurrcutlass.cute.runtimer cutlass.cutlass_dslr r r r rrrrrcutlass._mlir.dialectsrflashinfer.utilsrflashinfer.api_loggingr.cutlass.utils.static_persistent_tile_schedulerrrrrrrrrr+r-r7r8rgrrrcrrrrcachereryrrrr r)r)r)r*s      ,   $  K]HS   R1       P