+
    "Ý³ja  ã                   ó†  € ^ RI t ^ RIt^ RIt^ RIt^ RIHt ^ RIHt ^ RIH	t	H
t
 ^ RIHt ^ RIHt ^ RIt^ RIt]! ]4      P%                  4       P&                  ^,          t]R,          R,          t]P,                  P/                  R	4      t]e
   ]! ]4      tM	]R
,          t]R,          tR R ltR R lt]! ]4       ]! ]4      '       d
   ]! ]4       M]! R] R24      h^ RIHtHtH t H!t!H"t"H#t#H$t$H%t%H&t&H't'H(t(H)t)H*t* Rt+Rt,Rt-]!P\                  t/] P\                  t0]	! RR7       ! R R4      4       t1]	! RR7       ! R R4      4       t2]	! RR7       ! R R4      4       t3]	 ! R R4      4       t4R  R! lt5R" R# lt6R$ R% lt7R& R' lt8R( t9R) R* lt:R+ R, lt;R- t<R.]Pz                  R/]'R0])R1](R2]*R3]R4]"R5]:R6]6R7]7R8];R9]<R:]8/t>R; t?]?! ]+4      t@]?! ],4      tA]?! ]-4      tBRGR< R= lltCR> R? ltDRHR@ RA lltE]FRB8X  dL   ] PŽ                  ! RCRD7      tH]HP“                  RE]RRF7       ]HP•                  4       tK]E! ]KP˜                  4       R# R# )Ié    N)ÚIterable)Údeepcopy)Ú	dataclassÚfields)Úreduce©ÚPathÚgptqmodel_extÚcutlass_extensionsÚGPTQMODEL_CUTLASS_DIRÚcutlassÚpythonc                ó(   € V ^8„  d   QhR\         RR/# )é   ÚpathÚreturnNr   )Úformats   "Ú_/home/luca/screen-checking/.venv/lib/python3.14/site-packages/gptqmodel_ext/machete/generate.pyÚ__annotate__r      s   € ÷ "ñ "œDð " Tñ "ó    c                 óÄ   € \        V 4      pV\        P                  9   d    \        P                  P                  V4       \        P                  P	                  ^ V4       R# )r   N)ÚstrÚsysr   ÚremoveÚinsert)r   Ú	path_texts   & r   Ú_prepend_sys_pathr      s9   € Ü�D“	€IØ”C—H‘HÔÜ�‰�‰˜	Ô"Ü‡H�H‡O�O�A�yÖ!r   c                ó0   € V ^8„  d   QhR\         R\        /# )r   r   r   )r	   Úbool)r   s   "r   r   r   %   s   € ÷ ñ ¬4ð ´Dñ r   c                 ó|   € V R ,          P                  4       ;'       g    V R,          R,          P                  4       # )zcutlass_library.pyÚcutlass_libraryz__init__.py)Úis_file)r   s   &r   Ú _cutlass_python_bindings_presentr#   %   s<   € à	Ð$Õ	$×-Ñ-Ó/÷ 	@ð 	@ØÐ$Õ$ }Õ4×=Ñ=Ó?ðr   z)CUTLASS python bindings not found under `z9`. Set GPTQMODEL_CUTLASS_DIR to a valid CUTLASS checkout.)ÚDataTypeÚEpilogueScheduleTagÚEpilogueScheduleTypeÚMixedInputKernelScheduleTypeÚTileSchedulerTagÚTileSchedulerTypeÚVLLMDataTypeÚVLLMDataTypeNamesÚVLLMDataTypeSizeÚVLLMDataTypeTagÚVLLMDataTypeTorchDataTypeTagÚVLLMDataTypeVLLMScalarTypeTagÚVLLMKernelScheduleTagaj  
#include "machete_mm_launcher.cuh"

