
    )`i                     `&   d Z ddlZddlZ	 ddlmZ n# e$ r dadefdZY nw xY w G d dej                  Z	e	j
        diZ G d	 d
ej                  Zej        dej        dej        dej        dej        dej        dej        dej        dej        dej        dej        diZi ej        dej        dej        dej        dej        dej        dej        dej         dej!        dej"        dej#        d ej$        d!ej        d"ej%        d#ej        dej        dej        di ej        dej        dej&        d$ej'        d%ej(        d&ej)        d'ej*        d(ej        d)ej+        d*ej        d+ej,        d,ej        d-ej-        d.ej.        d/ej        d0ej/        d1ej        d2ej0        d3ej1        d4ej2        d5ej3        d6ej4        d7ej5        d8ej6        d9ej7        d:ej8        d;ej9        d<ej:        d=ej;        d>iZ<i ej        dej        d?ej        d@ej        dAej        dBej        dCej        dDej         dEej!        dFej"        dGej#        dHej$        dIej        dJej%        dKej        dLej        dMej        dNi ej        dOej        dPej&        dQej'        dRej(        dSej)        dTej*        dUej        dVej+        dWej        dXej,        dYej        dZej-        d[ej.        d\ej        d]ej/        d^ej        d_ej0        d`ej1        daej2        dbej3        dcej4        ddej5        deej6        dfej7        dgej8        dhej9        diej:        djej;        dkiZ=i ej        dej        dlej        dmej        dnej        doej        dpej        dqej         drej!        dmej"        dnej#        doej$        dpej        dqej%        drej        doej        doej        doi ej        dsej        dnej&        dsej'        dsej(        dnej)        doej*        doej        dpej+        dpej        dqej,        dqej        drej-        dqej.        dqej        drej/        dqej        dtej0        dnej1        doej2        dpej3        dqej4        drej5        dtej6        dnej7        doej8        dpej9        dqej:        drej;        dtiZ> G du dvej                  Z?e?j@        dwe?jA        dxiZB G dy dzej                  ZCeCjD        d{eCjE        d|iZFeCjD        d}eCjE        d~iZGej        ej-        fej        ej        fej        ej        fgZHd ZId ZJd ZKd ZLd ZMd ZNd ZO G d dej                  ZP G d dej                  ZQeQjR        deQjS        deQjT        deQjU        deQjV        deQjW        deQjX        deQjY        deQjZ        deQj[        deQj\        deQj]        diZ^ G d dej                  Z_i e_j`        de_ja        de_jb        de_jc        de_jd        de_je        de_jf        de_jg        de_jh        de_ji        de_jj        de_jk        de_jl        de_jm        de_jn        de_jo        de_jp        de_jq        de_jr        de_js        diZte_j`        e_ja        e_ja        e_j`        e_jb        e_jc        e_jc        e_jb        e_jd        e_je        e_je        e_jd        e_jf        e_jg        e_jg        e_jf        e_ji        e_ji        i	Zui e_j`        de_jb        de_jd        de_jf        de_ja        de_jc        de_je        de_jg        de_jh        de_ji        de_jj        de_jk        de_jl        de_jm        de_jo        de_jn        de_jp        de_jq        de_jr        de_js        diZve_j`        eCjD        fde_j`        eCjE        fde_ja        eCjD        fde_ja        eCjE        fdiZw G d dej                  Zxi exjy        dexjz        dÓexj{        dēexj|        dœexj}        dƓexj~        dǓexj        dȓexj        dɓexj        dʓexj        d˓exj        d̓exj        d͓exj        dΓexj        dϓexj        dГexj        dѓexj        dғi exj        dӓexj        dԓexj        dՓexj        d֓exj        dדexj        dؓexj        dٓexj        dړexj        dۓexj        dܓexj        dݓexj        dޓexj        dߓexj        dexj        dexj        dexj        di exj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        diZi exjy        dexjz        dexj{        dexj|        dexj}        dexj~        dexj        d exj        dexj        dexj        dexj        dexj        dexj        d exj        dexj        dexj        dexj        di exj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        d	exj        dexj        dexj        dexj        dexj        d
exj        dexj        dexj        di exj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        d
exj        dexj        d
exj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        dexj        diZ G d dej                  Zi ejy        dej        dej        dej        dej        dej        dej        dej        dej        dej        d ej        d!ej        d"ej        d#ej        d$ej        d%ej        d&Zi ejy        dej        dej        d'ej        d'ej        d'ej        d'ej        d'ej        d'ej        d(ej        d(ej        dej        d(ej        d)ej        d*ej        d(ej        d(Z G d+ d,ej                  Zej        d-ej        d.iZĐd/ ZŐd0 Z G d1 d2ej                  Zej        dej        d3ej        d4iZej        dej        dej        d5iZ G d6 d7ej                  Zej        d8ej        d9iZej        d:ej        d;iZ G d< d=ej                  Zej        d>ej        d?iZej        d@ej        dAiZ G dB dCej                  Zej        dDej        dEiZej        dFej        dGiZ G dH dIej                  Zej        dJej        dKej        dLej        dMej        dNiZej        dOej        dPej        dQej        dRej        dSiZ G dT dUej                  Zej        dVej        dWej        dXej        dYej        dZej        d[ej        d\iZ G d] d^ej                  Zd_d`d`dadbdcdddedfZdgdgdrdhdidhdidjdkZdl Z G dm dnej                  Zej        dVej        doej        dVej        dVej        doej        dpej        dqej        drej        dVej        drej        drej        dVej        driZ G ds dtej                  Zej        dWiZ  G du dvej                  Zej        dYiZ G dw dxej                  Zej        dZiZ G dy dzej                  Zej        d{ej        d|iZ G d} d~ej                  Z G d dej                  Z	e	j
        de	j        de	j        de	j        de	j        de	j        de	j        de	j        de	j        di	Z G d dej                  Zej        dej        diZej        dej        diZ G d dej                  Zej        dej        dej        diZej        dej        dej        diZ G d dej                  Z G d dej                  Z e j!        de j"        de j#        de j$        de j%        diZ&e j!        de j"        de j#        de j$        de j%        diZ' G d dej                  Z(e(j)        de(j*        de(j+        diZ,e(j)        de(j*        de(j+        diZ- G d dej                  Z.e.j/        de.j0        de.j1        de.j2        diZ3e.j/        de.j0        de.j1        de.j2        diZ4g dZ5 G d d          Z6 G d d          Z7 G d d          Z8 G d d          Z9 G d d          Z: G d d¦          Z;dÄ Z< G dĄ dej                  Z= G dƄ dej                  Z>dS (  z;
