o پitG @sJddlmZmZmZmZddlZddlmZddl Z ddl m Z ddl m Z ddlmZddlm mZddlm mZddlmZmZddlmZmZmZmZmZmZmZddeddfddZdd l m!Z!dd l"m#Z#dd l$m%Z%m&Z&m'Z'e j(ddd ded edede)ddf ddZ* ddededdfddZ+ GdddZ,dS))OptionalTupleTypeUnionN)cpasynctcgen05)PointerInt32Float16BFloat16Float32 Float8E4M3FN Float8E5M2lock_ptrreturncCstj|||ddS)zE arrive a spin lock when the lock_ptr is a multimem address. locipN) distributedmultimem_red_relaxed_gpu_add1)rrrr_/home/ubuntu/.local/lib/python3.10/site-packages/flashinfer/cute_dsl/gemm_allreduce_two_shot.pyspin_lock_multimem_arrivesr)nvvm)T) MemOrderKind MemScopeKind AtomicOpKindr expected_val reset_valscopec Cs|dkr4d}||kr2tjttj|jt|j||dt|j||dt j t j ||d }||ks dSdS|dkrfd}||krhtjttj|jt|j||dt|j||dt j t j ||d }||ks>dSdSdS)z| wait on a spin lock until the expected count is reached. Reset flag to reset_val if the expected count is reached. gpurr)b mem_order syncscoperrsysN)r atomicrmwri32rCASllvm_ptrr ir_valuerACQUIRErGPUSYS)rrrr rrresultrrrspin_lock_atom_cas_acquire_wait*s@   r/barrier barrier_mcc Csztj\}}}tj\}} } |||||| } tj|| ||dtjtjjjt || |dd||ddS)z' barrier for inter-gpu sm-wise rrr%)rrr rrN) cutearch block_idxgrid_dimrmultimem_red_release_sys_add1 fence_proxy ProxyKindaliasr/) r0r1 num_ranksrrbidxbidybidzbdimxbdimy_pidrrr"sm_wise_inter_gpu_multimem_barrierUs rBc#@seZdZdZ  d\deejdedee e fdee e fdef d d Z d d Z d dZ e jdddddfde jde jde jdejdejde jde jde jfddZe jde jde jde jde jd e jd!ee jd"e jd#e jd$e jd%e jd&ee je jdfd'e jd(ejdejde jde jde jf"d)d*Zd+ejd,e jd-e jd'e jdeejefd.ee je je jff d/d0Zd1e jd2e jd+ejd3e jd.ee je je jff d4d5Z d+ejd6ee je jfd-e jd'e jd3e jd.ee je je jff d7d8Z!e"de jd9ee e e fd:eejd;eejd'e jde d?e ded.ee e e ffd@dAZ$e"de jdBee e e fdee e fdejd.eejee e e fff dCdDZ%e"de jdEee e e fdFe d.e fdGdHZ&e" d]dIeejdeejd>> gemm = PersistentDenseGemmKernel( ... acc_dtype=cutlass.Float32, ... use_2cta_instrs=True, ... mma_tiler_mn=(128, 128), ... cluster_shape_mn=(2, 2) ... ) >>> gemm(a_tensor, b_tensor, c_tensor, max_active_clusters, stream) nonesm_100 acc_dtypeuse_2cta_instrs mma_tiler_mncluster_shape_mn use_tma_storecCs||_||_||_||_g|dR|_||_|rtjjntjj |_ ||_ d|_ d|_ d|_d|_d|_d|_ |dkrA||_ d|_dt|j|jg|j |jR|_d |_d|_d |_d |_t||_d|_d |_|dkr}tj|_tj|_d Sd S) arInitializes 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. - mma_tiler_mn: The (M, N) shape of the MMA instruction tiler. - use_2cta_instrs: Boolean indicating if the tcgen05 MMA variant with cta_group=2 should be used. 2. Cluster Shape: - cluster_shape_mn: The (ClusterM, ClusterN) shape of the CTA cluster. 3. Output C tensor store mode: - use_tma_store: Boolean indicating whether to use Tensor Memory Access (TMA) for storing results. :param acc_dtype: Data type of the accumulator. :type acc_dtype: type[cutlass.Numeric] :param mma_tiler_mn: Tuple (M, N) shape of the MMA instruction. :type mma_tiler_mn: Tuple[int, int] :param use_2cta_instrs: Boolean, True to use cta_group=2 MMA variant. :type use_2cta_instrs: bool :param cluster_shape_mn: Tuple (ClusterM, ClusterN) shape of the cluster. :type cluster_shape_mn: Tuple[int, int] :param use_tma_store: Use Tensor Memory Access (TMA) or normal store for output C tensor. :type use_tma_store: bool :param all_reduce: All-reduce mode, can be "none", "two_shot" :type all_reduce: str )rrKrrD) rrLrMN) rFrGrIrH mma_tilerrJrCtaGroupTWOONE cta_group all_reduce occupancyepilog_warp_id mma_warp_id tma_warp_idall_reduce_warp_idlenthreads_per_ctacta_sync_bar_idepilog_sync_bar_idtmem_ptr_sync_bar_idall_reduce_sync_bar_idutilsget_smem_capacity_in_bytes smem_capacityr:rank_idtorchrget_world_sizeget_rank)selfrFrGrHrIrJrZ sm_versionrrr__init__sL(   z"PersistentDenseGemmKernel.__init__cCsh|j\}}||jr dnddvrdS|jd|jrdnddkr"dS|jddkr2|jddkr2dSdS)NrLrK@FrrNT)mma_tile_shape_mnrGrI)rmmma_mmma_nrrris_validHs z"PersistentDenseGemmKernel.is_validc Cst|j|j|j|j|j|jdd}tj |j dgd}d}|jd|jd||f|_|jdt |j j |jd|jdf|_ ttg|jdR|j j f|_t |jj d|_t |jj d|_|jdk|_|jdk|_t|jrt|j |j|j|j|_n|j dd|_|||j|j|j|j|j|j|j |j!|j \|_"|_#|_$t%||j|j|j#|_&t'||j|j|j#|_(|jrt)|j|j|j|j$nd|_*|+||j|j"|_,dS)a2Set 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 - Computing epilogue subtile - Setting up A/B/C stage counts in shared memory - Computing A/B/C shared memory layout - Computing tensor memory allocation columns NrLmoderNrrK)- sm100_utilsmake_trivial_tiled_mmaa_dtype a_major_mode b_major_moderFrYrUr2size shape_mnkthr_idshapecta_tile_shape_mnk tiled_divide make_layoutrIcluster_layout_vmnknum_mcast_ctas_anum_mcast_ctas_b is_a_mcast is_b_mcastcutlass const_exprrJcompute_epilogue_tile_shaperGc_layoutc_dtypeepi_tile_compute_stagesb_dtyperhr[ num_acc_stage num_ab_stage num_c_stagemake_smem_layout_aa_smem_layout_stagedmake_smem_layout_bb_smem_layout_stagedmake_smem_layout_epic_smem_layout_staged_compute_num_tmem_alloc_colsnum_tmem_alloc_cols)rm tiled_mmamma_inst_shape_kmma_inst_tile_krrr_setup_attributesRs        z+PersistentDenseGemmKernel._setup_attributescCs|S)Nrxrrrsz"PersistentDenseGemmKernel.Nar"cmax_active_clusters epilogue_opc_mc barrier_flagbarrier_flag_mcc s|j_|j_|j_tj|_tj|_ tj|_ t jjkr:t djdjtjjj jjjdd} t| jj} tj| j} tjd} tjj| || j| jj|jt jur}t j ndd\}}t!j| j}tj"d}tjj#|||j| jj|jt jurt j ndd\}}t$j| }t$j|}||| _%d}d}t j&rtj'd}t()t(*||j+\}},|j-j|\_.}d_/j&rt0j'j1nd tj2Gfd d d }|_34| |||||j&r|n|jjj"j'j+j.|||| j5|j6d d ggjd R|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: Input tensor A :type a: cute.Tensor :param b: Input tensor B :type b: cute.Tensor :param c: Output tensor C :type c: cute.Tensor :param c_mc: Output symmetric tensor C_mc, any write or read to a multicast tensor will be broadcasted to all GPUs :type c_mc: 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 epilogue_op: Optional elementwise lambda function to apply to the output tensor :type epilogue_op: cutlass.Constexpr :raises TypeError: If input data types are incompatible with the MMA instruction. :raises AssertionError: If OOB (Out-Of-Bounds) tiles are present when TMA store is disabled. zType must match: z != NrL)NNNr) internal_typeNNrrcseZdZUejjejjfe d<ejjejjfe d<ejjejj fe d<ejjejj fe d<eje d<ej e d<ejj ejjj fjfe d<ejj ejjjejjfjfe d<ejj ejjjejjfjfe d <d S) z9PersistentDenseGemmKernel.__call__..SharedStorageab_full_mbar_ptrab_empty_mbar_ptracc_full_mbar_ptracc_empty_mbar_ptrtmem_dealloc_mbar_ptrtmem_holding_bufsCsAsBN)__name__ __module__ __qualname__r2structMemRangerInt64r__annotations__rr Alignrbuffer_align_bytesr{cosizerouterrrr c_smem_sizermrr SharedStorage6s<    rrK)gridblockclusterstream)7 element_typer{rrrf LayoutEnum from_tensormma_major_moder|r}rrr TypeErrorrryrzrFrYrUr2r~rrcluster_shape_to_tma_atom_ArIslice_rnvgpumake_tiled_tma_atom_Arr TFloat32cluster_shape_to_tma_atom_Brmake_tiled_tma_atom_B size_in_bytesnum_tma_load_bytesrJrrmake_tiled_tma_atomCopyBulkTensorTileS2GOpr _compute_gridrtile_sched_paramsrrrrshared_storagekernellaunchra)rmrr"rrrrrrrr atom_thr_sizea_op a_smem_layout tma_atom_a tma_tensor_ab_op b_smem_layout tma_atom_b tma_tensor_b a_copy_size b_copy_size tma_atom_c tma_tensor_cepi_smem_layoutrrrrr__call__s&         z"PersistentDenseGemmKernel.__call__rrmA_mklrmB_nklrmC_mnlrrrrrrc CsLtj}tj|}||jkr%t|t|t|j r%t|t |j j dk}tj \}}}|t |j j }|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 r|j-j.| j/| j0dnd}*|j1j.| j/| j0d}+|j2j.| j/| j0d},d}-d}.t|j3p#|j4p#|r6tj5||dd }-tj5||dd }.t6|t7|j8d d }/t6|t7|j8d d }0t6|t7|j8d d }1tj |/dgd}2|9|}3|3:|/}4|3;|0}5|3<|1}6t=t7|dj }7t>||d|7t?|+ddt?|4dd\}8}9t=t7|dj }:t>||d|:t?|,ddt?|5dd\};}<|@|+}=|A|,}>|B|j8dd}?|CtD|?|j'}@t |j+dkrtjEn tjjF|jG|jHd||jkrtjI| tj tjJ}A|AK}BtLtjMjN|j!}C|BjOr|BjP}D|Ddt |j j |Dd|Ddf}E|9d|Edd|Edf}F||S|O|ftgjhjidqr|$l|M|MZtRd}P|MjS|2kr|r|$e|M}PqX|r|(m|N|NZ|A[|A\}B|BjOs|(]|N||j^kr||j$dkrtjjn|jo| |ddt#|j^g|j$R}JtjjF|j_|Jdtjj`|jad| d}Ktb|K|@jc}L|}T|p|T|L|6| |\}U}V}Wd}Xd}Yd}Zd}[d}\d}]d}^d}_t|j rPtq|Wj |jr}X|s|U|X|T|*\}Y}[}\|t|T||6| |*\}}]}^n |t|T|U|6| |*\}Z}X}_tjI| tj tjJ}A|AK}BtLtjMjd|j'}`d}at|j rttjjdt#|j$dt#|j$}btjuj|jv|bd}a|BjOrI|BjP}D|Ddt |j j |Dd|Ddf}Ed}cd}dt|j r|^dddg|ER}cn |_dddddg|ER}d|Vddddd|`jXf}e|(j|`t?|edtw|e}et|j rt?|cdtw|c}cn t?|ddtw|d}dtj |ej dgd}f|Ajx|f}gtU|fD]}h|eddd|hf}itW|U|i|Wt|j r|Yy|Wz}j||j{|jr}j|[||j|g|h|jv}ktW|Y|[|\ddd|kftjj}tjj~jtjjjddt#|j$}ltjjF|j|ld||j$dkrtW||]d|kf|cd|hf|am|aVtjjF|j|ldq |Wz}j||j{|jr}j|X||jtW|Z|X|dddd|hfq tj( |(l|`Wdn 1swY|`Zt|jdkr=t|Ajt |j+tj}m||j$dkr=tjjdddtj(|j|m}ntjt|ntj}tjj~jWdn 1s8wY|A[|A\}B|BjOs||j$dkrXtjj|ddt#|j$}ltjjF|j|ld||j$dkr|rtj||dAtj|dtjj|K|jo|dt|j r|a]t|jdkr||jdkr|j}ot|j}ptj}qtjI| tj tjJ}A|AK}Bd|jj}r|j8d|r}st#|jtjj|s}ttj=|t|sf|sdfd }utj=d|rf|rdfd }vttj|j}wt|w|u|v}x|x9||jdd}y|BjOrW|BjP}Dt|Ajt |j+tj}m|Ddt |j j |Dd|Ddf}E||jdkrhtj(|j|m}ntj|n|pdd!d"Wdn 1scwYtjjF|jdt#|jdt6|t7|j8d d }z|3<|z}{|{d#ddg|ER}||j8dt |j j }}t|}|j}~t|||~|j8df}t7|d#|odff}|y|}|j \}}}t|D]}t|D]x}|d||fj}d$\}}}}t|jrtkrt|\}}}}nKt|jrtkrt|\}}}}n8t|jrtkrt|\}}}}n%t|jrtkr,t|\}}}}nt|jrtkr>t|\}}}}t|||||qАq|A[|A\}B|BjOstjjF|jdt#|jdt |Ajjt |j+}||jdkrtj(t|j||j||jWddS1swYdSdSdSdS)%zW GPU device kernel performing the Persistent batched GEMM computation. rLrrK)barrier_storage num_stagesproducer_groupconsumer_grouptx_countcta_layout_vmnk)rrrrrrTN)swizzle) mcast_mode)NrNNNN)rNNrrMrw)rrNr)rNrr) barrier_idnumber_of_threads)unroll) tma_bar_ptr mcast_mask) alignmentptr_to_buffer_holding_addrFT) unroll_full) is_two_cta)rr)spacetwo_shot)readrr)strider!)rrr NN)rrrr)r2r3warp_idxmake_warp_uniformr^rprefetch_descriptorrrrJr~rrr4block_idx_in_clusterget_flat_coord thread_idxrf SmemAllocatorallocaterrrpipelineCooperativeGroupAgentThreadrrPipelineTmaUmmacreaterdata_ptrrrr`r\PipelineUmmaAsyncrr elect_one mbarrier_initmbarrier_init_fencerIcluster_arrive_relaxedr get_tensorrinnerrrrrcreate_tma_multicast_mask local_tilerrU get_slice partition_A partition_B partition_Cr tma_partition group_modesmake_fragment_Amake_fragment_Bpartition_shape_Cmake_fragment_Cappend cluster_waitr0rbraStaticPersistentTileSchedulerr5initial_work_tile_infomake_pipeline_statePipelineUserTypeProducer is_valid_tiletile_idx reset_countBooleancountproducer_try_acquirerangeproducer_acquirecopyindexproducer_get_barrieradvanceadvance_to_next_workget_current_work producer_tailr]rdretrieve_tmem_ptrrF make_tensorlayoutConsumerconsumer_try_waitsetrField ACCUMULATE consumer_waitgemmconsumer_releaseproducer_commit alloc_tmemrepilog_tmem_copy_and_partition make_fragmentrepilog_smem_copy_and_partitionepilog_gmem_copy_and_partitionPipelineTmaStorerranknum_tiles_executedretileloadtostorer7r8 async_shared SharedSpace shared_ctarcrZr _current_work_linear_idxcp_async_bulk_wait_groupiteratorfence_acq_rel_gpurr9relinquish_tmem_alloc_permitmbarrier_arrive mbarrier_wait dealloc_tmemr_rir:lane_idxrwidth WARP_SIZEmake_copy_atomrCopyUniversalOpmake_tiled_copy_tvrspin_lock_atom_cas_relaxed_waitreint zipped_divide partition_Srange_constexprr multimem_ld_reduce_8xf16r multimem_ld_reduce_4xf32r multimem_ld_reduce_8xbf16r multimem_ld_reduce_16xe4m3rmultimem_ld_reduce_16xe5m2multimem_st_4xb32paramsproblem_layout_ncluster_mnlrB)rmrrrrrrrrrrrrrrrrrrrGr;r<r=mma_tile_coord_v is_leader_ctacta_rank_in_clusterblock_in_cluster_coord_vmnktidxr@smemstoragerrab_pipeline_producer_groupnum_tma_producerab_pipeline_consumer_group ab_pipelineacc_pipeline_producer_groupnum_acc_consumer_threadsacc_pipeline_consumer_group acc_pipelinenum_tmem_dealloc_threadsrrra_full_mcast_maskb_full_mcast_maskgA_mklgB_nklgC_mnl k_tile_cntthr_mmatCgAtCgBtCgC a_cta_layouttAsAtAgA b_cta_layouttBsBtBgBtCrAtCrB acc_shape tCtAcc_fake tile_sched work_tileab_producer_statecur_tile_coordmma_tile_coord_mnl tAgA_slice tBgB_slicepeek_ab_empty_statusk_tiletmem_ptr_read_threadstmem_ptr tCtAcc_baseab_consumer_stateacc_producer_statetCtAccpeek_ab_full_status num_kblocks kblock_idx kblock_coordepi_tidxtiled_copy_t2r tTR_tAcc_basetTR_rAcctTR_rCtiled_copy_r2s simt_atomtRS_rCtRS_sCbSG_sCbSG_gC_partitionedtTR_gC_partitionedacc_consumer_state c_pipelinec_producer_groupbSG_gCtTR_gCtTR_tAcc subtile_cntnum_prev_subtiles subtile_idx tTR_tAcc_mnacc_vecc_bufferepilog_threadstile_idflagrir:lane_idatom_val atom_thr_n atom_thr_m thr_layout val_layoutcopy_atom_loadtiled_copy_fake thr_copy_fakegC_mctCgC_mc tCgC_mc_slicecta_mma_tile_m m_local_ranktCgC_mc_slice_partitionedtCgC_mc_local_rankfrgC_mcatomloop_mloop_nijmc_ptrryzw last_flag_idxrrrrssh                                        F           Z                                            O   $zz PersistentDenseGemmKernel.kernelrztAccrrcCst|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] )rrrN)NNrrrrrrNNNNNNrrrrr)ryget_tmem_load_oprrrrFr2 flat_dividermake_tmem_copyr rl partition_DrNr)rmrzrrrrG copy_atom_t2rtAcc_epir thr_copy_t2rr gC_mnl_epirrrrrrMs2      z8PersistentDenseGemmKernel.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] ) ryget_smem_store_oprrrFr2make_tiled_copy_Dr rrT) rmrrrzr copy_atom_r2sr thr_copy_r2srrrrrrOs     z8PersistentDenseGemmKernel.epilog_smem_copy_and_partitionrcCst|d|}t|jr1|}t|dd}t|dd} t|dtd|| \} } || | fS|} | |} | |}t |dj |j }ttj|j }|||fS)aKMake tiledCopy for global memory store, then use it to: - partition register array (source) and global memory (destination) for none TMA store version; - 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 either: - For TMA store: (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 - For non-TMA store: (simt_atom, tTR_rC, tTR_gC) where: - simt_atom: The SIMT copy atom - tTR_rC: The register tensor C - tTR_gC: The partitioned global tensor C :rtype: Tuple[cute.CopyAtom, cute.Tensor, cute.Tensor] rrrLrKr)r2rrrrJr%rr$rr rrNrrrfrrg)rmrzrrrrgC_epirsC_for_tma_partitiongC_for_tma_partitionrrrrrrrrrrrP-s.#      z8PersistentDenseGemmKernel.epilog_gmem_copy_and_partition mma_tiler_mnkr{rrrrhr[c Csd} | rdnd} t|||d} t|||d} | r"t|||dnd}t|| t|| }d}| r:t||nd}|| }|||||}| r`| |||||||||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 b_dtype: Data type of operand B. :type b_dtype: type[cutlass.Numeric] :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 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 :param use_tma_store: Whether TMA store is enabled. :type use_tma_store: bool :return: A tuple containing the computed number of stages for: (ACC stages, A/B operand stages, C stages) :rtype: tuple[int, int, int] rLrrKNr)ryrrrr2r)rrr{rrrrrhr[rJrra_smem_layout_stage_oneb_smem_layout_staged_onec_smem_layout_staged_oneab_bytes_per_stagembar_helpers_bytesc_bytes_per_stagec_bytesrrrrrms^)       z)PersistentDenseGemmKernel._compute_stagesrc CsTt|d}tj||d}|dj}g|dR}t||}tj||} || fS)a Use 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[utils.PersistentTileSchedulerParams, tuple[int, int, int]] r)tiler)rrrK)r2rrkrrfPersistentTileSchedulerParamsr,get_grid_shape) rrrIrc_shapegc num_ctas_mnlcluster_shape_mnlrrrrrrs  z'PersistentDenseGemmKernel._compute_gridrUrcCs2||dd}|t||}t|}|S)a Compute the number of tensor memory allocation columns. :param tiled_mma: The tiled MMA object defining the core computation. :type tiled_mma: cute.TiledMma :param mma_tiler: The shape (M, N, K) of the MMA tile. :type mma_tiler: tuple[int, int, int] :param num_acc_stage: The stage of the accumulator tensor. :type num_acc_stage: int :return: The number of tensor memory allocation columns. :rtype: int NrL)r(r)r2r*rfget_num_tmem_alloc_cols)rrUrrrrrrrrs z6PersistentDenseGemmKernel._compute_num_tmem_alloc_colsab_dtyperZc Cs(d}|tjtjtjtjtjtjtjhvrd}|tjtjtj hvs<|tjkr/|tjtjtjhvs<|tj kr>|tjtjhvr>d}|tjkrW|tjtjtjtjtjtj tjtjhvsy|tjkrd|tjtjhvsy|tj kr{|tjtjtjtj tjtjhvr{d}t |dko|tjtjtjtjtjhvrd}|S)a Check if the dtypes are valid :param ab_dtype: The data type of the A and B operands :type ab_dtype: Type[cutlass.Numeric] :param acc_dtype: The data type of the accumulator :type acc_dtype: Type[cutlass.Numeric] :param c_dtype: The data type of the output tensor :type c_dtype: Type[cutlass.Numeric] :return: True if the dtypes are valid, False otherwise :rtype: bool TFrD) rr r rUint8Int8r rr r r)rrFrrZrvrrris_valid_dtypessz        z)PersistentDenseGemmKernel.is_valid_dtypescCsd}|s |ddvs|r|ddvsd}|dtdddvr d}|d|r'd nddkr.d}d d }|d|dd ksT|ddksT|ddksT||drT||dsVd}|S) a  Check if the mma tiler and cluster shape are valid :param use_2cta_instrs: Whether to use 2 CTA groups :type use_2cta_instrs: bool :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 Trrp)rrFrKrTirLcSs|dko ||d@dkS)NrrKrrrrrrszPPersistentDenseGemmKernel.is_valid_mma_tiler_and_cluster_shape..r)r7)rGrHrIrv is_power_of_2rrr$is_valid_mma_tiler_and_cluster_shapees,      z>PersistentDenseGemmKernel.is_valid_mma_tiler_and_cluster_shapemnkla_majorb_majorc_majorc Cszd} dd} | ||dk|||fr'| ||dk|||fr'| ||dk|||fs)d} | dkr;|dd kr;|dd kr;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)NrrKrr)rd)dtypeis_mode0_major tensor_shapemajor_mode_idxnum_major_elementsnum_contiguous_elementsrrrcheck_contigous_16B_alignments   zZPersistentDenseGemmKernel.is_valid_tensor_alignment..check_contigous_16B_alignmentr r FrDrrrr) r r rrrrrrrrZrvrrrris_valid_tensor_alignments& z3PersistentDenseGemmKernel.is_valid_tensor_alignmentcCsLd}|d|r dnd|df}|s$||ddkr"||ddks$d}|S)ah Check if the epilogue store option is valid :param use_2cta_instrs: Whether to use 2 CTA groups :type use_2cta_instrs: bool :param use_tma_store: Whether to use TMA store :type use_tma_store: bool :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 mma_tiler_mn: The (M, N) shape of the MMA instruction tiler :type mma_tiler_mn: Tuple[int, int] :return: True if the epilogue store option is valid, False otherwise :rtype: bool TrrLrKFr)rGrJr r rHrvcta_tile_shape_mnrrris_valid_epilog_store_options z6PersistentDenseGemmKernel.is_valid_epilog_store_optionc Cszd}t|||s d}t|||sd}t||| | ||| | | | s$d}t|||||s/d}tdvr;|dkr;d}|S)a 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 acc_dtype: The data type of the accumulator :type acc_dtype: Type[cutlass.Numeric] :param c_dtype: The data type of the output tensor :type c_dtype: Type[cutlass.Numeric] :param use_2cta_instrs: Whether to use 2 CTA groups :type use_2cta_instrs: bool :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 use_tma_store: Whether to use TMA store :type use_tma_store: bool :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)rLrNrRrD)rCrr rrdistrk)rrFrrGrHrIrJr r rrrrrrZ can_implementrrrrs$5 z'PersistentDenseGemmKernel.can_implement)rDrE)rD)-rrr__doc__rrNumericboolrrjrorvrr2jitTensor ConstexprrrTiledMmaCopyAtomrLayoutComposedLayoutrTilerfrr r4 TiledCopyrMrOrP staticmethodrrrrstrrr rrrrrrrrCs <   Z h  9      [  ? % @    b   $   V   *    8  $       rCr)-typingrrrrrjtorch.distributedrrr cutlass.cuter2 cutlass.utilsrfcutlass.pipelinercutlass.utils.blackwell_helpersblackwell_helpersrycutlass.utils.distributedcutlass.cute.nvgpurrcutlass.cute.typingrr r r r r rrcutlass._mlir.dialectsrcutlass.cutlass_dslrcutlass._mlir.dialects.nvvmrrrr"r,r/rBrCrrrrsN    $   + P