namespace machete {

{% for impl_config in impl_configs %}
{% set type_sig = gen_type_sig(impl_config.types) -%}
{% for s in impl_config.schedules %}
extern torch::Tensor impl_{{type_sig}}_sch_{{gen_sch_sig(s)}}(MMArgs);
{%- endfor %}

torch::Tensor mm_dispatch_{{type_sig}}(MMArgs args) {
  [[maybe_unused]] auto M = args.A.size(0);
  [[maybe_unused]] auto N = args.B.size(1);
  [[maybe_unused]] auto K = args.A.size(1);
    
  if (!args.maybe_schedule) {
    {%- for cond, s in impl_config.heuristic %}
    {%if cond is not none%}if ({{cond}})
    {%- else %}else
    {%- endif %}
        return impl_{{type_sig}}_sch_{{ gen_sch_sig(s) }}(args);{% endfor %}
  }

  {%- for s in impl_config.schedules %}
  if (*args.maybe_schedule == "{{ gen_sch_sig(s) }}")
    return impl_{{type_sig}}_sch_{{ gen_sch_sig(s) }}(args);
  {%- endfor %}
  TORCH_CHECK_NOT_IMPLEMENTED(false, "machete_gemm(..) is not implemented for "
                                     "schedule = ", *args.maybe_schedule);
}
{%- endfor %}


static inline std::optional<at::ScalarType> maybe_scalartype(
    std::optional<at::Tensor> const& t) {
    if (!t) {
      return std::nullopt;
    } else {
      return t->scalar_type();
    };
}

torch::Tensor mm_dispatch(MMArgs args) {
  auto out_type = args.maybe_out_type.value_or(args.A.scalar_type());
  auto a_type = args.A.scalar_type();
  auto maybe_g_scales_type = maybe_scalartype(args.maybe_group_scales);
  auto maybe_g_zeros_type = maybe_scalartype(args.maybe_group_zeros);
  auto maybe_ch_scales_type = maybe_scalartype(args.maybe_channel_scales);
  auto maybe_tok_scales_type = maybe_scalartype(args.maybe_token_scales);

  {% for impl_config in impl_configs %}
  {% set t = impl_config.types -%}
  {% set type_sig = gen_type_sig(t) -%}
  if (args.b_type == {{VLLMScalarTypeTag[t.b]}}
      && a_type == {{TorchTypeTag[t.a]}}
      && out_type == {{TorchTypeTag[t.out]}}
      && {%if t.b_group_scale != void -%}
      maybe_g_scales_type == {{TorchTypeTag[t.b_group_scale]}}
      {%- else %}!maybe_g_scales_type{%endif%}
      && {%if t.b_group_zeropoint != void -%}
      maybe_g_zeros_type == {{TorchTypeTag[t.b_group_zeropoint]}}
      {%- else %}!maybe_g_zeros_type{%endif%}
      && {%if t.b_channel_scale != void -%}
      maybe_ch_scales_type == {{TorchTypeTag[t.b_channel_scale]}}
      {%- else %}!maybe_ch_scales_type{%endif%}
      && {%if t.a_token_scale != void -%}
      maybe_tok_scales_type == {{TorchTypeTag[t.a_token_scale]}}
      {%- else %}!maybe_tok_scales_type{%endif%}
  ) {
      return mm_dispatch_{{type_sig}}(args);
  }
  {%- endfor %}
  
  TORCH_CHECK_NOT_IMPLEMENTED(
    false, "machete_mm(..) is not implemented for "
    "a_type=", args.A.scalar_type(),
    ", b_type=", args.b_type.str(),
    ", out_type=", out_type,
    ", with_group_scale_type=", maybe_g_scales_type
        ? toString(*maybe_g_scales_type) : "None",
    ", with_group_zeropoint_type=", maybe_g_zeros_type
        ? toString(*maybe_g_zeros_type) : "None",
    ", with_channel_scale_type=", maybe_ch_scales_type
        ? toString(*maybe_ch_scales_type) : "None",
    ", with_token_scale_type=", maybe_tok_scales_type
        ? toString(*maybe_tok_scales_type) : "None",
    "; implemented types are: \n",
    {%- for impl_config in impl_configs %}
    {% set t = impl_config.types -%}
    "\t{{gen_type_option_name(t)}}\n",
    {%- endfor %}
    "");
}

std::vector<std::string> supported_schedules_dispatch(
    SupportedSchedulesArgs args) {
    auto out_type = args.maybe_out_type.value_or(args.a_type);
    
    {% for impl_config in impl_configs %}
    {% set t = impl_config.types -%}
    {% set schs = impl_config.schedules -%}
    if (args.b_type == {{VLLMScalarTypeTag[t.b]}}
        && args.a_type == {{TorchTypeTag[t.a]}}
        && out_type == {{TorchTypeTag[t.out]}}
        && {%if t.b_group_scale != void -%}
        args.maybe_group_scales_type == {{TorchTypeTag[t.b_group_scale]}}
        {%- else %}!args.maybe_group_scales_type{%endif%}
        && {%if t.b_group_zeropoint != void-%}
        args.maybe_group_zeros_type == {{TorchTypeTag[t.b_group_zeropoint]}}
        {%- else %}!args.maybe_group_zeros_type{%endif%}
    ) {
        return {
            {%- for s in impl_config.schedules %}
            "{{gen_sch_sig(s)}}"{% if not loop.last %},{% endif %}
            {%- endfor %}
        };
    }
    {%- endfor %}
    
    return {};
};

}; // namespace machete
a:  
#include "machete_mm_launcher.cuh"