Data types and tags used for emitting CUTLASS C++ kernels
    N)autoreturnc                  (    t           } t           dz  a | S )N   )__cutlass_library_auto_enum)is    /home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/flashinfer/jit/gemm/cutlass/cutlass_library.py	enum_autor
   2   s    '#q(#    c                   "    e Zd Z e            ZdS )GeneratorTargetN)__name__
__module____qualname__r
   Library r   r	   r   r   =           ikkGGGr   r   libraryc                      e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
 e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z  e            Z! e            Z" e            Z# e            Z$ e            Z% e            Z& e            Z' e            Z( e            Z) e            Z* e            Z+ e            Z, e            Z- e            Z. e            Z/ e            Z0 e            Z1 e            Z2dS )DataTypeN)3r   r   r   r
   voidb1u2u4u8u16u32u64s2s4s8s16s32s64e4m3e5m2f8f6f4e3m2e2m3e2m1ue8m0ue4m3f16bf16f32tf32f64cf16cbf16cf32ctf32cf64cs2cs4cs8cs16cs32cs64cu2cu4cu8cu16cu32cu64invalidr   r   r	   r   r   I   s       9;;D	B	B	B	B
)++C
)++C
)++C	B	B	B
)++C
)++C
)++C9;;D9;;D	B	B	B9;;D9;;D9;;DIKKEIKKE
)++C9;;D
)++C9;;D
)++C9;;DIKKE9;;DIKKE9;;D
)++C
)++C
)++C9;;D9;;D9;;D
)++C
)++C
)++C9;;D9;;D9;;DikkGGGr   r   r   r%   r&   hsdczr'   r(   r)   r   r   r   r   r   r   r   r   r   r    r!   r"   r#   r$   r+   r*   r,   r-   r.   r/   r0   r1   r2   r3   r4   r5   r6   r7   r8   r?   r@   rA   rB   rC   rD   r9   r:   r;   r<   r=   r>   zcutlass::uint1b_tzcutlass::uint2b_tzcutlass::uint4b_tuint8_tuint16_tuint32_tuint64_tzcutlass::int2b_tzcutlass::int4b_tint8_tint16_tint32_tint64_tzcutlass::float_e4m3_tzcutlass::float_e5m2_tz%cutlass::type_erased_dynamic_float8_tz%cutlass::type_erased_dynamic_float6_tz%cutlass::type_erased_dynamic_float4_tzcutlass::float_e2m3_tzcutlass::float_e3m2_tzcutlass::float_e2m1_tzcutlass::float_ue8m0_tzcutlass::float_ue4m3_tzcutlass::half_tzcutlass::bfloat16_tfloatzcutlass::tfloat32_tdoublez!cutlass::complex<cutlass::half_t>z%cutlass::complex<cutlass::bfloat16_t>zcutlass::complex<float>z%cutlass::complex<cutlass::tfloat32_t>zcutlass::complex<double>z#cutlass::complex<cutlass::uint2b_t>z#cutlass::complex<cutlass::uint4b_t>z"cutlass::complex<cutlass::uint8_t>z#cutlass::complex<cutlass::uint16_t>z#cutlass::complex<cutlass::uint32_t>z#cutlass::complex<cutlass::uint64_t>z"cutlass::complex<cutlass::int2b_t>z"cutlass::complex<cutlass::int4b_t>z!cutlass::complex<cutlass::int8_t>z"cutlass::complex<cutlass::int16_t>z"cutlass::complex<cutlass::int32_t>z"cutlass::complex<cutlass::int64_t>r                   @         c                   6    e Zd Z e            Z e            ZdS )BlasModeN)r   r   r   r
   	symmetric	hermitianr   r   r	   r^   r^   !  s$        	I	IIIr   r^   zcutlass::BlasMode::kSymmetriczcutlass::BlasMode::kHermitianc                   6    e Zd Z e            Z e            ZdS )ComplexTransformN)r   r   r   r
   noneconjr   r   r	   rb   rb   .  s$        9;;D9;;DDDr   rb   z cutlass::ComplexTransform::kNonez%cutlass::ComplexTransform::kConjugatezcute::identityzcute::conjugatec                 D     t           fdt          D                       S )Nc              3   *   K   | ]\  }}|k    V  d S Nr   ).0_rrI   	data_types      r	   	<genexpr>zis_complex.<locals>.<genexpr>I  s+      @@%"ayA~@@@@@@r   )anyRealComplexBijectionrj   s   `r	   
is_complexro   H  s'    @@@@+?@@@@@@r   c                 6    | t           j        t           j        fv S rg   )GemmKindBlockScaledUniversal3xGroupedBlockScaledUniversal3x	gemm_kinds    r	   is_block_scaledrv   L  s     '.  r   c                 6    | t           j        t           j        fv S rg   )rq   BlockwiseUniversal3xGroupedBlockwiseUniversal3xrt   s    r	   is_blockwiserz   S  s     %,  r   c                 L    | t           j        t           j        t           j        fv S rg   )rq   GroupedUniversal3xrs   ry   rt   s    r	   
is_groupedr}   Z  s&    #.,  r   c                 H    t           D ]\  }}| |k    r|c S t          j        S rg   rm   r   rE   )	real_typerrI   s      r	   get_complex_from_realr   c  s5    $  1>>HHH r   c                 H    t           D ]\  }}| |k    r|c S t          j        S rg   r   )complex_typer   rI   s      r	   get_real_from_complexr   k  s7    $  11HHH r   c                 l    | t           j        k    rdS t          |          dk    rdS dt          |          z  S )Nr   r[   r\   )r   r   DataTypeSizern   s    r	   get_tma_alignmentr   s  s9    HM!!q	i	 A	%	%sl9---r   c                   6    e Zd Z e            Z e            ZdS )ComplexMultiplyOpN)r   r   r   r
   multiply_addgaussianr   r   r	   r   r   }  s$        9;;Ly{{HHHr   r   c                       e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
 e            Z e            Z e            Z e            Z e            ZdS )MathOperationN)r   r   r   r
   r   multiply_add_saturatemultiply_add_mixed_input_upcastxor_popcand_popcmultiply_add_fast_bf16multiply_add_fast_f16multiply_add_fast_f32multiply_add_complex_fast_f32multiply_add_complexmultiply_add_complex_gaussianmultiply_add_fast_accumr   r   r	   r   r     s        9;;L%IKK&/ikk#y{{Hy{{H&Y[[%IKK%IKK$-IKK!$9;;$-IKK!'ikkr   r   zcutlass::arch::OpMultiplyAddz$cutlass::arch::OpMultiplyAddSaturatez,cutlass::arch::OpMultiplyAddMixedInputUpcastzcutlass::arch::OpXorPopczcutlass::arch::OpAndPopcz$cutlass::arch::OpMultiplyAddFastBF16z#cutlass::arch::OpMultiplyAddFastF16z#cutlass::arch::OpMultiplyAddFastF32z*cutlass::arch::OpMultiplyAddComplexFastF32z#cutlass::arch::OpMultiplyAddComplexz+cutlass::arch::OpMultiplyAddGaussianComplexz%cutlass::arch::OpMultiplyAddFastAccumc                      e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
 e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            ZdS )
