o پi!@s ddlZddlmZmZmZmZmZmZddlm Z ddl m m Z ddlZddlZddlmZddlmZmZddlmZddlmZddlmmZddlmmZddl m!Z!m"Z"m#Z#m$Z$ddlm%Z%m&Z&m'Z'm(Z(ddl)m*Z*ddl+m,Z,m-Z-dd l.m/Z/dd l0m1Z1m2Z2dd l3m4Z4dd l5m6Z6m7Z7dd l8m9Z9m:Z:ddl;m Gddde9Z?      d[dee@e@e@e@fdeejAdeejAdeeejAdeejAdeBdeBdeBdeBdee@e@fdee@e@fd eCd!e@d"e@d#eDd$eDf d%d&ZEeFd'krd(eBd)ee@d*ffd+d,ZGejHd-d.ZIeIjJd/eGd0d1d2eIjJd3eGd4d5d2eIjJd6eGd7d8d2eIjJd9ejKejLd:eIjJd;ejKejLd:eIjJdd?d@geBd?dAeIjJdBd?dCgeBd?dAeIjJdDdCd@geBdCdAeIjJdEdCd@geBdCdAeIjJdFeCdGdHd2eIjJdIe@dJdKd2eIjJdLe@dMdNd2eIjJdOdPdQdReIjJdSdPdTdReIMZNeOeNjPdUkreIQdVeOeNjRdWkreIQdXeOeNjSdWkreIQdYeEeNjPeNjTeNjUeNjVeNjWeNjXeNjYeNjZeNj[eNjReNjSeNj\eNj]eNj^eNj_eNj`eadZdSdS)\N)OptionalTypeTupleUnionCallableLiteral)partial)cpasynctcgen05)LdMatrix8x8x16bOpLdMatrix16x16x8bOpStMatrix8x8x16bOpStMatrix16x8x8bOp)Int32Float32Boolean const_expr) LayoutEnum) from_dlpackmake_ptr)PipelineTmaCpAsyncUmma) ParamsBase ArgumentsBase)TileSchedulerOptions)VarlenArguments VarlenManager)GemmSm90NamedBarrierGemmc:@seZdZdZdZejZejZejZ  dde e j de e j de e e fde e e e fd ee d ef d d Zd edefddZej  ddejdejdeejdeejd ededeedejdeejdeejfddZejdejdeejdeejdejdejd ejd!eejd"eejd#eejd$eejd%eejd&eejd'eejd(eejd)ed*ej d+ej!d,eej!d-ej"d.ej"d/ej"d0eej!d1eej!d2e#ej!ej"dfd3e#ej!ej"dfd4ej$d5ed6e j%e&f8d7d8Zejd9e j'j(d:e j'j)d;ee j'j)de*d?e e j'j)ee j'j)ffd@dAZ+ej        ddBe j'j(dCe j'j(dDe j'j)dEe j'j)dejdFejdGejdHejd>e*dIe,dJe*dKeejdLeejdMeej-dNeej-dOeejdPeejdQeejdReejd?e e j'j)e j'j)ejff(dSdTZ.ej ddUej-dVej-dWejdXejdYejdZe d[e,fd\d]Z/d^ejd_ejd?e ej-ejejffd`daZ0dbe*dcejd4ej$dde#e,efd?e ej-ejejff dedfZ1dUej-dgee2dhee e j diejdjejdbe*d?e ej-ejejffdkdlZ3dUej-dme2dhe e j dnejdoe j!dbe*d?e ej-ejejffdpdqZ4ejdejd+ej!drej5dIe,d?e'j(f dsdtZ6d+ej!duej5d?e'j(fdvdwZ7 ddxej!dyej5dzed?e'j(fd{d|Z8ejd}ej5d?e'j(fd~dZ9e:dejde e e e fde e e e fd4ej$de e j de e j dee e j d ee dee e j dee e j dgee2dmee2d ede;dde de d?e e e e ff"ddZe=de e j de e j de e j dee e j de?de?d?efddZ@e=de e j de e j d e de e j d?ef ddZAe=de e e fde e e fded?efddZBe=de de de de de e j de e j de?de?de?d?efddZCe=de e j de e j de e j de e e fde e e fde de de de de?de?de?d?efddZDdS) GemmSm100aThis class implements batched matrix multiplication (C = A x B) with support for various data types and architectural features specific to Blackwell GPUs with persistent tile scheduling and warp specialization. :param acc_dtype: Data type for accumulation during computation :type acc_dtype: type[cutlass.Numeric] :param mma_tiler_mn: Shape of the Matrix Multiply-Accumulate (MMA) tile (M,N) :type mma_tiler_mn: Tuple[int, int] :param cluster_shape_mn: Cluster dimensions (M,N) for parallel processing :type cluster_shape_mn: Tuple[int, int] :note: In current version, A and B tensor must have the same data type - i.e., Float8E4M3FN for A and Float8E5M2 for B is not supported :note: Supported A/B data types: - TFloat32 - Float16/BFloat16 - Int8/Uint8 - Float8E4M3FN/Float8E5M2 :note: Supported accumulator data types: - Float32 (for all floating point A/B data types) - Float16 (only for fp16 and fp8 A/B data types) - Int32 (only for uint8/int8 A/B data types) :note: Supported C data types: - Float32 (for float32 and int32 accumulator data types) - Int32 (for float32 and int32 accumulator data types) - Float16/BFloat16 (for fp16 and fp8 accumulator data types) - Int8/Uint8 (for uint8/int8 accumulator data types) - Float8E4M3FN/Float8E5M2 (for float32 accumulator data types) :note: Constraints: - MMA tiler M must be 64/128 (use_2cta_instrs=False) or 128/256 (use_2cta_instrs=True) - MMA tiler N must be 32-256, step 32 - Cluster shape M must be multiple of 2 if use_2cta_instrs=True - Cluster shape M/N must be positive and power of 2, total cluster size <= 16 Example: >>> gemm = GemmSm100( ... acc_dtype=Float32, ... mma_tiler_mn=(128, 128), ... cluster_shape_mn=(2, 2) ... ) >>> gemm(mA, mB, mD, max_active_clusters, stream) dNF acc_dtypea_dtype mma_tiler_mncluster_shape_mnk sf_vec_sizegather_AcCs||_|ddko|ddv|_||_|ddksJdg|dR|_||_|du|_d|_d|_||_|rB|ddksBJd |jrIt j j nt j j |_ |jsSdnd |_d|_d |_d |_d |_|j|j|_|jd|_t|j|_tjj|jt|j|j|jg|jR|_dS) aInitializes 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_mnk: The (ClusterM, ClusterN) shape of the CTA cluster. :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 cluster_shape_mnk: Tuple (ClusterM, ClusterN) shape of the cluster. :type cluster_shape_mnk: Tuple[int, int] r)zCluster shape K must be 1NTFz'Cluster shape N must be 1 for gather A )rr(r&)r use_2cta_instrsr# mma_tilerr$ blockscaled is_persistentpingpongr%r CtaGroupTWOONE cta_groupnum_ab_load_warps occupancyepilog_warp_id mma_warp_idab_load_warp_idepi_load_warp_idscheduler_warp_idlen num_epi_warpscutearch WARP_SIZEthreads_per_cta)selfr r!r"r#r$r%rCD/home/ubuntu/.local/lib/python3.10/site-packages/quack/gemm_sm100.py__init__sB    zGemmSm100.__init__ epilogue_args varlen_argscCsd}|jd|jd||jjf|_|jd|jrdndt|jdd|jdf|_t|j rKt |j|j |j |j|j|jdd|_d|_n/t |j|j |j |j|j|j|jdd|_t |j|j |j |j|jtjj|jdd|_d}|jd|jd|jd|f|_t|j r|jd|jd|jd|f|_nd|_|jdt|jjj|jd|jdf|_tt|j|jjjf|_ t|j rtt|j|jjjf|_!nd|_!t|j jd|_"|j#r|j"dksJt|j jd|_$|j"dk|_%|j$dk|_&t|j r%t|j!jd|_'|j'dk|_(t j)|j|j|j*dur4|j*nt+j,|j-dur?|j-nt.j/|j0|j1d|_2|j#sOdn |j3durWd nd }|4|j|j|j|j2|j|j5|j|j|j-|j1|j*|j0||t.j67d |j8|j9\|_:|_;|_<|_=d|_>|j#sdn |j3durdn|j;|_?t @|j|j|j|j;|_A|jA|_Bt|j#rtCD|j|j|j|j;|_Bt E|j|j|j5|j;|_Fd|_Gt|j-durt H|j-|j*|j2|j<|_Gd|_It|j1durt H|j1|j0|j2|j=|_It|j r!tJK|j|j|j|j;|_LtJM|j|j|j|j;|_Nnd \|_L|_Nt|j r;|O|j|j|j:|_PdSd }||_PdS)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 r'rr(r&Nr+)layout_c elem_ty_cvarlen_mvarlen_ksm_NN)Qr-r!widthmma_inst_shape_mnkr,r>round_upmma_inst_shape_mnk_sfbrr. sm100_utilsmake_trivial_tiled_mma a_major_mode b_major_moder r4 tiled_mma tiled_mma_sfb"make_blockscaled_trivial_tiled_mmasf_dtyper$r r1r3 mma_tiler_sfbsizethr_idshapecta_tile_shape_mnk tiled_divide make_layoutr#cluster_layout_vmnkcluster_layout_sfb_vmnknum_mcast_ctas_ar%num_mcast_ctas_b is_a_mcast is_b_mcastnum_mcast_ctas_sfb is_sfb_mcastcompute_epilogue_tile_shaped_layoutr ROW_MAJORd_dtypecutlassBFloat16c_layoutc_dtypeepi_tile mCuSeqlensM_compute_stagesb_dtypeutilsget_smem_capacity_in_bytesr?r6 num_acc_stageab_stage epi_stage epi_c_stage sched_stagea_prefetch_stagemake_smem_layout_aa_smem_layout_stageda_smem_load_layout_stagedquack_sm100_utilsmake_smem_layout_cpasync_amake_smem_layout_bb_smem_layout_stagedepi_smem_layout_stagedmake_smem_layout_epiepi_c_smem_layout_stagedblockscaled_utilsmake_smem_layout_sfasfa_smem_layout_stagedmake_smem_layout_sfbsfb_smem_layout_staged_compute_num_tmem_alloc_colsnum_tmem_alloc_cols)rBrFrGmma_inst_bits_kmma_inst_tile_kprefetch_A_idxSM100_TMEM_CAPACITY_COLUMNSrCrCrD_setup_attributess0                           zGemmSm100._setup_attributesmAmBmDmCscheduler_argsstreammSFAmSFBc *stjr| dur | dusJ|j_|j_|dur|jnd_|dur(|jnd_| dur2| jnd_t |_ t |_ |durJt |nd_ |durVt |nd_ t |_t |_tjjkr}tdjdjt|durt}|jdujksJddfdd||fD\}}||tjrt|jj} t| j| } t|jj} t| j| } tjjj} t j!d}t j"d}d \}}tj rt#$j%jj}tj&j'|||j(jj)j|jt*ur t+j,ndd \}}t#-j%jj}tj&j.|||j(jj)j|jt*ur/t+j,ndd \}}d \}}d \}}tjrt#$j%jj}t j/d}tj&j'|| |j(jj)jt+j0d \}}t#1j%jj}t j2d}tj&j.|| |j3j4j5jt+j0d \}}t6j|_7tj rj7t6j|7_7tjrt6j|}t6j|}j7||7_7j7| 9_7d \} }!t|durj8|j9j:t;|d r|j<sd nd d\} }!d \}"}#t|durj8|j=j:dd\}"}#>|t?@|}$jA|jBdud}%C|||||}&|%@|&}'|%D|'|jE}(d_F|dur is_staticassume element_typerP).0strCrD s " z7GemmSm100.__call__....)tuplestriderrCrrDs z$GemmSm100.__call__..c s6g|]}|durt|jtj|j|dndqS)N)r)r> make_tensoriteratorrbr_)rr) new_striderCrD s z&GemmSm100.__call__..NNNrrN) internal_type add_to_outputstoreadd)op_typeload)rKrr&cseZdZUejjejjdfe d<ejjejj dfe d<ejjejj dfe d<ejjejj dfe d<ejjejj dfe d<ejjej fe d<eje d<ee d <ejjejjefd fe d <ejjejjjd ur}jnefjfe d <ejjejjjd urjnefjfe d<e d<ejjejjjejjfjfe d<ejjejjjejjfjfe d<ejjejjfjfe d<ejjejjfjfe d<d S)z)GemmSm100.__call__..SharedStorager&ab_pipeline_array_ptrepi_pipeline_array_ptracc_pipeline_array_ptrsched_pipeline_array_ptra_prefetch_pipeline_array_ptr tile_counttmem_dealloc_mbar_ptrtmem_holding_bufsAIdxNsDsCepisAsBsSFAsSFB)__name__ __module__ __qualname__r>structMemRangeroInt64rz__annotations__r|ryr}r~rAlignrnbuffer_align_bytesrrepi_get_smem_structr!cosizerouterrvrrC)a_idx_smem_sizeepi_c_smem_size epi_smem_sizeepilogue_paramsrBr[ sfa_smem_size sfb_smem_sizerCrD SharedStorageqs^    rr()gridblockclusterrmin_blocks_per_mp)Qrr.rr!rvrnrrr[r from_tensora_layoutb_layoutrlrqmma_major_moderVrW TypeErrorrmAIdxr%rrtile_atom_to_shape_SFr_r$r>rrr]rXr^slice_rrrTcluster_shape_to_tma_atom_Ar#nvgpumake_tiled_tma_atom_Ar-rcrroTFloat32cluster_shape_to_tma_atom_Bmake_tiled_tma_atom_BrInt16cluster_shape_to_tma_atom_SFBrr\rYrd size_in_bytesnum_tma_load_bytes_make_tma_epi_atoms_and_tensorsrrshasattrrrepi_to_underlying_argumentsrto_underlying_argumentsget_scheduler_classrtget_scheduler_argumentsget_grid_shapemax_active_clustersrr Float8E8M0FNUr~r`rshared_storagekernelrlaunchrA)*rBrrrrrFrrGrrr sfa_layout sfb_layout atom_thr_size a_smem_layout b_smem_layout tma_atom_a tma_tensor_aa_opb_op tma_atom_b tma_tensor_b tma_atom_sfatma_tensor_sfa tma_atom_sfbtma_tensor_sfbsfa_opsfa_smem_layoutsfb_opsfb_smem_layout sfa_copy_size sfb_copy_size tma_atom_d tma_tensor_d tma_atom_c tma_tensor_c varlen_paramsTileSchedulerClstile_sched_argstile_sched_paramsrrrC) rrrrrrBr[rrrD__call__s` "                          "/ $zGemmSm100.__call__rXrYrmA_mklrmB_nklrmSFA_mklrmSFB_nklrmD_mnlrmC_mnlrrrcrdra_smem_load_layoutrr r epi_smem_layoutepi_c_smem_layoutrsrrcCst|jdu}t|jdu}|r|rJt|jr|s|sJt| du}t|du} tjtj}!|!|jkrO|||| | | fD] }"t|"durNt |"qAt |j j dk}#tj\}$}%}%|$t |j j }&|&dk}'tjtj}(tj\})}%}%tj}*|*|j}+|+j},|+j}-|#r|!|jkrd}.tj|,|.|j|||+j|'d}/d}0t| r|jt|d|+jd}0|j||+j d}1d}2d}3t|j!dur|j"|j#|+j$| d }2|+j%&|j'f}3d}4t|jr|(|+j)}4|+j*j&|j+|j,d }5|+j*j&|j+|j,d }6|+j-j&|j+|j,d }7d}8t|jr8|r%|j.dn|j.d}9t/|9|j0f}:|+j1&|:}8d \};}t| rp|+j6j&|j+|j,d }>|7||+}?|8|&}@t|j2r|8|&nd}A|9|j:dd}B|;t<|B|j=}Ct>j?|||j@tA|s|jBdur|j dn|jBj dtA|j d d }DtC|j?||3|2}tDjEtFtGjHtjjItJ|jKg|jLRd}Ed}Ft| rtDjEtFtGjMdtjjId}F|!|jkrd}G|DN|||G|DO}H|DP}I|Q|(}Jd}Kt|j2r|Q|(}Kd \}L}Md \}N}Ot|jRp.|jSp.|#rWt jT||Jdd}Lt jT||Jd d}Mt|j2rWt jT||Jdd}Nt jT||Kd d}O|}P|PU}QtDVtDjWjX|jY}Rt|rp|DZt[d}S|Qj\r|Qj]}T|Td}U|D^|U|j_|j`|G|Tdt |j j |Td |Tdf}Vd}Wt|j r|Da||U}Xtb|Xtc|j:ddg|Vddf}Wtb|Dd||Utc|j:d dg|Vd df}Yt|j2rtb|Da||Utc|j:ddg|Vddf}Ztb|Dd| |Utc|j:d dg|Vd df}[|De|G|Df|U}\t/t|dj }]d}^t|j r7|@g|W}_thji||Jd|]|_|6|L|Hd\}^}%}%|@j|Y}`t|j2rL|@g|Z}a|Aj|[}bthji||Jd t/t|dj |`|7|M|Id\}c}%}%d \}d}et|j2rthji||Jd|]|a|;d|Nd\}d}%}%t/t|dj }fthji| |Kd |f|b||T\}}%}th||0}|r|Ft[d}ttj |d gd}tj|d dD]}|0|~|||~d|0|~|~}q|Pn|Po}Q|Qj\s|0p|~|!|jKkr |Etjj|jd |-d!}||5}||7}t||Cj}t|j2rtj|t||jd"}t||j:|jt|d#}t||}tj|t|t||jd"}t||j:|jt|d#}t||}||;|\}}}||<|\}}}nd \}}d\}}}d\}}}|}P|PU}QtDVtDjWjr|jY}tDVtDjWjX|j=}|Qj\r|Qj]}T|Td}U|Df|U}tk||j:d}g|ddd|jyf}||/|1|||||||g|'|(||||||||\}}}|Pn|Po}Q|Qj\s|1p||!|jKk r|!|jLdkr%tjj|j|-|#d$|Et[|!|jLdk}G|D| |||G|D}|D}tjj|jd |-d!}t||Cj}tDjEtFtGj|jtjjId}|)}|||||#\}}}t|j |j}|||j|j||=|\}}}d\}}}d}t|dur|||j|j|>|j|\}}}}|}P|PU}QtDVtDjWjr|j=}|}tDVtDjWjr|j}t|r|DZ|Qj\ r|Qj]}T|Td}U|||j|U\}}|D|U|j|||G|ddddd|jyf}|1w||D|Gd}t| r |j| |D| |U|j.dd||=|T|d%\}}%}%d}t|dt|}|Df|U}tC|j||||| o>|dkd&}|||?||0||d|||||||||||||T|D||P||G\}}%tj{ |1||Wdn 1 svwY|}|Pn|Po}Q|Qj\s|!|jLdk rtjj|#d$||!|jLdk rt|# rtj|,|(d Atj|,dtjj||j|#d$|G r|pdSdSdS)'zW GPU device kernel performing the Persistent batched GEMM computation. Nr&r )rXrcab_pipeline_mbar_ptr is_leader_cta)NNr) c_smem_layoutepi_pipeline_mbar_ptr)rcacc_pipeline_mbar_ptr)sched_pipeline_mbar_ptrhas_C)swizzlerNr() len_m_static len_k_static) barrier_id num_threadsT) mcast_moder*)rrNr) cta_coord cta_layout src_tensor dst_tensor mcast_mask tma_desc_ptr)rNrr)r0r1r2r3 filter_zerosr4F)limit_mlimit_kNNN)r.is_async)is_scheduler_warp unroll_fullpredunrollmode)src_idxproducer_stater) alignmentptr_to_buffer_holding_addrdtyper) is_two_cta)r5) clear_acc)r cu_seqlens_m cu_seqlens_kr%r>r?make_warp_uniformwarp_idxr9r prefetch_descriptorr]r^r_ block_idxblock_idx_in_cluster thread_idxrorw SmemAllocatorallocaterrr mbarrier_initmake_ab_pipelinerdata_ptrmake_epi_pipelinerrmake_acc_pipelinertile_count_semaphoremake_sched_pipeliner#rr get_tensorr}make_a_prefetch_pipelinerrrinnerrr`rbr~rr.rrrrepi_get_smem_tensors get_slicepartition_shape_Cr-make_fragment_Cappendryrcreatenum_epi_tensormapsrrrpipeline NamedBarrierintrTmemPtrr@r<r8r7 EpilogueLoadinit_tensormap_ABget_tma_desc_a_ptrget_tma_desc_b_ptrget_flat_coordrgrhcreate_tma_multicast_maskinitial_work_tile_infomake_pipeline_statePipelineUserTypeProducerrzfence_tensormap_initr is_valid_tiletile_idxupdate_tensormap_ABrroffset_batch_A local_tileselectoffset_batch_Bfence_tensormap_update_ABlen_k partition_A copy_utilstma_get_copy_fn partition_Bceil_divload_ABarriveadvance_to_next_workget_current_work producer_tailr5Consumeroffset_batch_AIdxlen_m_make_gmem_tiled_copy_Ar consumer_waitgather_m_get_copy_fnindex sync_warp elect_oneconsumer_releaseadvancegather_k_get_copy_fnload_A_gather_Ar; tiled_copy_1dlane_idx partition_D partition_Smake_identity_tensor make_fragmentrangeproducer_acquirecopyproducer_commit flat_dividefetch_next_workr:r|epilog_gmem_copy_and_partitionoffset_batch_epitma_producer_copy_fnarrive_and_waitretrieve_tmem_ptrr make_fragment_Amake_fragment_Brlayout recast_ptrr find_tmem_tensor_col_offsetr[rmake_tmem_layout_sfar$make_tmem_layout_sfbmainloop_s2t_copy_and_partitionmma alloc_tmemrinit_tensormap_epiepi_get_tma_atomsget_tma_desc_d_ptrget_tma_desc_epi_ptrsEpiloguer=epilog_tmem_copy_and_partitionepilog_smem_store_and_partitionrlrnepilog_smem_load_and_partitionrqrrmake_epi_store_pipeline&epi_get_tensormap_update_shapes_ordersupdate_tensormap_epifence_tensormap_update_epi group_modesrankepi_load_acc_subtileepiloguerelinquish_tmem_alloc_permitmbarrier_arrive mbarrier_wait dealloc_tmem)rBrXrYrrrrrrrrrrrrrrrcrdrrrr r r r!rsrrrKrLhas_Dr)rOtma_atomr,bidx_mma_tile_coord_vr$cta_rank_in_clustertidxsmemstoragerrnum_tmem_dealloc_threads ab_pipeline epi_pipeline acc_pipelinesched_pipelinera_prefetch_pipelinesA_mmarrra_idx_smem_dima_idx_smem_layoutrrrrepi_smem_tensorsthr_mma thr_mma_sfb acc_shape tCtAcc_fakevarlen_managertmem_alloc_barrierepi_load_barrier is_tma_warptma_desc_a_ptrtma_desc_b_ptrblock_in_cluster_coord_vmnkblock_in_cluster_coord_sfb_vmnk a_mcast_mask b_mcast_masksfa_mcast_masksfb_mcast_masktile_scheduler work_tileab_producer_statedo_epi_load_barrier_arrivetile_coord_mnkl batch_idxmma_tile_coord_mnlgA_mkmA_mkgB_nkgSFA_mklgSFB_nklr~ a_cta_layoutcopy_AtCgAtCgBtCgSFAtCgSFBcopy_Bcopy_SFAcopy_SFBsfb_cta_layout k_tile_cnta_prefetch_consumer_statemAIdx_mkr tiled_copy_A thr_copy_A prefetch_Ar;tile_Mtile_K thr_copy_AIdxtAsAIdxtAcAIdxtiled_copy_AIdxa_prefetch_producer_stategAIdxtAgAIdxm_limit tApAIdx_mmk_tilek_limit tApAIdx_kepi_producer_statedo_epi_load_barrier_wait copy_C_fnbGS_gCcopy_C epi_tile_numepi_idx acc_tmem_ptrtCrAtCrB 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_statek_lentCtAcctma_desc_d_ptrtma_desc_epi_ptrsepilogue_barrierepi_tidxtiled_copy_t2r tTR_tAcc_basetTR_rAcctTR_rDtiled_copy_r2stRS_rDtRS_sDtRS_rCtSR_rCtSR_sCtiled_copy_s2racc_consumer_stateepi_store_pipelineepi_read_state epi_shapes epi_orderstTR_tAcccopy_Dload_acc_subtilerCrCrDrs$                                                                             B                               :                            #                     V  zGemmSm100.kernel a_pipelinea_producer_staterrrrreturnc Cs2td}d|kr ||}d}tj|dddD]@} |j} d} t|dur0|| | |f} ||||||| | g| R|||td}| d|krX||}qd|kr|d} |j} d} t|dury|| | |ddf} ||||||| | g| Rddi|||||fS) NTrFr(r@rCr>r?) rproducer_try_acquirerorrrrrproducer_cpasync_commit) rBrErFrrrrpeek_a_empty_statusrrsmem_idx prefetch_outrCrCrDrs<        zGemmSm100.load_A_gather_Arrr*r+rraccr$rr r#r$r'r%r(r&r)c Cst|du}t|r>tdd| | fDsJtdd||fDs$Jtdd||fDs1Jtdd||fDs>Jt|joD|j}td}d| krW| sR|rW||}| r^|||tj j dt j |d gd }t j| d d D]}t|r| s|||t j|jj|j| d @dWdn1swY| r|||t|rdddd|jf}||}||}t |||t |||t j|ddD]=}dd||jf}t|rdd|f}|tj j| |j|tj j| |jt ||||||||tj j dq|||td}|d | kr2| s-|r2||}qu| r;||||||fS)Ncs|]}|duVqdSNrCrxrCrCrDrYz GemmSm100.mma..csrNrOrCrPrCrCrDrZrRcsrNrOrCrPrCrCrDr[rRcsrNrOrCrPrCrCrDr\rRTrFr&rBr(r@)dst_rankr<)rallr%r,rconsumer_try_waitrsetr Field ACCUMULATEr>r]rorrr?rsync_object_fullarrive_mbarrierrrSFArSFBgemmrrr)rBrrr*r+rXrrrMrr$rr r#r$r'r%r(r&r)r.need_nonleader_ctapeek_ab_full_status num_k_blocksrs2t_stage_coordtCsSFA_compact_s2t_stagedtCsSFB_compact_s2t_staged k_blk_idx k_blk_coordsf_kblock_coordrCrCrDr@sb            z GemmSm100.mmar2r6rBr4r7rrKc CsH|st||ddd|f|||}||dS|ddS)Ng)r>rretilerrfill) rBr2r6rBr4r7rrKtRS_rAccrCrCrDrs   zGemmSm100.epi_load_acc_subtilesSFtSFc 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) r>r6make_copy_atomr Cp4x32x128bOpr4r[ make_s2t_copyrarget_s2t_smem_desc_tensorr) rBrkrl tCsSF_compact tCtSF_compact copy_atom_s2ttiled_copy_s2t thr_copy_s2ttCsSF_compact_s2t_tCsSF_compact_s2ttCtSF_compact_s2trCrCrDrs        z)GemmSm100.mainloop_s2t_copy_and_partitionrtAccr,cCst|j|jdur |jntj|jdur|jntj|j ||}t |d|}t ||d}||}||} t |jd|jdf} t | |} || } t | dj|j } || | fS)aw 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: Int32 :param tAcc: The accumulator tensor to be copied and partitioned :type tAcc: 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] N)rNrrN)NNrrrrr()NNNrr)rTget_tmem_load_opr`rlrrmrnrorpr r>rr make_tmem_copyrarrrrr_)rBrryrsr, copy_atom_t2rtAcc_epir2 thr_copy_t2rrBcAcccAcc_epitTR_cAccr4rCrCrDrs"      z(GemmSm100.epilog_tmem_copy_and_partitionrlrIr5rc Csnt|dur|ntj|dur|ntj|j|}t||}| |} |dur+| |nd} | |} || | 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_rD: The partitioned accumulator tensor :type tTR_rD: cute.Tensor :param tidx: The thread index in epilogue warp groups :type tidx: Int32 :param sD: The shared memory tensor to be copied and partitioned :type sD: cute.Tensor :type sepi: cute.Tensor :return: A tuple containing (tiled_copy_r2s, tRS_rD, tRS_sD) where: - tiled_copy_r2s: The tiled copy operation for register to smem copy(r2s) - tRS_rD: The partitioned tensor C (register source) - tRS_sD: The partitioned tensor C (smem destination) :rtype: Tuple[cute.TiledCopy, cute.Tensor, cute.Tensor] N) rTget_smem_store_oprrmrorpr r>make_tiled_copy_Drarrh) rBr2rlrIr5rr copy_atom_r2sr6 thr_copy_r2sr8r7rCrCrDrs    z)GemmSm100.epilog_smem_store_and_partitionrqr tRS_rD_layoutcCst|||j|}|j}t|trt|j|jd} nt|t r-|jdvr-t |jdd} nt j } t | |} t | |} | |} | |} t ||}| |}| ||| fS)N) num_matrices transpose)r&r+r&)r)rTrr op isinstancer r rrrr r>rCopyUniversalOprmrrarrrh)rBr2rqrIrrrrstore_opr copy_atom_s2rr< thr_copy_s2rr;r9r:rCrCrDr+s           z(GemmSm100.epilog_smem_load_and_partitionr#c Cst|j r d}n|jddt|j rdnd}ttjj|}|j|j d}|}ttjj|} t|j rJtj j ||j || |j |d} | Stj ||j || |j ||jsYdn|s]dndd} | S)Nr(r"r&)barrier_storage num_stagesproducer_groupconsumer_grouptx_countcta_layout_vmnkr)rrrrrrproducer_drop_count)rr%r5r,rgCooperativeGroupAgentThreadrerfPipelineTmaUmmarerzrr) rBrXrcr#r$ producer_cntab_pipeline_producer_group mcast_sizeconsumer_arrive_cntab_pipeline_consumer_group pipeline_abrCrCrDrWLsB     zGemmSm100.make_ab_pipeliner'cCsJttjj}|j|jrdnd}ttjj|}tjj||j|||dS)Nr&r()rrrrr) rgrrrr=r,PipelineUmmaAsyncrery)rBrcr'acc_pipeline_producer_groupnum_acc_consumer_threadsacc_pipeline_consumer_grouprCrCrDrZszGemmSm100.make_acc_pipelinecluster_layout_mnkr(r)c Csttjj}t|}|jt|jg|j |j R}|r"|d7}||d}ttjj|}tj j ||j ||t|dkrBddSddS)Nr(r)rrrr consumer_mask)rgrrrr>r]r5r<r8r7r; PipelineAsyncrer}r) rBrr(r)sched_pipeline_producer_group cluster_size warps_per_ctarsched_pipeline_consumer_grouprCrCrDr\s(  zGemmSm100.make_sched_pipelinea_prefetch_pipeline_mbar_ptrcCsHd}tjtjj||d}|jd}ttjj|}tjj||j||dS)Nr")rFr()rrrr)rgrrrr5PipelineCpAsyncrer~)rBrra_prefetch_producer_groupra_prefetch_consumer_grouprCrCrDr^s  z"GemmSm100.make_a_prefetch_pipeline mma_tiler_mnkr`rvr[rnrrr)NrKrL smem_capacityr6c#Cs|du}t| r d}n |ddkrdnd}t|ddkr!dnd}| dur)dn t|ddkr4dnd}t|||d}t|||d}| durRt| | |dnd}| dur`t| | |dnd}t|rvt|||d}t |||d}t ||t ||}t|dkr|t j d |d7}t|r|t ||t ||7}d }t|d kr|t j d |dd7}| durt | |nd}|| | ||}||}t| durt | |} || |7}||||}!|!|}"||!||"|7}||"||fS) a~Computes 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 d_dtype: Data type of operand C (output). :type d_dtype: type[cutlass.Numeric] :param d_layout: Layout enum of operand D. :type d_layout: 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 :return: A tuple containing the computed number of stages for: (ACC stages, A/B operand stages, C stages) :rtype: tuple[int, int, int] Nr&r(r'rr+rrLrrK)rr>r]rTrrrrrrrrrPepi_smem_bytes_per_stage)#clsrXrr`rsr!rvr[r$rnrrrlrqrFrrr6r.ryr{r|a_smem_layout_staged_oneb_smem_layout_staged_oned_smem_layout_staged_onec_smem_layout_staged_onesfa_smem_layout_staged_onesfb_smem_layout_staged_oneab_bytes_per_stagembar_helpers_bytesd_bytes_per_stageepi_bytes_per_stage epi_bytesc_bytes_per_stageremaining_bytesrzrCrCrDrus- &        zGemmSm100._compute_stagesr-rycCs4||dd}|t||}tj|}|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 Nr&)rbrcr>rdrorwget_num_tmem_alloc_cols)rXr-ryrrrrCrCrDr?s z&GemmSm100._compute_num_tmem_alloc_colsa_majorb_majorc Csd}||krd}|}|tjtjtjtjtjtjtjhvrd}|ttjt hvsA|tjkr5|tjtjtjhvsA|t krC|tjtjhvrCd}|dur~|tkr]|ttjtjtjtjt tjtjhvs||tjkrj|tjtjhvs||t kr~|tjtjtt tjtjhvr~d}|tj ur|dkr|dksd}|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 d_dtype: The data type of the output tensor :type d_dtype: Type[cutlass.Numeric] :return: True if the dtypes are valid, False otherwise :rtype: bool TFNk) roFloat16rprUint8Int8 Float8E4M3FN Float8E5M2rr Float4E2M1FN)r!rvr rnrris_validab_dtyperCrCrDis_valid_dtypesWsj    zGemmSm100.is_valid_dtypesrcCsd}|tjtjtjhvrd}|dvrd}|tjtjhvrd}|tjkr)|dkr)d}|tjtjhvr7|dkr7d}|ttjtjtjtjhvrFd}|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 d_dtype: The data type of the output tensor :type d_dtype: Type[cutlass.Numeric] :return: True if the dtypes and sf_vec_size are valid, False otherwise :rtype: bool TF>rr"r"r)rorrrrrrrp)rr[r$rnrrCrCrD)is_valid_dtypes_and_scale_factor_vec_sizes(z3GemmSm100.is_valid_dtypes_and_scale_factor_vec_sizecluster_shape_mnr.cCsd}|ddvr d}|s|dtdddvrd}n|ddvr!d}d d }|d|dd ksG|ddksG|ddksG||drG||dsId}|rY|dd ksW|dd krYd}|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)@rHr'Fr(r"i)rHr'cSs|dko ||d@dkS)Nrr(rC)rQrCrCrDrsz@GemmSm100.is_valid_mma_tiler_and_cluster_shape..rr+)r)r"rr.r is_power_of_2rCrCrD$is_valid_mma_tiler_and_cluster_shapes,      z.GemmSm100.is_valid_mma_tiler_and_cluster_shapernrld_majorc 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 d_dtype: The data type of the output tensor :type d_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 d_major: The major axis of the C tensor :type d_major: str :return: True if the problem shape is valid, False otherwise :rtype: bool TcSs*|rdnd}||}d|j}||dkS)Nrr(rH)rP)rIis_mode0_major tensor_shapemajor_mode_idxnum_major_elementsnum_contiguous_elementsrCrCrDcheck_contigous_16B_alignment/ s   zJGemmSm100.is_valid_tensor_alignment..check_contigous_16B_alignmentrrFrC) rrrrrrnrrrrrrCrCrDis_valid_tensor_alignment s%z#GemmSm100.is_valid_tensor_alignmentc CsRd} t||||| | sd} tj||ddsd} t||||||| | | s'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 d_dtype: The data type of the output tensor :type d_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 d_major: The major axis of the C tensor :type d_major: str :return: True if the gemm can be implemented, False otherwise :rtype: bool TF)r.)rrrr) rr rnr"rrrrrrrr can_implementrCrCrDr= s.zGemmSm100.can_implement)NFrN)NNNNNNNN)F)Errr__doc__r?rrfEpilogueArgumentsEpilogueParamsrroNumericrrirboolrErrr>jitTensorrrcudaCUstreamrrTiledMmaCopyAtomrrParamsLayoutComposedLayoutrTile Constexprrrgr PipelineStaterrr TiledCopyrrrrrrrPointerrWrZr\r^ classmethodrru staticmethodrstrrrrrrrCrCrCrDrhs<.   CG    &      L  -      V   #  4  * ! 2                 N 5   +  4       rr'r'r&r(皙?r(Fmnklrrnrrr rrrc_majorr"r tolerancewarmup_iterations iterationsskip_ref_checkdynamic_persistentc6 sDtdtd|td|d|d|td|d|d|td | d | td | td | td | td||rQJd|\}}}t|||| | |||||| std|d|d|d| d| d|d|dd|d|d|d|tjstdtdd<fdd }||||dk|dd\}}||||dk|dd\}}|||||dk|dd\} }}|dur|||||dk|\} }}nd\} }g| dR}t||| |}t j } | | d| d}!|rtj dtjdd}"nd}"tt|!|"dur.tt|"tjjd d!ndd" |t tj}#t|#jt|   |s  |t jt jt jt jhvrxt d#|!|!}$nt d#||}$|dur|$|}$|$!}$|!}%|t"kr|$}&nY|t jt jhvr|dkrd$nd%}'|dkr|||fn|||f}(t#j$|(tj%|'t#j&j'd&})t(|)d'd!j)|dkrdndd(}*||*_*t#j+|$|*|dd}*|)!}&n|$,t#-|}&tj.j/|%|&| d)d*dd+l0m1}+ttjjd,|||},| }-| }.ddl2}/|/3d-|j4d.kr@|dks.Jtj5d/tj6dd  fd0d1}0n|durMfd2d1}0n|,j-fd3d1}0|+|0|.|-d4}1|,|1d5}2td6|1d7d8|2d9|/3d- f d:d1}3|+|3|.|-d4}4|,|4d5}5td;|4d7d8|5d9dS)=a"Execute a persistent batched dense GEMM operation on Blackwell architecture with performance benchmarking. This function prepares input tensors, configures and launches the persistent GEMM kernel, optionally performs reference validation, and benchmarks the execution performance. :param mnkl: Problem size (M, N, K, L) :type mnkl: Tuple[int, int, int, int] :param ab_dtype: Data type for input tensors A and B :type ab_dtype: Type[cutlass.Numeric] :param d_dtype: Data type for output tensor C :type d_dtype: Type[cutlass.Numeric] :param acc_dtype: Data type for accumulation during matrix multiplication :type acc_dtype: Type[cutlass.Numeric] :param a_major/b_major/d_major: Memory layout of tensor A/B/C :type a_major/b_major/d_major: str :param mma_tiler_mn: MMA tiling size. If not specified in the decorator parameters, the autotuner will use the default value of (256, 256). Otherwise, the autotuner will use the value specified in the decorator parameters. :type mma_tiler_mn: Tuple[int, int], optional :param cluster_shape_mn: Cluster shape. If not specified in the decorator parameters, the autotuner will use the default value of (2, 1). Otherwise, the autotuner will use the value specified in the decorator parameters. :type cluster_shape_mn: Tuple[int, int], optional :param tolerance: Tolerance value for reference validation comparison, defaults to 1e-01 :type tolerance: float, optional :param warmup_iterations: Number of warmup iterations before benchmarking, defaults to 0 :type warmup_iterations: int, optional :param iterations: Number of benchmark iterations to run, defaults to 1 :type iterations: int, optional :param skip_ref_check: Whether to skip reference result validation, defaults to False :type skip_ref_check: bool, optional :raises RuntimeError: If CUDA GPU is not available :raises ValueError: If the configuration is invalid or unsupported by the kernel :return: Execution time of the GEMM kernel :rtype: float z2Running Blackwell Persistent Dense GEMM test with:zmnkl: z AB dtype: z , C dtype: z , Acc dtype: zMatrix majors - A: z, B: z, C: zMma Tiler (M, N): z, Cluster Shape (M, N): z Tolerance: zWarmup iterations: z Iterations: zSkip reference checking: z-Dynamic persistent mode is not supported yet.zUnsupported testcase z, z$GPU is required to run this example!iWTc s|r|||fn|||f}|rdnd}|tjhv}t|} |tjtjhvr'| ntj} tj|| |tjj j tjj dddd | } | } | j tjd} |tjtjhvrW| n| tj}t|dd }||_|rr|j|rnd ndd }tj| |||d }| || | fS) Nr&r(rr(r&rgr()stdscale) permute_order init_type init_configrHr assumed_alignr leading_dimis_dynamic_layout)ror cutlass_torchrIrrtorchbfloat16create_and_permute_torch_tensorTensorInitTypeGAUSSIANGaussianInitConfigtorfloat32viewuint8rrmark_layout_dynamicconvert_cute_tensor)rmode0mode1rrIr r_r is_unsigned torch_dtype gen_dtypetorch_tensor_cpu torch_tensorf32_torch_tensortorch_tensor_view cute_tensor)rrCrDcreate_and_permute_tensor sD        z&run..create_and_permute_tensorrr rNr9r(rr)rIdevicer+r)r[z mkl,nkl->mnlrr)rrrrgh㈵>)atolrtol)do_benchr&g?r)r(cs<tjdddddfdddddfjtjdS)Nr)scale_ascale_b out_dtype)r  _scaled_mmmTr rC)a_torchb_torchscale_abrCrDr szrun..cs"tddddddjSNr&rr()r matmulpermuter+rC)r,r-rCrDr s"c s.tdddddddddjSr/)r baddbmmr1r+rC)r,r-c_torch_convertrCrDr s   )warmuprepgeAzCuBLAS Average time: z.3fz ms, TFLOPS: z.1fc sSrOrCrC) compiled_gemmcurrent_streamepi_argsrrrrrrGrCrDr szCute-DSL Average time: )T)7printrrrr r is_available RuntimeError manual_seedrorw HardwareInfoget_max_active_clusterszerosint32rrrrXr> AddressSpacegmemrrr7r cuda_streamcompilerrrreinsumcpurr rrrSKIPrrrrrrItesting assert_closetriton.testingr&timesleeprPonesr)6rrrnrrr rrrrr"rrrrrrkwargsrrrr"a_ref a_torch_cpub_ref b_torch_cpurd_torch d_torch_cpucc_torchr#r^ hardware_inforr[ torch_streamrefgpu_dref_drr_f8_torch_tensor ref_d_tensorr&flopsrepeatsr4rK fn_cublas timing_cublas tflops_cublasfntimingtflopsrC)r,r-r3r6r7r8rrrrrr.rrGrDrun| s5  J  3                  rf__main__rrG.cCs4z tdd|dDWStytdw)Ncss|] }t|VqdSrO)ristriprPrCrCrDr sz-parse_comma_separated_ints..,z2Invalid format. Expected comma-separated integers.)rsplit ValueErrorargparseArgumentTypeError)rrCrCrDparse_comma_separated_ints s   rnz.Example of Dense Persistent GEMM on Blackwell.) descriptionz--mnkl)r'r'rOr(z!mnkl dimensions (comma-separated))typedefaulthelpz--mma_tiler_mn)rHrHz Mma tile shape (comma-separated)z--cluster_shape_mn)r(r(zCluster shape (comma-separated)z --ab_dtype)rprqz --d_dtypez --c_dtypez --acc_dtypez --a_majorrr)choicesrprqz --b_majorrz --d_majorz --c_majorz --tolerancegQ?zTolerance for validationz--warmup_iterationsr)zWarmup iterationsz --iterationsz&Number of iterations to run the kernelz--skip_ref_check store_truezSkip reference checking)actionrrz--dynamic_persistentzDynamic persistent kernelr+z$--mnkl must contain exactly 4 valuesr&z,--mma_tiler_mn must contain exactly 2 valuesz0--cluster_shape_mn must contain exactly 2 valuesPASS)rrrrr(FF)brltypingrrrrrr functoolsrcuda.bindings.driverbindingsdriverrr ro cutlass.cuter>cutlass.cute.nvgpur r cutlass.torchr cutlass.pipelinergcutlass.utils.blackwell_helpersrwblackwell_helpersrT cutlass.utils.blockscaled_layoutblockscaled_layoutrcutlass.cute.nvgpu.warpr r r rrrrr cutlass.utilsrcutlass.cute.runtimerrquack.pipelinerquack.cute_dsl_utilsrrquack.tile_schedulerrquack.varlen_utilsrrquack.gemm_sm90rrquack.copy_utilsrquack.sm100_utilsrrrirrfloatrrfrrnArgumentParserparser add_argumentrIrp parse_argsargsr<rerrorr"rrrnrrr rrrrrrrrrr9rCrCrCrDs2          B0        0