namespace machete {
    
{% for sch in unique_schedules(impl_configs) %}
{% set sch_sig = gen_sch_sig(sch) -%}
struct sch_{{sch_sig}} {
  using TileShapeNM = Shape<{{
      to_cute_constant(sch.tile_shape_mn)|join(', ')}}>;
  using ClusterShape = Shape<{{
      to_cute_constant(sch.cluster_shape_mnk)|join(', ')}}>;
  // TODO: Reimplement
  // using KernelSchedule   = {{KernelScheduleTag[sch.kernel_schedule]}};
  using EpilogueSchedule = {{EpilogueScheduleTag[sch.epilogue_schedule]}};
  using TileScheduler    = {{TileSchedulerTag[sch.tile_scheduler]}};
  using EpilogueTileType = cutlass::epilogue::collective::EpilogueTileAuto;
};
{% endfor %}
    
{% for impl_config in impl_configs %}
{% set t = impl_config.types -%}
{% set schs = impl_config.schedules -%}
{% set type_sig = gen_type_sig(t) -%}

template<typename Sch>
using Kernel_{{type_sig}} = MacheteKernelTemplate<
  {{DataTypeTag[t.a]}},  // ElementA
  {{DataTypeTag[t.b]}},  // ElementB
  {{DataTypeTag[t.out]}},  // ElementD
  {{DataTypeTag[t.accumulator]}}, // Accumulator
  {{DataTypeTag[t.b_group_scale]}}, // GroupScaleT
  {{DataTypeTag[t.b_group_zeropoint]}}, // GroupZeroT
  {{DataTypeTag[t.b_channel_scale]}}, // ChannelScaleT
  {{DataTypeTag[t.a_token_scale]}}, // TokenScaleT
  cutlass::gemm::KernelTmaWarpSpecializedCooperative,
  Sch>;

{% for sch in schs %}
{% set sch_sig = gen_sch_sig(sch) -%}
torch::Tensor 
impl_{{type_sig}}_sch_{{sch_sig}}(MMArgs args) {
  return run_impl<Kernel_{{type_sig}}<sch_{{sch_sig}}>>(args);
}
{%- endfor %}
{%- endfor %}

}; // namespace machete
ak  
#include "machete_prepack_launcher.cuh"

namespace machete {

torch::Tensor prepack_B_dispatch(PrepackBArgs args) {
  auto convert_type = args.maybe_group_scales_type.value_or(args.a_type);
  {%- for t in types %}
  {% set b_type = unsigned_type_with_bitwidth(t.b_num_bits) %}
  if (args.a_type == {{TorchTypeTag[t.a]}}
      && args.b_type.size_bits() == {{t.b_num_bits}} 
      && convert_type == {{TorchTypeTag[t.convert]}}) {
    return prepack_impl<
      PrepackedLayoutBTemplate<
        {{DataTypeTag[t.a]}}, // ElementA
        {{DataTypeTag[b_type]}}, // ElementB
        {{DataTypeTag[t.convert]}}, // ElementConvert
        {{DataTypeTag[t.accumulator]}}, // Accumulator
        cutlass::layout::ColumnMajor,
        cutlass::gemm::KernelTmaWarpSpecializedCooperative>
    >(args.B); 
  }
  {%- endfor %}
  
  TORCH_CHECK_NOT_IMPLEMENTED(false, 
    "prepack_B_dispatch(..) is not implemented for "
    "atype = ", args.a_type,
    ", b_type = ", args.b_type.str(),
    ", with_group_scales_type= ", args.maybe_group_scales_type ? 
        toString(*args.maybe_group_scales_type) : "None");
}

}; // namespace machete
T)Úfrozenc                   ó,   a € ] tR tRt o V 3R ltRtV tR# )ÚScheduleConfigi   c                ó†   <€ V ^8„  d   Qh/ S[ S[S[3,          ;R&   S[ S[S[S[3,          ;R&   S[;R&   S[;R&   S[;R&   # )r   Útile_shape_mnÚcluster_shape_mnkÚkernel_scheduleÚepilogue_scheduleÚtile_scheduler)ÚtupleÚintr'   r&   r)   )r   Ú__classdict__s   "€r   r   ÚScheduleConfig.__annotate__   sW   ø‡ ‚ á™™c˜•?Ñ"ñ ñ ™S¡#¡s˜]Õ+Ñ+ñ ñ 2Ñ1ñ	 ñ
 ,Ñ+ñ ñ &Ñ%ò r   © N©Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__annotate_func__Ú__static_attributes__Ú__classdictcell__©r<   s   @r   r3   r3      ó   ø‡ ‡ „ r   r3   c                   ó,   a € ] tR tRt o V 3R ltRtV tR# )Ú