LayoutTypeN)r   r   r   r
   ColumnMajorRowMajorColumnMajorInterleaved2RowMajorInterleaved2ColumnMajorInterleaved32RowMajorInterleaved32ColumnMajorInterleaved64RowMajorInterleaved64	TensorNWC
TensorNHWCTensorNDHWC
TensorNCHWTensorNGHWCTensorNC32HW32TensorNC64HW64TensorC32RSK32TensorC64RSK64	TensorKCS
TensorKCSRTensorKCSRTr   r   r	   r   r     s        )++Ky{{H'ikk$9;;(y{{%IKK(y{{%IKK	IJ)++KJ)++KY[[NY[[NY[[NY[[N	IJ)++KKKr   r   zcutlass::layout::ColumnMajorzcutlass::layout::RowMajorz*cutlass::layout::ColumnMajorInterleaved<2>z'cutlass::layout::RowMajorInterleaved<2>z+cutlass::layout::ColumnMajorInterleaved<32>z(cutlass::layout::RowMajorInterleaved<32>z+cutlass::layout::ColumnMajorInterleaved<64>z(cutlass::layout::RowMajorInterleaved<64>zcutlass::layout::TensorNWCzcutlass::layout::TensorNHWCzcutlass::layout::TensorNDHWCzcutlass::layout::TensorNCHWzcutlass::layout::TensorNGHWCz!cutlass::layout::TensorNCxHWx<32>z!cutlass::layout::TensorCxRSKx<32>z!cutlass::layout::TensorNCxHWx<64>z!cutlass::layout::TensorCxRSKx<64>zcutlass::layout::TensorKCSzcutlass::layout::TensorKCSRzcutlass::layout::TensorKCSRTnn2n32n64tt2t32t64nwcnhwcndhwcnchwnghwcnc32hw32nc64hw64c32rsk32c64rsk64kcskcsrkcsrtc                   n   e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
 e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z  e            Z! e            Z" e            Z# e            Z$ e            Z% e            Z& e            Z' e            Z( e            Z) e            Z* e            Z+ e            Z, e            Z- e            Z. e            Z/ e            Z0 e            Z1 e            Z2 e            Z3 e            Z4 e            Z5 e            Z6 e            Z7 e            Z8 e            Z9 e            Z: e            Z;dS )KernelScheduleTypeN)<r   r   r   r
   ScheduleAuto
MultistageCpAsyncWarpSpecializedCpAsyncWarpSpecializedPingpong!CpAsyncWarpSpecializedCooperativeTmaTmaWarpSpecializedTmaWarpSpecializedPingpongTmaWarpSpecializedCooperativeTmaWarpSpecializedFP8FastAccum)TmaWarpSpecializedCooperativeFP8FastAccum&TmaWarpSpecializedPingpongFP8FastAccumImplicitTmaWarpSpecializedSm90%PtrArrayTmaWarpSpecializedCooperative1PtrArrayTmaWarpSpecializedCooperativeFP8FastAccum"PtrArrayTmaWarpSpecializedPingpong.PtrArrayTmaWarpSpecializedPingpongFP8FastAccum&BlockwiseTmaWarpSpecializedCooperative.PtrArrayBlockwiseTmaWarpSpecializedCooperativeTmaWarpSpecialized1SmSm100TmaWarpSpecialized2SmSm100"ImplicitTmaWarpSpecialized1SmSm100"ImplicitTmaWarpSpecialized2SmSm100"PtrArrayTmaWarpSpecialized1SmSm100"PtrArrayTmaWarpSpecialized2SmSm100-PtrArrayTmaWarpSpecialized1SmBlockScaledSm100-PtrArrayTmaWarpSpecialized2SmBlockScaledSm100&PtrArrayNvf4TmaWarpSpecialized1SmSm100&PtrArrayNvf4TmaWarpSpecialized2SmSm100&PtrArrayMxf4TmaWarpSpecialized1SmSm100&PtrArrayMxf4TmaWarpSpecialized2SmSm100*PtrArrayMxf8f6f4TmaWarpSpecialized1SmSm100*PtrArrayMxf8f6f4TmaWarpSpecialized2SmSm100 SparseTmaWarpSpecialized1SmSm100 SparseTmaWarpSpecialized2SmSm100%BlockScaledTmaWarpSpecialized1SmSm100%BlockScaledTmaWarpSpecialized2SmSm100"Mxf8f6f4TmaWarpSpecialized1SmSm100"Mxf8f6f4TmaWarpSpecialized2SmSm100#BlockwiseTmaWarpSpecialized1SmSm100#BlockwiseTmaWarpSpecialized2SmSm100+PtrArrayBlockwiseTmaWarpSpecialized1SmSm100+PtrArrayBlockwiseTmaWarpSpecialized2SmSm100Mxf4TmaWarpSpecialized1SmSm100Mxf4TmaWarpSpecialized2SmSm100Nvf4TmaWarpSpecialized1SmSm100Nvf4TmaWarpSpecialized2SmSm100*Mxf8f6f4TmaWarpSpecializedCooperativeSm120'Mxf8f6f4TmaWarpSpecializedPingpongSm120&Nvf4TmaWarpSpecializedCooperativeSm120#Nvf4TmaWarpSpecializedPingpongSm120&Mxf4TmaWarpSpecializedCooperativeSm120#Mxf4TmaWarpSpecializedPingpongSm120.F8f6f4SparseTmaWarpSpecializedCooperativeSm120+BlockwiseTmaWarpSpecializedCooperativeSm120(BlockwiseTmaWarpSpecializedPingpongSm120r   r   r	   r   r     s       9;;LJ&Y[[%.Y[["(1	%
)++C"!*$-IKK!%.Y[["09	--6Y[[*%.Y[[",5IKK)8A	5)2&5>Y[[2-6Y[[*5>Y[[2!*!*)2&)2&)2&)2&4=IKK14=IKK1-6Y[[*-6Y[[*-6Y[[*-6Y[[*1:.1:.'0y{{$'0y{{$,5IKK),5IKK))2&)2&*3)++'*3)++'2;)++/2;)++/%.Y[["%.Y[["%.Y[["%.Y[["1:..7ikk+-6Y[[**3)++'-6Y[[**3)++'5>Y[[22;)++//8y{{,,,r   r   z-cutlass::gemm::collective::KernelScheduleAutozcutlass::gemm::KernelMultistagez+cutlass::gemm::KernelCpAsyncWarpSpecializedz3cutlass::gemm::KernelCpAsyncWarpSpecializedPingpongz6cutlass::gemm::KernelCpAsyncWarpSpecializedCooperativezcutlass::gemm::KernelTmaz'cutlass::gemm::KernelTmaWarpSpecializedz/cutlass::gemm::KernelTmaWarpSpecializedPingpongz2cutlass::gemm::KernelTmaWarpSpecializedCooperativez3cutlass::gemm::KernelTmaWarpSpecializedFP8FastAccumz>cutlass::gemm::KernelTmaWarpSpecializedCooperativeFP8FastAccumz;cutlass::gemm::KernelTmaWarpSpecializedPingpongFP8FastAccumz3cutlass::conv::KernelImplicitTmaWarpSpecializedSm90zEcutlass::gemm::KernelTmaWarpSpecializedCooperativeFP8BlockScaledAccumz/cutlass::gemm::KernelTmaWarpSpecialized1SmSm100z/cutlass::gemm::KernelTmaWarpSpecialized2SmSm100z7cutlass::conv::KernelImplicitTmaWarpSpecialized1SmSm100z7cutlass::conv::KernelImplicitTmaWarpSpecialized2SmSm100z7cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmSm100z7cutlass::gemm::KernelPtrArrayTmaWarpSpecialized2SmSm100z5cutlass::gemm::KernelSparseTmaWarpSpecialized1SmSm100z5cutlass::gemm::KernelSparseTmaWarpSpecialized2SmSm100z:cutlass::gemm::KernelTmaWarpSpecialized1SmBlockScaledSm100z:cutlass::gemm::KernelTmaWarpSpecialized2SmBlockScaledSm100z7cutlass::gemm::KernelTmaWarpSpecialized1SmMxf8f6f4Sm100z7cutlass::gemm::KernelTmaWarpSpecialized2SmMxf8f6f4Sm100z8cutlass::gemm::KernelTmaWarpSpecializedBlockwise1SmSm100z8cutlass::gemm::KernelTmaWarpSpecializedBlockwise2SmSm100z@cutlass::gemm::KernelPtrArrayTmaWarpSpecializedBlockwise1SmSm100z@cutlass::gemm::KernelPtrArrayTmaWarpSpecializedBlockwise2SmSm100z3cutlass::gemm::KernelTmaWarpSpecialized1SmMxf4Sm100z3cutlass::gemm::KernelTmaWarpSpecialized2SmMxf4Sm100z3cutlass::gemm::KernelTmaWarpSpecialized1SmNvf4Sm100z3cutlass::gemm::KernelTmaWarpSpecialized2SmNvf4Sm100z:cutlass::gemm::KernelPtrArrayTmaWarpSpecializedCooperativezFcutlass::gemm::KernelPtrArrayTmaWarpSpecializedCooperativeFP8FastAccumz7cutlass::gemm::KernelPtrArrayTmaWarpSpecializedPingpongzCcutlass::gemm::KernelPtrArrayTmaWarpSpecializedPingpongFP8FastAccumzMcutlass::gemm::KernelPtrArrayTmaWarpSpecializedCooperativeFP8BlockScaledAccumzBcutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmBlockScaledSm100zBcutlass::gemm::KernelPtrArrayTmaWarpSpecialized2SmBlockScaledSm100z;cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmNvf4Sm100z;cutlass::gemm::KernelPtrArrayTmaWarpSpecialized2SmNvf4Sm100z;cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmMxf4Sm100z;cutlass::gemm::KernelPtrArrayTmaWarpSpecialized2SmMxf4Sm100z?cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmMxf8f6f4Sm100z?cutlass::gemm::KernelPtrArrayTmaWarpSpecialized2SmMxf8f6f4Sm100z4cutlass::gemm::KernelTmaWarpSpecializedMxf8f6f4Sm120z<cutlass::gemm::KernelTmaWarpSpecializedPingpongMxf8f6f4Sm120z0cutlass::gemm::KernelTmaWarpSpecializedNvf4Sm120z8cutlass::gemm::KernelTmaWarpSpecializedPingpongNvf4Sm120z0cutlass::gemm::KernelTmaWarpSpecializedMxf4Sm120z8cutlass::gemm::KernelTmaWarpSpecializedPingpongMxf4Sm120z.cutlass::gemm::KernelScheduleSparseF8f6f4Sm120z@cutlass::gemm::KernelTmaWarpSpecializedBlockwiseCooperativeSm120z=cutlass::gemm::KernelTmaWarpSpecializedBlockwisePingpongSm120 _cpasync_cpasync_warpspecialized!_cpasync_warpspecialized_pingpong$_cpasync_warpspecialized_cooperative_unspecialized_warpspecialized_warpspecialized_pingpong_warpspecialized_cooperative_warpspecialized_fp8_fastaccum*_warpspecialized_cooperative_fp8_fastaccum'_warpspecialized_pingpong_fp8_fastaccum_1sm_2sm_q_1sm_q_2sm_o_vs32_1sm_o_vs32_2sm_o_vs16_1sm_o_vs16_2sm_cooperative_q_pingpong_q_cooperative_o_vs16_pingpong_o_vs16_cooperative_o_vs32_pingpong_o_vs32_qc                   N   e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
 e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            Z e            ZdS )EpilogueScheduleTypeN)r   r   r   r
   r   EpilogueTransposedNoSmemWarpSpecializedPtrArrayNoSmemWarpSpecializedNoSmemWarpSpecialized1SmNoSmemWarpSpecialized2Sm PtrArrayNoSmemWarpSpecialized1Sm PtrArrayNoSmemWarpSpecialized2Smr   r   TmaWarpSpecialized1SmTmaWarpSpecialized2SmPtrArrayTmaWarpSpecialized1SmPtrArrayTmaWarpSpecialized2Smr   r   r   r   r	   r  r    s        9;;L"%IKK$-IKK!(y{{(y{{'0y{{$'0y{{$"$-IKK!%IKK%IKK$-IKK!$-IKK!)2&,5IKK)))r   r  z3cutlass::epilogue::collective::EpilogueScheduleAutoz!cutlass::gemm::EpilogueTransposedz(cutlass::epilogue::NoSmemWarpSpecializedz0cutlass::epilogue::PtrArrayNoSmemWarpSpecializedz+cutlass::epilogue::NoSmemWarpSpecialized1Smz+cutlass::epilogue::NoSmemWarpSpecialized2Smz3cutlass::epilogue::PtrArrayNoSmemWarpSpecialized1Smz3cutlass::epilogue::PtrArrayNoSmemWarpSpecialized2Smz%cutlass::epilogue::TmaWarpSpecializedz0cutlass::epilogue::TmaWarpSpecializedCooperativez(cutlass::epilogue::TmaWarpSpecialized1Smz(cutlass::epilogue::TmaWarpSpecialized2Smz0cutlass::epilogue::PtrArrayTmaWarpSpecialized1Smz0cutlass::epilogue::PtrArrayTmaWarpSpecialized2Smz8cutlass::epilogue::PtrArrayTmaWarpSpecializedCooperativez5cutlass::epilogue::PtrArrayTmaWarpSpecializedPingpong_epi_nosmem_epi_tma_tma_1sm_tma_2smc                   6    e Zd Z e            Z e            ZdS )EpilogueFunctor3xN)r   r   r   r
   LinearCombination!LinearCombinationBlockScaleFactorr   r   r	   r+  r+    s(        !	(1	%%%r   r+  z,cutlass::epilogue::fusion::LinearCombinationz2cutlass::epilogue::fusion::LinCombBlockScaleFactorc           
          | t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j	        f	v S rg   )
r  r   r   r   r"  r#  r$  r%  r   r   )epilogue_schedule_types    r	   is_tma_epiloguer0    sJ    !)/:22::B?
& 
 
r   c                    |s| S i t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j	        t           j
        t          j        t          j        t          j        t          j        t          j        t          j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j         t           j!        t           j"        t          j#        t          j$        t          j%        t          j&        i}||          S rg   )'r   r   r   r   r   r   r   r   r   r   r   r  r   r  r  r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r"  r$  r#  r%  )schedulegroupedgroup_schedule_maps      r	   to_grouped_scheduler5    s    8:L:r 	ACU  DE 	57I7l	
 	DFX  GK 	ACU  DE 	/1E1h 	:<P<v 	24H4f 	57I7l 	57I7l 	9;M;t 	9;M;t 	9;M;t  	9;M;t!" 	=?Q?|#$ 	=?Q?|%& 	>@R@~'( 	>@R@~24H4f24H4f- 2 h''r   c                   J    e Zd Z e            Z e            Z e            ZdS )TileSchedulerTypeN)r   r   r   r
   Default
PersistentStreamKr   r   r	   r7  r7  =  s/        ikkGJikkGGGr   r7  z"cutlass::gemm::PersistentSchedulerzcutlass::gemm::StreamKScheduler	_stream_kc                   6    e Zd Z e            Z e            ZdS )SideModeN)r   r   r   r
   LeftRightr   r   r	   r=  r=  U  s$        9;;DIKKEEEr   r=  zcutlass::SideMode::kLeftzcutlass::SideMode::kRightlsrsc                   6    e Zd Z e            Z e            ZdS )FillModeN)r   r   r   r
   LowerUpperr   r   r	   rC  rC  g  s$        IKKEIKKEEEr   rC  zcutlass::FillMode::kLowerzcutlass::FillMode::kUpperluc                   6    e Zd Z e            Z e            ZdS )DiagTypeN)r   r   r   r
   NonUnitUnitr   r   r	   rI  rI  y  s$        ikkG9;;DDDr   rI  zcutlass::DiagType::kNonUnitzcutlass::DiagType::kUnitnuunc                   r    e Zd Z e            Z e            Z e            Z e            Z e            ZdS )OpcodeClassN)	r   r   r   r
   SimtTensorOpWmmaTensorOpSparseTensorOpBlockScaledTensorOpr   r   r	   rO  rO    sH        9;;Dy{{H9;;LY[[N#)++r   rO  simttensoropwmma_tensorop
sptensorop
bstensoropzcutlass::arch::OpClassSimtzcutlass::arch::OpClassTensorOpz"cutlass::arch::OpClassWmmaTensorOpz$cutlass::arch::OpClassSparseTensorOpz)cutlass::arch::OpClassBlockScaledTensorOpc                       e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
dS )OperationKindN)r   r   r   r
   GemmRankKRank2KTrmmSymmConv2dConv3dr   r   r	   r[  r[    s[        9;;DIKKEY[[F9;;D9;;DY[[FY[[FFFr   r[  gemmrank_krank_2ktrmmsymmconv2dconv3dc                   "    e Zd Z e            ZdS )TargetN)r   r   r   r
   r   r   r   r	   rk  rk    r   r   rk  maxwellpascalvoltaturingampereadahopper)2   <   =   F   K   P   Y   Z   `      c      )rv  H   rw  rx  V   W   ry  rz  c                     | }d}|rCd}|                                 D ]*\  }}d|z  }t          j        |||          }||k    rd}|}+|C|S )NTFz\$\{%s\})itemsresub)templatevaluestextchangedkeyvalueregexnewtexts           r	   SubstituteTemplater    sw    DG
  ,,.. 	 	JC!C'EfUE400G$DD   Kr   c                      e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
 e            Z e            Z e            Z e            Z e            Z e            ZdS )rq   N)r   r   r   r
   r\  Sparse	UniversalUniversal3xSparseUniversal3xPlanarComplexPlanarComplexArrayGroupedrr   r|   rs   rx   ry   r   r   r	   rq   rq     s        9;;DY[[F	I)++K!	IKKM"ikkG&Y[["$-IKK!$9;;"+)++r   rq   spgemmgemm_planar_complexgemm_planar_complex_arraygemm_groupedc                   "    e Zd Z e            ZdS )	RankKKindNr   r   r   r
   r  r   r   r	   r  r            	IIIr   r  c                   "    e Zd Z e            ZdS )TrmmKindNr  r   r   r	   r  r    r  r   r  c                   "    e Zd Z e            ZdS )SymmKindNr  r   r   r	   r  r  %  r  r   r  c                   6    e Zd Z e            Z e            ZdS )EpilogueFunctorN)r   r   r   r
   r,  LinearCombinationClampr   r   r	   r  r  .  s(        !	&Y[[r   r  z,cutlass::epilogue::thread::LinearCombinationz1cutlass::epilogue::thread::LinearCombinationClampc                   J    e Zd Z e            Z e            Z e            ZdS )MixedInputModeN)r   r   r   r
   ConvertOnly	ScaleOnlyScaleWithZeroPointr   r   r	   r  r  ;  s2        )++K	I"r   r  c                       e Zd Z e            Z e            Z e            Z e            Z e            Z e            Z	 e            Z
 e            Z e            ZdS )SwizzlingFunctorN)r   r   r   r
   	Identity1	Identity2	Identity4	Identity8
HorizontalStridedDgradIdentity1StridedDgradIdentity4StridedDgradHorizontalr:  r   r   r	   r  r  B  st        	I	I	I	IJ%IKK%IKK&Y[[ikkGGGr   r  z=cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<1>z=cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<2>z=cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<4>z=cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<8>z<cutlass::gemm::threadblock::GemmHorizontalThreadblockSwizzlezEcutlass::conv::threadblock::StridedDgradIdentityThreadblockSwizzle<1>zEcutlass::conv::threadblock::StridedDgradIdentityThreadblockSwizzle<4>zDcutlass::conv::threadblock::StridedDgradHorizontalThreadblockSwizzlez5cutlass::gemm::threadblock::ThreadblockSwizzleStreamKc                   8    e Zd Z e            fZ e            ZdS )GroupScheduleModeN)r   r   r   r
   DeviceHostr   r   r	   r  r  ]  s&        ikk^F9;;DDDr   r  z5cutlass::gemm::kernel::GroupScheduleMode::kDeviceOnlyz9cutlass::gemm::kernel::GroupScheduleMode::kHostPrecomputer  r  c                       e Zd ZdZdZdZdS )ConvKindr   r   rU   N)r   r   r   FpropDgradWgradr   r   r	   r  r  r  s        EEEEEr   r  zcutlass::conv::Operator::kFpropzcutlass::conv::Operator::kDgradzcutlass::conv::Operator::kWgradfpropdgradwgradc                       e Zd ZdZdZdS )ConvModer   r   N)r   r   r   CrossCorrelationConvolutionr   r   r	   r  r    s        KKKr   r  c                   "    e Zd ZdZdZdZdZdZdS )IteratorAlgorithmr   r   rU      rV   N)r   r   r   Analytic	OptimizedFixedChannelsFewChannelsFixedStrideDilationr   r   r	   r  r    s*        HIMKr   r  z+cutlass::conv::IteratorAlgorithm::kAnalyticz,cutlass::conv::IteratorAlgorithm::kOptimizedz0cutlass::conv::IteratorAlgorithm::kFixedChannelsz.cutlass::conv::IteratorAlgorithm::kFewChannelsz6cutlass::conv::IteratorAlgorithm::kFixedStrideDilationanalytic	optimizedfixed_channelsfew_channelsfixed_stride_dilationc                       e Zd ZdZdZdZdS )StrideSupportr   r   rU   N)r   r   r   StridedUnityFixedr   r   r	   r  r    s        GEEEEr   r  z&cutlass::conv::StrideSupport::kStridedz$cutlass::conv::StrideSupport::kUnityz$cutlass::conv::StrideSupport::kFixedunity_stridefixed_stridec                   ^    e Zd Z e            Z e            Z e            Z e            ZdS )	GroupModeN)r   r   r   r
   	NoneGroupSingleGroupMultipleGroup	Depthwiser   r   r	   r  r    s:        	I)++KIKKM	IIIr   r  zcutlass::conv::GroupMode::kNonez&cutlass::conv::GroupMode::kSingleGroupz(cutlass::conv::GroupMode::kMultipleGroupz$cutlass::conv::GroupMode::kDepthwisesingle_groupmultiple_group	depthwise)r   r   r   c                   $    e Zd Zej        dfdZdS )MathInstructionNc                 h    || _         || _        || _        || _        || _        || _        || _        d S rg   )instruction_shape	element_a	element_belement_accumulatoropcode_classmath_operationelement_scale_factor)selfr  r  r  r  r  r  r  s           r	   __init__zMathInstruction.__init__  s@     "3""#6 (,$8!!!r   )r   r   r   r   r   r  r   r   r	   r  r    s3         %1!9 9 9 9 9 9r   r  c                        e Zd Z	 	 ddZd ZdS )TileDescriptionr   r   r   Nc	                     || _         || _        || _        || _        || _        || _        || _        || _        || _        d S rg   )	threadblock_shape
tile_shapestages
warp_countmath_instructionminimum_compute_capabilitymaximum_compute_capabilitycluster_shapeexplicit_vector_sizes)	r  r  r  r  r  min_computemax_computer  r  s	            r	   r  zTileDescription.__init__  sO     "3+$ 0*5'*5'*%:"""r   c           	      <   | j         dk    rcd                    | j        d         | j        d         | j        d         | j        d         | j        d         | j        d         | j                  S d| j        d         | j        d         | j        d         | j        fz  S )Nrz  z${tbm}x{tbn}x{tbk}_{cm}x{cn}x{ck}_{s}r   r   rU   )tbmtbntbkcmcnckrG   z%dx%d_%dx%d)r  formatr  r  r  )r  s    r	   procedural_namezTileDescription.procedural_name  s    *b009@@*1-*1-*1-%a(%a(%a(+ A    !&q)&q)&q)	$  r   )r  Nr   r   r   r  r  r   r   r	   r  r    s?          "; ; ; ;*    r   r  c                       e Zd Zd Zd ZdS ).Direct2dConvFixedStrideDilationTileDescriptionc
                     |d         |d         z  |d         z  |d         |d         |d         z  g| _         || _        || _        || _        || _        || _        || _        || _        || _        |	| _	        d S )Nr   r   rU   r  )
r  threadblock_output_shapefilter_shaper  r  stridedilationr  r  r  )
r  r
  r  r  r  r  r  r  r  r  s
             r	   r  z7Direct2dConvFixedStrideDilationTileDescription.__init__  s     %Q'&q)*&q)* %Q'Ol1o-"
 )A%($  0*5'*5'''r   c                    d| j         d         | j         d         | j         d         | j        d         | j        d         | j        d         | j        d         | j        | j        d         | j        d         f
z  }| j        ddgk    rE| j        ddgk    r8|d| j        d         | j        d         | j        d         | j        d         fz  z  }|S )Nz#%dx%dx%d_%dx%dx%dx%d_%d_filter%dx%dr   r   rU   r  z_stride%dx%d_dilation%dx%d)r  r
  r  r  r  r  )r  str_names     r	   r  z>Direct2dConvFixedStrideDilationTileDescription.procedural_name5  s    8"1%"1%"1%)!,)!,)!,)!,Ka a <
 
 ;2r(""t}R'@'@4AAa a 	8  H r   Nr  r   r   r	   r  r    s2        6 6 6:    r   r  c                   $    e Zd Zdej        fdZdS )TensorDescriptionr   c                 >    || _         || _        || _        || _        d S rg   )elementlayout	alignmentcomplex_transform)r  r  r  r  r  s        r	   r  zTensorDescription.__init__O  s'     "!2r   Nr   r   r   rb   rc   r  r   r   r	   r  r  N  s1        )*>N>S3 3 3 3 3 3r   r  c                   0    e Zd Zdej        ej        fdZdS )SymmetricTensorDescriptionr   c                 Z    || _         || _        || _        || _        || _        || _        d S rg   )r  r  	fill_moder  r  	side_mode)r  r  r  r  r  r  r  s          r	   r  z#SymmetricTensorDescription.__init__Z  s3     ""!2"r   N)r   r   r   rb   rc   r=  r>  r  r   r   r	   r  r  Y  s8         */-# # # # # #r   r  c                   $    e Zd Zdej        fdZdS )TriangularTensorDescriptionr   c                 h    || _         || _        || _        || _        || _        || _        || _        d S rg   )r  r  r  r  	diag_typer  r  )r  r  r  r  r  r!  r  r  s           r	   r  z$TriangularTensorDescription.__init__m  s<     """"!2r   Nr  r   r   r	   r  r  l  s3         */3 3 3 3 3 3r   r  c                    | j         j        }| j         j        }| j        t          j        k    r| j        t          j        k    rt          | j
        j                 dk    rd}n t          | j
        j                 dk    rd}nd}t          | j
        j                 |d         z  |d         dz  z  dz  t          | j        j                 |d         z  |d         z  dz  z   |d         |d         dz  z  |z  z   }nt          | j
        j                 }t          | j
        j                 }|                                 rt          | j        j                 }||d         z  |d         z  dz  ||d         z  |d         z  dz  z   }||z  }|dz	  S )NrY   rU   rV   rW   r   r   
   )tile_descriptionr  r  operation_kindr[  r\  ru   rq   r  r   Ar  Bis_mixed_input)	operation	cta_shaper  elements_per_8b_mdsmem_per_stagedata_type_size_adata_type_size_b
smem_usages           r	   CalculateSmemUsager0    s   *<I'.F 	 M$6668?22 	+,22!")+-.!33!"!" ,-	!<	!PQ@QRVWW9;./)A,>1MQRRSlila/04FFG 	 (	(;<'	(;<##%% 	A+IK,?@ y|+il:a?1-	!<AB 	
  &(Jr   c                   "    e Zd ZdZdZdZdZdZdS )GemmUniversalModez2
    Types corresponding to GemmUniversalMode
    r   r   rU   r  N)r   r   r   __doc__r\  GemmSplitKParallelBatchedArrayr   r   r	   r2  r2    s/          DGEEEr   r2  c                       e Zd ZdZdZdZdZdS )
SplitKModez+
    Types corresponding to SplitKMode
    r   r   rU   N)r   r   r   r3  
NoneSplitKSerialParallelr   r   r	   r8  r8    s)          JFHHHr   r8  (?  r3  enumr  r   r
   ImportErrorr   intEnumr   r   GeneratorTargetNamesr   r#   r%   r&   r/   r1   r3   r6   r8   r'   r(   r)   ShortDataTypeNamesr   r   r   r   r   r   r   r   r   r    r!   r"   r$   r+   r*   r,   r-   r.   r0   r2   r4   r5   r7   r?   r@   rA   rB   rC   rD   r9   r:   r;   r<   r=   r>   DataTypeNamesDataTypeTagr   r^   r_   r`   BlasModeTagrb   rc   rd   ComplexTransformTagComplexTransformTag3xrm   ro   rv   rz   r}   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   MathOperationTagr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   	LayoutTagTransposedLayoutShortLayoutTypeNamesShortComplexLayoutNamesr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   KernelScheduleTagKernelScheduleSuffixesr  r  r  r  r  r  r   r!  r"  r#  r$  r%  EpilogueScheduleTagEpilogueScheduleSuffixesr+  r,  r-  EpilogueFunctor3xTagr0  r5  r7  r8  r9  r:  TileSchedulerTagTileSchedulerSuffixesr=  r>  r?  SideModeTagShortSideModeNamesrC  rD  rE  FillModeTagShortFillModeNamesrI  rJ  rK  DiagTypeTagShortDiagTypeNamesrO  rP  rQ  rR  rS  rT  OpcodeClassNamesOpcodeClassTagr[  r\  r]  r^  r_  r`  ra  rb  OperationKindNamesrk  ArchitectureNamesSharedMemPerCCr  rq   r  r  r  r  r  r  r  rr   r|   rs   rx   ry   GemmKindNamesr  RankKKindNamesr  TrmmKindNamesr  SymmKindNamesr  r  EpilogueFunctorTagr  r  r  r  r  r  r  r  r  r  SwizzlingFunctorTagr  r  r  GroupScheduleModeTagShortGroupScheduleModeNamesIntEnumr  r  r  r  ConvKindTagConvKindNamesr  r  r  r  r  r  r  IteratorAlgorithmTagIteratorAlgorithmNamesr  r  r  r  StrideSupportTagStrideSupportNamesr  r  r  r  r  GroupModeTagGroupModeNamesDynamicClusterShaper  r  r  r  r  r  r0  r2  r8  r   r   r	   <module>rp     s.!  B   					&&&&&&&   "#s          di   
 (/; / / / / /ty / / /h L#M6M6L#L#L#M3M3KKK /M6/K/ K/ K	/
 K/ L%/ L%/ L%/ K/ K/ K/ L%/ L%/ L%/ M6/  M6!/" K#/ /$ K%/& K'/( M6)/* M6+/, M6-/. NG//0 NG1/2 L%3/4 M65/6 L%7/8 M69/: L%;/< M6=/> NG?/@ M6A/B NGC/D M6E/ /F L%L%L%M6M6M6L%L%L%M6M6M6]/ /b/M6/K$/ K$/ K$	/
 K/ L*/ L*/ L*/ K#/ K#/ K/ L)/ L)/ L)/ M*/  M*!/" K8#/ /$ K8%/& K8'/( M*)/* M*+/, M*-/. N,//0 N,1/2 L#3/4 M(5/6 L'7/8 M(9/: L(;/< M6=/> N;?/@ M,A/B N;C/D M-E/ /F L7L7L6M8M8M8L6L6L5M7M7M7]/ /b/M1/K/ K/ K	/
 K/ L"/ L"/ L"/ K/ K/ K/ L"/ L"/ L"/ M1/  M1!/" K#/ /$ K%/& K'/( M1)/* M1+/, M1-/. NA//0 NA1/2 L"3/4 M25/6 L"7/8 M29/: L";/< M2=/> NB?/@ M2A/B NBC/D M3E/ /F L!L!L"M2M2M3L!L!L"M2M2M3]/ /h    ty    77    ty    =B  +,  \8=!\8=!\8=! A A A          . . .    	   * * * * *DI * * *"  >')O13a66(*P')N')N/1]&(M/1^)+R &       0:4 &(T #%N	
 ')V $&P ')V $&P 6 8 : 8 : B B  B!" B#$ 68:) 	2 J//&
(G#Z%G')I$j&I')I$j&I:0
 C& ' '	
  #T $e $e % 6 G 6 G z z  z!" z#$ %6G)  2 -23S-23S*/0#*/0#	 D; D; D; D; D; D; D; D;N9#%T9!#D9 -/\9 57l	9
 8:r9 69 )+T9 13d9 46j9 57l9 @  CC9 =?|9 57l9 =  @G9 13d9  13d!9" 9;t#9 9$ 9;t%9& 9;t'9( 9;t)9* 79p+9, 79p-9. <>z/90 <>z192 9;t394 9;t596 :<v798 :<v99: B  EG;9< B  EG=9> 57l?9@ 57lA9B 57lC9D 57lE9 9 9F <>zG9H H  KSI9J 9;tK9L E  HMM9N E  HWO9P D  GKQ9R D  GKS9T =?|U9V =?|W9X =?|Y9Z =?|[9\ A  DE]9^ A  DE_9` ACya9b >@~c9d =?qe9f :<vg9 9h =?q:<vEGwB  EG?  BAq9 9 x9#R9!:9 -/I9 57Z	9
 8:`9 ,9 )+=9 13N9 46T9 57W9 @Bn9 =?h9 57I9 =?]9 169  16!9" 96#9 9$ 96%9& 96'9( 96)9* 7+9, 7-9. <f/90 <f192 98394 98596 :F798 :F99: BF;9< BF=9> 5}?9@ 5}A9B 5}C9D 5}E9 9 9F <>\G9H HJvI9J 9;VK9L EGpM9N EGeO9P DfQ9R DfS9T =}U9V =}W9X =}Y9Z =}[9\ A=]9^ A=_9` ACSa9b >c9d =?Te9f :<Ng9 9h =?T:<NEtBDT?q9 9 x8 8 8 8 849 8 8 8(%'\+-P .0Z 68j	
 13` 13` 9;p 9;p +-T 68j .0Z .0Z 68j 68j >@z  ;=t! (%r+R . 6	
 1= 1= 9= 9= +Z 6
 . .
 6
 6
 >
  ;Z! (4 4 4 4 4	 4 4 4 ')W79m   ( ( (@    	    v "F@  r "{     ty    M-N/ mT8>4@     ty    N/N/ nc8>3?     ty    3M- &hmTB & & & & &$) & & & f*o#\  2:B F#%P    DI    )((     TY    		 	  	



	 	  $. . . . .ty . . .$ M6OX&1!<n#V*N!6(.$    	   
 %x0    ty   
 #V,    ty   
 #V,) ) ) ) )di ) ) ) %'U*,_ % % % % %TY % % %	 	 	 	 	ty 	 	 	  _ _ _ _!_*,s*,s+-sU
     	    UW  hF     t|    N5N5N5 NGNGNG    t|       	     M!O#%W!#S)+c  
#%5!>)+B     DI    C??  2     	    :CG?	 >-	  ii 9 9 9 9 9 9 9 9*' ' ' ' ' ' ' 'T3 3 3 3 3 3 3 3n3 3 3 3 3 3 3 3# # # # # # # #&3 3 3 3 3 3 3 3*" " "J                s    ##