)`i`dZddlZddlZddlZddlZ ddlZeedredkr e dddl Tddl Tddl Tddl TddlTddlTddlTddlTn+#e $r#ddlTddlTddlTddlTddlTddlTddlTddlTYnwxYwejeZGddZGd d ZGd d ZGd dZ GddZ!dS)zh Utilities for filtering CUTLASS library kernels and emitting library intitialization and building code NCUTLASS_IGNORE_PACKAGETz+Disabling attempt to import cutlass_library)*c*eZdZdZdZdZdZdZdS)EmitOperationKindAlla  Emit the OperationKind-level CUTLASS library initialization code. The code is generated in the {generated_path}/{operation_kind} directory (e.g., tools/library/generated/gemm in the build directory, for OperationKind=Gemm), in the all_{operation_kind}_operations.cu file (e.g., all_gemm_operations.cu for OperationKind=Gemm). That file declares several functions in namespace cutlass::library. The functions all have this form, void initialize_{configuration_name}(Manifest& manifest); The file also _defines_ the following function in that namespace. void initialize_all_{operation_kind}_operations(Manifest& manifest); That function calls all of the functions declared in this file. Those functions are defined in subdirectories (which this class does not create). cv||_||_||_d|_d|_d|_d|_d|_dS)N( /* Generated by manifest.py - Do not edit. */ #include "cutlass/cutlass.h" #include "cutlass/library/library.h" #include "cutlass/library/manifest.h" namespace cutlass { namespace library { /////////////////////////////////////////////////////////////////////////////////////////////////// zv // // Entry point to construct operations // void initialize_all_${operation_name}_operations(Manifest &manifest) { ;void initialize_${configuration_name}(Manifest &manifest); . initialize_${configuration_name}(manifest); } /////////////////////////////////////////////////////////////////////////////////////////////////// } // namespace library } // namespace cutlass )generated_pathkindargsheader_templateentry_template configuration_prototype_templateconfiguration_templateepilogue_template)selfr r rs /home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/flashinfer/data/cutlass/python/cutlass_library/manifest.py__init__zEmitOperationKindAll.__init__ZsO(DDIDID D-kD)!RDDctdtj|jt |j|_tdt|jztj |jdtj|jdt |jd|_ tdt|j t|j d|_ |j |j|j g|_g|_|S) Nz#*** EmitOperationKindAll::__enter__z,*** operation_path (directory to create): T)exist_okall__operations.cu&*** top_level_path (file to write): w)_LOGGERdebugospathjoinr OperationKindNamesr operation_pathstrmakedirstop_level_pathopentop_level_filewriter source_filesconfigurationsrs r __enter__zEmitOperationKindAll.__enter__s MM7888',,t':z,EmitOperationKindAll.emit..s&.Z.Z)&!v.Z.Z.Z.Z.Z.Zrz *** min_cc=z*** configuration_name=configuration_name) rrlensorteditemsr,appendr)r*SubstituteTemplater)r operationsr4r,r7r5s remitzEmitOperationKindAll.emitsk MM2333 MM=C OO==>>> MM\.Z.ZzGWGWGYGY.Z.Z.Z(Z(Z\\]]]"()9)9););"<"<JJ mm,F,,---#1#7#7#9#9JJ  a H4FHHIII ""#5666 !!"4T5Z]qtF]G#I#I J J J JJJJrctd|jt |jdt |ji|jD]1}|jt |j d|i2|j|j |j dS)Nz"*** EmitOperationKindAll::__exit__operation_namer7) rrr)r*r<rr#r r,rrclose)rexception_typeexception_value tracebackr7s r__exit__zEmitOperationKindAll.__exit__s MM677701DGWYklpluYvFwxxyyy"1}}  243NQegyPz { {||||d4555rN__name__ __module__ __qualname____doc__rr.r>rEr2rrrrEs^(&&&R* J J J      rrc*eZdZdZdZdZdZdZdS)EmitOperationKindLibrarya\ Emit the CUTLASS library initialization code for each OperationKind. The code is generated in the directory {generated_path}/{operation_kind}/{min_cc} (e.g., tools/library/generated/gemm/90 in the build directory, for min_cc=90 and OperationKind=Gemm), in the file all_sm{min_cc}_{operation_kind}_operations.cu (e.g., all_sm90_gemm_operations.cu for min_cc=90 and OperationKind=Gemm). The min_cc variable here indicates the minimum GPU architecture version that the things to be initialized require. For example, min_cc=90 indicates sm90. That file declares several functions in namespace cutlass::library. The functions all have this form, void initialize_all_sm{min_cc}_{subclass_name}_{extended_name}_operations(Manifest& manifest); where extended_name is operation.extended_name() for all the operations given to the emit method (which see below). (All operations for a given configuration_name are guaranteed to have the same extended_name().) The file also _defines_ the following function in that namespace. void initialize_all_sm{min_cc}__{operation_kind}_operations(Manifest& manifest); That function calls all of the functions declared in this file. Those functions are defined in subdirectories. The mapping from OperationKind to emitter handles the details of what happens in each of those subdirectories. c||_||_||_||_tjt tjttj ttj ttj ttjt tjt$i|_d|_d|_d|_d|_d|_d|_d|_dS)Nrz // // Entry point to construct operations // void initialize_all_sm${min_cc}_${subclass_name}_${operation_name}_operations(Manifest &manifest) { r r zV initialize_all_sm${min_cc}_${subclass_name}_${operation_name}_operations(manifest); zcvoid initialize_all_sm${min_cc}_${subclass_name}_${operation_name}_operations(Manifest &manifest); r )r r4r r OperationKindGemmEmitGemmConfigurationLibraryConv2dEmitConv2dConfigurationLibraryConv3dEmitConv3dConfigurationLibraryRankKEmitRankKConfigurationLibraryRank2KEmitRank2KConfigurationLibraryTrmmEmitTrmmConfigurationLibrarySymmEmitSymmConfigurationLibraryemittersrrrrsubclass_call_templatesubclass_prototype_templater)rr r4r rs rrz!EmitOperationKindLibrary.__init__s(DDKDIDI6::8:66DMDD-kD)"SD"{D(ND$Drctdtdt|jtdt|jtd|jtj |jt|jt|j|_ tdt|j tj |j tj |j d|jdt|jd|_ td t|j t|j d |_|j|ji|_i|_i|_|S) Nz'*** EmitOperationKindLibrary::__enter__z*** generated_path: z *** OperationKindNames[kind]: z*** min_cc: z**** operation_path (directory to make): all_smr5rrr)rrr%r r#r r4r r!r"r$r&r'r(r)r*rr+subclass_filessubclass_configurationsr-s rr.z"EmitOperationKindLibrary.__enter__s MM;<<< MME3t/B+C+CEEFFF MMT5G 5RTTUUU MM04;00111',,t':configuration_pathr;r<r) rr7r=re subclass_pathsubclass_top_level_path subclass_dirconfiguration_emitter operations rr>zEmitOperationKindLibrary.emit"sF MM6777 MMC/ACCDDD z??Q     qM//11M MM7-GHHHD///gll4#6 FFm mmBc-.@.@BBCCCh}46d"=1!# k kkmkk>PQUQZ>[kkk!m!m mm'),-D)E)EFGGG,00G+M+Md-( -(..t/CDDD*A)Bd &7??4#6}#E#JKKL MM(3|+<+<<=== !ty !,0B C CXG\!..)""9---- mmF-@AABCCC  &--.C.VWWW XXXXXXXXXXXXXXX  /667IJJJ &,,-?@eh|QhR.T.TUUUUUs!A/JJ #J c ttdt|jD]Z\}}t |j|t|jd}|j t|j |[|j t|j t |jdt|jdt|jD]\}}t |j|t|jd}| t|j ||j|D],}| t|jd|i-| |j||j t|j||j |j|j dS)Nz&*** EmitOperationKindLibrary::__exit__)r4 subclass_namer@r7)rrr9rbr:r%r4r#r r)r*r<r_rrcrrrAr^)rrBrCrDrp subclass_file subclass_cfg configurations rrEz!EmitOperationKindLibrary.__exit__Ps8 MM:;;;(.t/B/H/H/J/J(K(Kdd$ }dk""&,TY7l   243SUa b bcccc,dk"",TY7//     )/t/B/H/H/J/J(K(K__$ }dk""&,TY7l ,T-@,OOPPP7 F- T8 -;        $0111  243NP\ ] ]^^^^d4555rNrFr2rrrLrLs_>000f:+U+U+U\' ' ' ' ' rrLc*eZdZdZdZdZdZdZdS)EmitInterfaceLibraryad Emit the topmost-level CUTLASS library initialization code. The code is generated in the generated_path directory (e.g., tools/library/generated in the build directory), in the initialize_all.cpp file. That file declares several functions in namespace cutlass::library. The functions all have this form, void initialize_all_{operation_kind}_operations(Manifest& manifest); where {operation_kind} abbreviates the "kind" of operation (e.g., gemm for matrix-matrix multiply, conv2d for 2-d convolution, or trmm for triangular solve with multiple right-hand sides). The definitions of these functions live in subdirectories. The file also _defines_ the following function in that namespace. void initialize_all(Manifest& manifest); That function first prepares the manifest, and then calls all of the functions declared in this file. c||_||_g|_g|_t ||_d|_d|_d|_d|_ d|_ dS)Nz0 /* Generated by manifest.py - Do not edit. */ z #include "cutlass/library/library.h" #include "cutlass/library/manifest.h" namespace cutlass { namespace library { ${prototypes} zP void initialize_all_${kind}_operations(Manifest &manifest) { ${fn_calls} } zg void initialize_all(Manifest &manifest) { manifest.reserve(${operation_count}); ${fn_calls} } z1 } // namespace library } // namespace cutlass ) r r prototypesfn_callsr%operation_counttop_level_hdr_templatetop_level_prologuetop_level_initialize_kindtop_level_initializetop_level_suffix)rr rzrs rrzEmitInterfaceLibrary.__init__sg(DDIDODM//D#D D&D" !DDrcxtdtj|jd|_tdt|jzt|jd|_ |j |j |jg|_ |S)Nz#*** EmitInterfaceLibrary::__enter__zinitialize_all.cppz*** top_level_path: r) rrr r!r"r r'r%r(r)r*r{r+r-s rr.zEmitInterfaceLibrary.__enter__s MM7888',,t':zEmitInterfaceLibrary.emits MM2333 MM*^;<<<O-R.)++,,, M+D(**+++++rc td|jt |jdd|ji|jt |j|j d|j d|j|j |j dS)Nz"*** EmitInterfaceLibrary::__exit__rx )rzry) rrr)r*r<r|r"rxr~rzryrrA)rrBrCrDs rrEzEmitInterfaceLibrary.__exit__s MM677701H. sg6O6O`dc$**S//!:L:R:RSV:W:WXY:Z6[6[6O6O6O6O6O6Or)falseFalse0Fallcbg|]+}t|jdv)|,S)r)r#r=r)r3xrs r z%Manifest.__init__..s="u"u"uASTUAVZ^ZiZoZopsZtZtAtAt1AtAtAtrcg|] }|dk| Srqr2r3rs rrz%Manifest.__init__..!sKKK1177Q777rcg|] }|dk| Srr2rs rrz%Manifest.__init__..#s!W!W!WqTVww!wwwrcg|] }|dk| Srr2rs rrz%Manifest.__init__..$s$"Y"Y"YQRVXQXQX1QXQXQXrz6Using {filter_count} kernel filters from {filter_file}) filter_count filter_filemax))r=rrzoperations_by_name kernel_filterkernel_filter_list kernel_namesoperations_enabledselected_kernelsignore_kernel_namesexclude_kernel_namescompute_capabilities_baseline compute_capabilities_feature_setcurr_build_dir filter_by_cckernels architectures RuntimeErrorr8rr9setrNrOrQrSrUrYr[ignore_kernelsexclude_kernelskernel_filter_fileget_kernel_filtersrrformatdisable_full_archs_compilationinstantiation_levelis_kernel_filter_set_to_allr ValueError)rroperations_lists ` rrzManifest.__init__siDODID DD DD DD!D "D*,D&-1GD)DD y3'9,d /d " " "TW[Wiijj jORSWSeOfOf.sd.@.F.Fs.K.K.Kmqlsd++1#6O6OhliN6O6O6O3O3O,P,Pd( 5 5 5! E ! !"$     !   #v"u"u"uo"u"u"u   KK (:(:3(?(?KKK!W!WT-@-F-Fs-K-K!W!W!Wd"Y"Yd.B.H.H.M.M"Y"Y"Yd ($&$ ! !$($;$;D.Us2VVVtQTAUAUVT[[]]VVVrcFg|]}|t|Sr2)rrrs rrz/Manifest.get_kernel_filters..Ws)<< ? ? FFFFGGG TTTTUUU0ZZ $ $[$ 7 7Z UUUkUUUVVV  XXXXXXYYYY Nrc:||rE|j|||j|<|}|j}|j|j vri|j|j<||j|jvri|j|j|<||j|j| vrg|j|j||<|j|j||||xj dz c_ dSt d |dS)zT Inserts the operation. operation_kind -> configuration_name -> [] zCulled {} from manifestN)rrr;rrr7rrr=rrzrrr)rrnr7r4s rr;zManifest.appends} {{9S ""9#<#<#>#>???=Fdi7799:%7799~f  !)=)=)?)? ? ?46 01 ty'?@ @ @<> 01&9 4?93K#LV#T#Y#Y#[#[ [ [PR 01&9:LM oi./78JKRRS\]]] a mm-44Y5N5N5P5PQQRRRRRrc t|d5}tdi}||dz|dt|ddzt j|jd}|j D]a}t|}t j||d|d dd} |d | d b|d |j D](}t|j | D]} t|||  D]} td t| t|| d}||dz||| | D];} |dt| ddz<|d|j r|||* ddddS#1swxYwYdS)Nrz:cutlass_target_sources(cutlass_library_objs PRIVATE z z %s \/ generatedrrz rz) zJcutlass_add_cutlass_library( SUFFIX ${kind}_sm${min_cc}_${subclass} )r4r subclassz) )r(r<r*r%replacer r!r"rr=rr#r9r#emit_disable_full_archs_compilation) r manifest_pathr'r+ manifest_file target_textr r kind_str all_kind_filer4r source_files remit_manifest_cmakezManifest.emit_manifest_cmakes mS ! !R]&( k+.///*s>+A+A$+L+L'M'MMNNNw||D$7EEn/&&((66$%d+ ^X?^h?^?^?^__gghlnqrr 4=4445555'"""/&&((RR$T_T27799:: R RF d!3F!;!@!@!B!BCC ' 'h,.V&8&>HVVXXK    f 4 5 5 5+D1&9(C T T !!*s;3F3FtS3Q3Q/R/R"RSSSS    & & & &  0 R  4 4]L Q Q Q RRRRRRRRRRRRRRRRRRRRsII99I=I=c d}d}d}d}d}d}t|D]}||D]} || r|| r|hd| } nU|| r|dh| } n<|| r|d d h| } n"td | |d t | d d| ȌdS)NcdSr1r2rhs r for_hopperz@Manifest.emit_disable_full_archs_compilation..for_hoppers $rc*d|vpd|vp d|vpd|vod|vS)N1681616832168641688tf32r2rs r for_amperez@Manifest.emit_disable_full_archs_compilation..for_amperes=D6T/6T/6T>4fn 6rcd|vrd|vpd|vS)Nr r 8816r2rs r for_turingz@Manifest.emit_disable_full_archs_compilation..for_turings$D.7V4%7!D. !rc d|vS)N884r2rs r for_voltaz?Manifest.emit_disable_full_archs_compilation..for_volta s $ rc,|dS)Nz.cpp)endswithrs ris_cppz.is_cppsv&& &rc$|ttjz}|tkr;td|ttj|dt t|S)Nz Empty archs set for file {} after taking the intersection of {} (global requested archs) and {} (per file requested archs)  )rrrrrr"mapr%)archsrintersected_archss r,get_src_archs_str_given_requested_cuda_archszbManifest.emit_disable_full_archs_compilation..get_src_archs_str_given_requested_cuda_archss#c$*L&M&MM  #%% ' ' {C0R,S,SUZ[[ ]]]XXc#'899:::r>PWZKFHzRPer file archs are not set {}, as there is no rule specified for this file patternz1cutlass_apply_cuda_gencode_flags({} SM_ARCHS {}) rr)r9rrrr*r%r) rr+rr rrrrr4r archs_strs rrz,Manifest.emit_disable_full_archs_compilations666 !!!''' ; ; ;<,,..// M M&'/ M MKvk"" MK(( MHHWbcc K(( MHH"{[[ ;'' MHH"bS^__ "#w#~#~@K$L$LMMM    T [ [\_`k`s`stxz}`~`~\\BK!L!L M M M M M M Mrctjti}tjti}tjti}t j|jd}t j |rtj |t j |||||j |j5}|j}|jD]"}|t&|# dddn #1swxYwYi} |jD]4} i| | <|j| D] } i| | | <5|jD]\}} t+| D] \} } |||| ||j5}| D]a\}}t,d|dt1|dt1|dkrdndd|||b|jD]X\}}|| || vrg| || |<| || ||j|Y dddn #1swxYwY"|||||j5}|| dddn #1swxYwYt j|d }|||| dS) Nrz Emitting z with z operationrrqsrzmanifest.cmake)GeneratorTargetLibraryrLrrvr r!r"rexistsshutilrmtreerfrzrr'r=rr>r#r:r9rinfor8r+extendr)rtargetoperation_emitters kind_emittersinterface_emittersr iface_emitterr'rr+r r4opsr,operation_kind_emitterr7r=rfilesrs rr>z Manifest.emit-sg7 3M 3W\\$"5{CCN w~~n%%$ mN###H^ # F #ND4H$) T T?Xe$3n O0022??.-n=>>>>???????????????? L$$&&((l4OD)..00((&%' T6""( $4466))$*399;;$7$7 q q &. '  'PTPY Z Z q^t0>0D0D0F0F H H, * LLA%7AAs:AA`cdn`o`ost`t`tZ\Z\z}AAA B B B " ' '(:J G G G G!7!D!J!J!L!L q qoh|N;FCCC?Al>*628<  ( 0 : A ABXBefnBo p p p p q  q q q q q q q q q q q q q q q !=  K K)Oe##C((()))))))))))))))GLL1ABBM]NLIIIIIs8ADDDC)KK K )L  L L r1)rrr)rGrHrIrrrrrrrr;rrr%r&r>r2rrrrsE'E'E'E'N... *GGGVSSSDRRR<-M-M-M`*15J5J5J5J5J5Jrr)"rJenumloggingos.pathr r(builtinshasattrr ImportErrorcutlass_library.librarycutlass_library.gemm_operation cutlass_library.rank_k_operation!cutlass_library.rank_2k_operationcutlass_library.trmm_operationcutlass_library.symm_operation cutlass_library.conv2d_operation cutlass_library.conv3d_operationlibrarygemm_operationrank_k_operationrank_2k_operationtrmm_operationsymm_operationconv2d_operationconv3d_operation getLoggerrGrrrLrvrrr2rrrKsoB   !/// WX/00E5Kt5S5S +C D DD''''....00001111........000000000!!!    !!!!          ! ' H % %j j j j j j j j ZE E E E E E E E Ni i i i i i i i \        tJtJtJtJtJtJtJtJtJtJsAA%A>=A>