TypeConfigi)  c                óŠ   <€ V ^8„  d   Qh/ S[ ;R&   S[ S[,          ;R&   S[ ;R&   S[ ;R&   S[ ;R&   S[ ;R&   S[ ;R&   S[ ;R&   # )	r   ÚaÚbÚb_group_scaleÚb_group_zeropointÚb_channel_scaleÚa_token_scaleÚoutÚaccumulator)r$   r*   )r   r<   s   "€r   r   ÚTypeConfig.__annotate__)  sn   ø‡ ‚ á�Kñ ñ ‘,ÕÑñ ñ Ññ	 ñ
  Ññ ñ Ññ ñ Ññ ñ 
�Mñ ñ Ñò r   r>   Nr?   rG   s   @r   rJ   rJ   )  rH   r   rJ   c                   ó,   a € ] tR tRt o V 3R ltRtV tR# )ÚPrepackTypeConfigi5  c                óJ   <€ V ^8„  d   Qh/ S[ ;R&   S[;R&   S[ ;R&   S[ ;R&   # )r   rL   Ú
b_num_bitsÚconvertrS   )r$   r;   )r   r<   s   "€r   r   ÚPrepackTypeConfig.__annotate__5  s5   ø‡ ‚ á�Kñ ñ �Oñ ñ Ññ	 ñ
 Ñò r   r>   Nr?   rG   s   @r   rV   rV   5  rH   r   rV   c                   ó,   a € ] tR tRt o V 3R ltRtV tR# )Ú
ImplConfigi=  c                ó‚   <€ V ^8„  d   Qh/ S[ ;R&   S[S[,          ;R&   S[S[S[R,          S[3,          ,          ;R&   # )r   ÚtypesÚ	schedulesNÚ	heuristic)rJ   Úlistr3   r:   r   )r   r<   s   "€r   r   ÚImplConfig.__annotate__=  sB   ø‡ ‚ áÑñ ñ ‘NÕ#Ñ#ñ ñ ‘E™# �*¡nÐ4Õ5Õ6Ñ6ò	 r   r>   Nr?   rG   s   @r   r\   r\   =  rH   r   r\   c                ó0   € V ^8„  d   QhR\         R\        /# ©r   Úschedule_configr   ©r3   r   )r   s   "r   r   r   D  s   € ÷ ñ ¤nð ¼ñ r   c                 ó$  € V P                   ^ ,           RV P                   ^,           2pV P                  ^ ,           RV P                  ^,           2,           RV P                  ^,           2,           p\        V P                  ,          P	                  R4      R,          p\
        V P                  ,          P	                  R4      R,          p\        V P                  ,          P	                  R4      R,          pV RV RV 2RV RV 2,           # )r   Úxz::Ú_éÿÿÿÿ)	r5   r6   r0   r7   Úsplitr%   r8   r(   r9   )re   Ú
tile_shapeÚcluster_shaper7   r8   r9   s   &     r   Úgenerate_sch_sigrn   D  s#  € à×(Ñ(¨Õ+Ð
,¨A¨o×.KÑ.KÈAÕ.NÐ-OÐPð ð ×,Ñ,¨QÕ/Ð
0Øˆo×/Ñ/°Õ2Ð3Ð
4õ	5àˆo×/Ñ/°Õ2Ð3Ð
4õ	5ð ô
 ,¨O×,KÑ,KÕL×RÑRØóàõ
€Oô ,¨O×,MÑ,MÕN×TÑTØóàõ
Ðô & o×&DÑ&DÕE×KÑKÈDÓQÐRTÕU€Nð ˆ,�a˜� a¨Ð'8Ð9ØÐÐ   .Ð!1Ð
2õ	3ðr   c                ó0   € V ^8„  d   QhR\         R\        /# rd   rf   )r   s   "r   r   r   \  s   € ÷ 
ñ 
¬Nð 
¼sñ 
r   c                 ó€   € R RRRRR/p\        V 4      pVP                  4        F  w  r4VP                  W44      pK  	  V# )Ú#KernelTmaWarpSpecializedCooperativeÚTmaMI_ÚTmaWarpSpecializedCooperative_ÚTmaCoop_ÚStreamKSchedulerÚstreamK)rn   ÚitemsÚreplace)re   Úkernel_terse_names_replaceÚsch_sigÚorigÚterses   &    r   Úgenerate_terse_sch_sigr}   \  sM   € à-¨xØ(¨*Ø˜Ið"Ðô ˜Ó/€GØ1×7Ñ7Ö9‰ˆØ—/‘/ $Ó.Šñ :à€Nr   c                ó$   € V ^8„  d   QhR\         /# ©r   Úkernel_types©rJ   )r   s   "r   r   r   j  s   € ÷ ñ ¬*ñ r   c                 ó¶   € \        R P                  \        \        4       Uu. uF#  p\        \        WP                  4      ,          NK%  	  up4      4      # u upi )Ú )r   Újoinr   rJ   r+   ÚgetattrÚname©r€   Úfields   & r   Úgenerate_type_signaturer‰   j  sS   € ÜØ
�‰ô $¤JÔ/óá/�Eô "¤'¨,¿
¹
Ó"C×DÐDÙ/ñó	
óð ùòs   ¢)Ac                ó$   € V ^8„  d   QhR\         /# r   r�   )r   s   "r   r   r   u  s   € ÷ ñ ¬Jñ r   c                 óú   € R P                  \        \        4       Uu. uFN  pVP                  P	                  RR4      R,            R2\
        \        WP                  4      ,          ,           NKP  	  up4      # u upi )z, Úb_Úwith_Ú_typeÚ=)r„   r   rJ   r†   rx   r+   r…   r‡   s   & r   Úgenerate_type_option_namer�   u  sw   € Ø�9‰9ô  ¤
Ô+ó	
ñ ,�ð �z‰z×!Ñ! $¨Ó0°7Õ:Ð;¸1Ð=Ü¤¨·j±jÓ AÕB÷Cð Cá+ñ	
óð ùò	
s   �AA8c                 ó@   € V ^ 8g  ;'       d    W ^,
          ,          ^ 8H  # )r   r>   )Úns   &r   Úis_power_of_twor“     s   € Ø�‰F×*Ð*˜ !�e�¨Ñ)Ð*r   c                ó:   € V ^8„  d   QhR\         \        ,          /# ©r   Úvalue)ra   r;   )r   s   "r   r   r   ƒ  s   € ÷ 
(ñ 
(œD¤�Iñ 
(r   c                 ó‚   € R  R lp\        V \        4      '       d   V  U u. uF
  q! V 4      NK  	  up # V! V 4      # u up i )c                ó$   € V ^8„  d   QhR\         /# r•   )r;   )r   s   "r   r   Ú&to_cute_constant.<locals>.__annotate__„  s   € ÷ #ñ #¤ñ #r   c                 ó:   € \        V 4      '       d   R V  2# RV  R2# )ri   zInt<Ú>)r“   )r–   s   &r   Ú_to_cute_constantÚ+to_cute_constant.<locals>._to_cute_constant„  s'   € Ü˜5×!Ò!Ø�u�g�;Ðà˜%˜ �?Ð"r   )Ú
isinstancer   )r–   rœ   s   & r   Úto_cute_constantrŸ   ƒ  sA   € õ#ô �%œ×"Ò"Ù6;Ó<±e¨UÐ! %Ö(±eÑ<Ð<á  Ó'Ð'ùò =s   ¡<c                ó:   € V ^8„  d   QhR\         \        ,          /# ©r   Úimpl_configs©ra   r\   )r   s   "r   r   r   �  s   € ÷ ñ ¤4¬
Õ#3ñ r   c           	      óˆ   € \        V  UUu/ uF  qP                   F  q"R bK  	  K  	  uppP                  4       4      # u uppi ©N)ra   r_   Úkeys)r¢   Úimpl_configÚschs   &  r   Úunique_schedulesr©   �  sA   € äá)5ô	
Ù)5˜+×AVÔAV¸#�ŠIÑAV‰C©ò	
ç
‰$‹&óð ùó	
s   ‹>
c           
      ó´   € ^\         P                  ^\         P                  ^\         P                  ^ \         P                  ^@\         P
                  /V ,          # )é   )r$   Úu4Úu8Úu16Úu32Úu64)Únum_bitss   &r   Úunsigned_type_with_bitwidthr²   ™  sC   € à	Œ8�;‰;Ø	Œ8�;‰;Ø
ŒH�L‰LØ
ŒH�L‰LØ
ŒH�L‰Lðð õð r   ÚvoidÚDataTypeTagÚVLLMScalarTypeTagÚTorchTypeTagÚKernelScheduleTagr%   r(   rŸ   Úgen_sch_sigÚgen_type_sigr©   r²   Úgen_type_option_namec                 óp   € \         P                  ! V 4      pVP                  P                  \        4       V# r¥   )Újinja2ÚTemplateÚglobalsÚupdateÚtemplate_globals)Útemplate_strÚtemplates   & r   Úcreate_templaterÃ   ´  s*   € Ü�Š˜|Ó,€HØ×Ñ×ÑÔ,Ô-Ø€Or   c                ó:   € V ^8„  d   QhR\         \        ,          /# r¡   r£   )r   s   "r   r   r   ¿  s   € ÷ Wñ W¤¤jÕ!1ñ Wr   c           
      ó$  € . pVP                  R \        P                  V R7      34       . pV  F½  pVP                  P                  \
        P                  8X  d   VP                  P                  MVP                  P                  pVP                  \        VP                  P                  \        VP                  P                  ,          VVP                  P                  R7      4       K¿  	  R R lp. p\        4       pV F5  p	V! V	4      p
W¨9  g   K  VP                  V	4       VP                  V
4       K7  	  VP                  R\        P                  VR7      34       \        R V ^ 4      p\         P"                  ! W±,          4      p. .p^ p^ p\%        \'        \)        V 4      4      4      pWë8  Ed   WÏ,
          pV^ 8X  d   VP                  . 4       Tp^ pVR
,          p\+        VP,                  4      V8¼  dM   \%        V4      pVP,                  RV Vn        VP,                  VR Vn        VR
,          P                  V4       M(VR
,          P                  V4       VP/                  4        V\+        VR
,          R
,          P,                  4      ,          pV\+        VR
,          R
,          P,                  4      ,          pEK  \1        V4       F7  w  ppVP                  R	V^,            2\2        P                  VR7      34       K9  	  V# )Úmachete_mm_dispatch)r¢   )rL   rX   rY   rS   c                ó$   € V ^8„  d   QhR\         /# )r   Úprepack_type)rV   )r   s   "r   r   Ú$create_sources.<locals>.__annotate__Ù  s   € ÷ Oñ OÔ):ñ Or   c                 óH   € V P                   V P                  V P                  3# r¥   )rL   rX   rY   )rÈ   s   &r   Úprepacked_type_keyÚ*create_sources.<locals>.prepacked_type_keyÙ  s!   € ð —‘ × 7Ñ 7¸×9MÑ9MÐNÐNr   Úmachete_prepack)r^   c                 ó:   € V \        VP                  4      ,           # r¥   )Úlenr_   )rh   Úys   &&r   Ú<lambda>Ú create_sources.<locals>.<lambda>ñ  s   €  A¬¨A¯K©KÓ(8Ö$8r   NÚmachete_mm_impl_partrj   )ÚappendÚmm_dispatch_templateÚrenderr^   rN   r$   r³   rL   rV   r,   rM   rS   ÚsetÚaddÚprepack_dispatch_templater   ÚmathÚceilr   ra   ÚreversedrÏ   r_   ÚpopÚ	enumerateÚmm_impl_template)r¢   Únum_impl_filesÚsourcesÚprepack_typesr§   Úconvert_typerË   Úunique_prepack_typesÚprepack_types_seenrÈ   ÚkeyÚ	num_implsÚnum_impls_per_fileÚfiles_implsÚcurr_num_impls_assignedÚcurr_impl_in_fileÚcurr_impl_configsÚroom_left_in_fileÚcurr_icÚtmp_icÚpartÚ
file_implss   &&                    r   Úcreate_sourcesrò   ¿  sÌ  € Ø€Gà‡N�Nà!Ü ×'Ñ'°\Ð'ÓBð	
ôð €MÛ#ˆð × Ñ ×.Ñ.´(·-±-Ô?ð ×Ñ×Òà×"Ñ"×0Ñ0ð 	ð
 	×ÑÜØ×#Ñ#×%Ñ%Ü+¨K×,=Ñ,=×,?Ñ,?Õ@Ø$Ø'×-Ñ-×9Ñ9ô	ö	
ñ $õOð ÐÜ›ÐÛ%ˆÙ  Ó.ˆØÖ(Ø ×'Ñ'¨Ô5Ø×"Ñ" 3Ö'ñ	 &ð ‡N�NàÜ%×,Ñ,Ø*ð -ó ð	
ôô Ñ8¸,ÈÓJ€IÜŸš 9Õ#=Ó>Ðà+-¨$€KàÐØÐÜ ¤¤h¨|Ó&<Ó!=Ó>Ðà
!Õ
-Ø.ÕBÐØ Ô!Ø×Ñ˜rÔ"Ø 2ÐØ !Ðà# BÕ'ˆÜˆw× Ñ Ó!Ð%6Ô6ä˜gÓ&ˆFØ&×0Ñ0Ð1CÐ2CÐDˆFÔØ '× 1Ñ 1Ð2CÐ2DÐ EˆGÔØ˜�O×"Ñ" 6Õ*à˜�O×"Ñ" 7Ô+Ø×!Ñ!Ô#Ø¤3 {°2¥°rÕ':×'DÑ'DÓ#EÕEÐØœS ¨R¥°Õ!4×!>Ñ!>Ó?Õ?Óä% kÖ2ÑˆˆjØ�‰à& t¨a¥x jÐ1Ü ×'Ñ'°ZÐ'Ó@ðö	
ñ 3ð €Nr   c                ó$   € V ^8„  d   QhR\         /# )r   r   r   )r   s   "r   r   r     s   € ÷ ;ñ ;œTñ ;r   c                  ó†  € \         P                  P                  R 4      p V '       d   \        V 4      P	                  4       pMj\         P                  P                  R4      pV'       d!   \        V4      P	                  4       R,          M"\        P
                  ! 4       R,          R,          pVR,          R,          R,          # )ÚGPTQMODEL_CACHE_DIRÚXDG_CACHE_HOMEÚ	gptqmodelz.cacheÚmacheteÚ	generatedÚmanual)ÚosÚenvironÚgetr	   Ú
expanduserÚhome)Ú
configuredÚ
cache_rootÚxdg_cache_homes      r   Ú_default_output_dirr    sŠ   € Ü—‘—‘Ð 5Ó6€JßÜ˜*Ó%×0Ñ0Ó2‰
äŸ™Ÿ™Ð(8Ó9ˆ÷ ô �Ó ×+Ñ+Ó-°Ö;ä—’“˜xÕ'¨+Õ5ð 	ð
 ˜	Õ! KÕ/°(Õ:Ð:r   c                ór   € V ^8„  d   QhR\         \        P                  \         ,          ,          R,          /# )r   Ú
output_dirN)r   rû   ÚPathLike)r   s   "r   r   r   '  s)   € ÷ A2ñ A2œœrŸ{™{¬3Õ/Õ/°$Õ6ñ A2r   c           
      óø  € \        \        \        \        P                  R 7      p/ RRbRR bRR!bRRbRR bRR"bRR!bRR#bR	RbR
R$bRR%bRR!bRR"bRR&bRR'bRR$bpVP                  4        UUu. uF  w  r4V\        V/ VB 3NK  	  pppR R lp. p\        R \        P                  \        P                  3 4       4      pT\        V\        P                  ! V! V4      4      \        P                  ! V4      4       U	u. uF%  p	\        V	^ ,          V	^,          V	^,          4      NK'  	  up	,          p\        R \        P                   \        P"                  3 4       4      p
T\        V
\        P                  ! V! V4      4      \        P                  ! V4      4       U	u. uF%  p	\        V	^ ,          V	^,          V	^,          4      NK'  	  up	,          pV f   \%        4       pM\'        V 4      P)                  4       pVP+                  4       '       dU   VP-                  4       '       g   \/        RV R24      h\1        VP3                  4       4      '       d   \/        RV R24      hMVP5                  RR7       \7        V4       FS  w  rÍW¼ R2,          pVP9                  RRR7      ;_uu_ 4       pVP;                  V4       RRR4       \=        RV 24       KU  	  R# u uppi u up	i u up	i   + '       g   i     L4; i)()r7   r8   r9   z"M > 256 && K <= 16384 && N <= 4096zM > 256z!M > 128 && K <= 4096 && N <= 4096z!M > 128 && K <= 8192 && N <= 8192zM > 128z M > 64 && K <= 4069 && N <= 4069z M > 64 && K <= 4069 && N <= 8192z!M > 64 && K >= 8192 && N >= 12288zM > 64z M > 32 && K <= 6144 && N <= 6144z"M > 32 && K >= 16384 && N >= 12288zM > 32z!M > 16 && K <= 12288 && N <= 8192zM > 16z
N >= 26624Nc                óF   € V ^8„  d   QhR\         \        \        3,          /# )r   r`   )Údictr   r3   )r   s   "r   r   Úgenerate.<locals>.__annotate__Q  s   € ÷ ñ ¬¬S´.Ð-@Õ(Añ r   c                 óR   € . pV  F  w  r#W19  g   K  VP                  V4       K   	  V# r¥   )rÔ   )r`   r_   ri   re   s   &   r   Úget_unique_schedulesÚ&generate.<locals>.get_unique_schedulesQ  s2   € ð ˆ	Û"+ÑˆAØÖ/Ø× Ñ  Ö1ñ #,ð Ðr   c              3   ó  "  € T Fw  p\         P                  \         P                  3 FP  p\        VVV\         P                  \         P                  \         P                  V\         P
                  R 7      x € KR  	  Ky  	  R# 5i©)rL   rM   rN   rO   rP   rQ   rR   rS   N©r$   Úf16Úbf16rJ   r³   Úf32©Ú.0rM   rL   s   &  r   Ú	<genexpr>Úgenerate.<locals>.<genexpr>]  si   é € ð $ñ :ˆAÜ—,‘,¤§¡Ó.ˆAô 	ØØØÜ&Ÿm™mÜ$ŸM™MÜ"Ÿ-™-ØÜ Ÿ™÷		
ñ 		
ñ /ñ		
ó :ùs   ‚A?Bc              3   óê   "  € T Fi  p\         P                  \         P                  3 FB  p\        VVVV\         P                  \         P                  V\         P
                  R 7      x € KD  	  Kk  	  R# 5ir  r  r  s   &  r   r  r  u  se   é € ð #ñ ,ˆAÜ—,‘,¤§¡Ó.ˆAô 	ØØØØÜ$ŸM™MÜ"Ÿ-™-ØÜ Ÿ™÷		
ñ 		
ñ /ñ		
ó ,ùs   ‚A1A3z)Machete output path is not a directory: `z`.z)Machete output directory must be empty: `T)Úparentsz.cuÚwzutf-8)ÚencodingzRendered template to ))é€   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	  ÚTmaMIÚTmaCoopr)   ÚStreamKrw   r3   ra   r*   Úu4b8Úu8b128ÚzipÚ	itertoolsÚrepeatr\   r$   r¬   r­   r  r	   rþ   ÚexistsÚis_dirÚRuntimeErrorÚanyÚiterdirÚmkdirrò   ÚopenÚwriteÚprint)r  Úsch_common_paramsÚdefault_tile_heuristic_configÚcondÚtile_configÚdefault_heuristicr  r¢   ÚGPTQ_kernel_type_configsrh   ÚAWQ_kernel_type_configsÚoutput_pathÚfilenameÚcodeÚfilepathÚoutput_files   &               r   ÚgeneraterA  '  sy  € ô ÜÜ!Ü(×0Ñ0ôÐð%à,Ð.Eð%ð 	Ð*ð%ð
 	,Ð-Cð%ð 	,Ð-Dð%ð 	Ð*ð%ð 	+Ð,Bð%ð 	+Ð,Bð%ð 	,Ð-Dð%ð 	Ð)ð%ð 	+Ð,Bð%ð 	-Ð.Dð%ð  	Ð(ð!%ð$ 	,Ð-Cð%%ð& 	Ð(ð'%ð* 	Ð,ð+%ð, 	Ð$ð-%Ð!ð: "?×!DÑ!DÔ!Fôá!FÑˆDð 
Œ~˜{Ð@Ð.?Ñ@ÓAÙ!Fð ñ õ
ð €Lä#ñ $ô ×#Ñ#¤\×%8Ñ%8Ñ9ó$ó  Ðð äØ$Ü×ÒÑ1Ð2CÓDÓEÜ×ÒÐ.Ó/ô
óñ
ˆAô 	�1�Q•4˜˜1�˜q �tÖ$ñ
ñõ €Lô #ñ #ô —+‘+œxŸ{™{Ñ+ó#ó Ðð äØ#Ü×ÒÑ1Ð2CÓDÓEÜ×ÒÐ.Ó/ô
óñ
ˆAô 	�1�Q•4˜˜1�˜q �tÖ$ñ
ñõ €Lðf ÒÜ)Ó+‰ä˜:Ó&×1Ñ1Ó3ˆØ×Ñ×ÒØ×!Ñ!×#Ò#ÜÐ!JÈ;È-ÐWYÐZÓ[Ð[Üˆ{×"Ñ"Ó$×%Ò%ÜÐ!JÈ;È-ÐWYÐZÓ[Ð[ð &ð 	×Ñ $ÐÔ'ô )¨Ö6‰ˆØ :¨SÐ!1Õ1ˆØ�]‰]˜3¨ˆ]×1Ô1°[Ø×Ñ˜dÔ#÷ 2äÐ% h ZÐ0Ö1ó	 7ùóqùò@ùò0÷D 2×1ús   Á*KÃ;+KÆ+K$Ê+K)Ë)K9Ú__main__z#Generate Machete CUDA source files.)Údescriptionz--output-dir)ÚtypeÚdefault)é   r¥   )MÚargparser*  rÚ   rû   Úcollections.abcr   Úcopyr   Údataclassesr   r   Ú	functoolsr   Úpathlibr	   r¼   r   Ú__file__Úresolver  Ú_ROOTÚ_CUTLASS_EXT_DIRrü   rý   Ú_CUTLASS_ROOTÚ_CUTLASS_PYTHON_DIRr   r#   r.  Úvllm_cutlass_library_extensionr$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r.   r/   r0   ÚDISPATCH_TEMPLATEÚIMPL_TEMPLATEÚPREPACK_TEMPLATEÚTmaWarpSpecializedCooperativer$  r%  r3   rJ   rV   r\   rn   r}   r‰   r�   r“   rŸ   r©   r²   r³   rÀ   rÃ   rÕ   rß   rÙ   rò   r  rA  r@   ÚArgumentParserÚparserÚadd_argumentÚ
parse_argsÚargsr  r>   r   r   Ú<module>r]     s¡  ðó Û Û Û 	Ý $Ý ß )Ý Ý ã ã 
áˆX‹×ÑÓ ×(Ñ(¨Õ+€Ø˜?Õ*Ð-AÕAÐ à—
‘
—‘Ð6Ó7€ØÒÙ˜Ó'�Mà˜IÕ%€Mà# hÕ.Ð õ"õñ Ð"Ô #Ù#Ð$7×8Ò8ÙÐ)Õ*á
ð	ØÐ Ð Yð	[óð ÷
÷ ÷ õ ð(|Ð ð|0€ðd!Ð ðF 	%×BÑB€Ø
×
<Ñ
<€ñ �$Ô÷&ð &ó ð&ñ �$Ô÷ð ó ðñ �$Ô÷ð ó ðð ÷7ð 7ó ð7õõ0
õõò+õ
(õòð ˆH�M‰MØ�?ØÐ6ØÐ0ØÐ.ØÐ.ØÐ(ØÐ(ØÐ)ØÐ+ØÐ(Ø!Ð#>ØÐ5ðÐ ò"ñ 'Ð'8Ó9Ð Ù" =Ó1Ð Ù+Ð,<Ó=Ð ÷Wõt;÷A2ðH ˆzÔØ×$Ò$Ð1VÔW€FØ
×Ñ˜¨T¸4ÐÔ@Ø×ÑÓ€DÙˆT�_‰_Öñ	 r   