o
    پi                  
   @   sD  d dl mZ d dlmZ d dlmZmZ d dlZzd dlm	Z	 W n e
y+   dZ	Y nw d dlZd dlmZ d dlmZmZmZmZmZ d dlmZ d dlmZ ejeeeeeedfZejjj j!Z"ej#Z$ej%eej&eej'eej(eej)eiZ*ed	d
 Z+eddej,deeef fddZ-eG dd dZ.eG dd deZ/dS )    )Tuple)	lru_cache)	dataclassfieldsN)extract)Int32Int64Float16BFloat16Float32)JitArgument)NumericMetac                 C   s   t j j| dS )Ncluster_size)cutlassutilsHardwareInfoget_max_active_clustersr    r   H/home/ubuntu/.local/lib/python3.10/site-packages/quack/cute_dsl_utils.pyr   %   s   r   devicereturnc                 C   s   t j| S N)torchcudaget_device_capability)r   r   r   r   get_device_capacity*   s   r   c                   @   s   e Zd Zdd Zdd ZdS )
ParamsBasec                    s`    fddt  D }dd |D }g g } _|D ]}t|}||7 } jt| q|S )Nc                       g | ]}t  |jqS r   getattrname.0fieldselfr   r   
<listcomp>2       z6ParamsBase.__extract_mlir_values__.<locals>.<listcomp>c                 S      g | ]	}t |ts|qS r   
isinstanceStaticTypesr#   fr   r   r   r'   3       )r   _values_posr   extract_mlir_valuesappendlen)r&   
all_fieldsnon_constexpr_fieldsvaluesobj
obj_valuesr   r%   r   __extract_mlir_values__1   s   
z"ParamsBase.__extract_mlir_values__c                        fddt  D }dd | D }dd | D }t|  jD ]\\}}}t||d | ||< ||d  }q% jdi ||S )Nc                       i | ]
}|j t |j qS r   r!   r    r"   r%   r   r   
<dictcomp><       z7ParamsBase.__new_from_mlir_values__.<locals>.<dictcomp>c                 S       i | ]\}}t |tr||qS r   r*   r#   nr.   r   r   r   r=   =        c                 S       i | ]\}}t |ts||qS r   r*   r@   r   r   r   r=   >   
    r   r   itemszipr0   r   new_from_mlir_values	__class__r&   r6   r4   constexpr_fieldsr5   r!   r$   n_itemsr   r%   r   __new_from_mlir_values__;      z#ParamsBase.__new_from_mlir_values__N)__name__
__module____qualname__r9   rM   r   r   r   r   r   /   s    
r   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )ArgumentsBasec                    sN    fddt  D }dd |D }g }|D ]}t|dr$||  q|S )Nc                    r   r   r   r"   r%   r   r   r'   J   r(   z0ArgumentsBase.__c_pointers__.<locals>.<listcomp>c                 S   r)   r   r*   r-   r   r   r   r'   K   r/   __c_pointers__)r   hasattrextendrS   )r&   r4   r5   c_ptrsr7   r   r%   r   rS   I   s   
zArgumentsBase.__c_pointers__c                    sx    fddt  D }dd |D }g g } _|D ]}t|dr3| }||  jt| q jd q|S )Nc                    r   r   r   r"   r%   r   r   r'   S   r(   z4ArgumentsBase.__get_mlir_types__.<locals>.<listcomp>c                 S   r)   r   r*   r-   r   r   r   r'   T   r/   __get_mlir_types__r   )r   r0   rT   rW   rU   r2   r3   )r&   r4   r5   typesr7   	obj_typesr   r%   r   rW   R   s   

z ArgumentsBase.__get_mlir_types__c                    r:   )Nc                    r;   r   r<   r"   r%   r   r   r=   `   r>   z:ArgumentsBase.__new_from_mlir_values__.<locals>.<dictcomp>c                 S   r?   r   r*   r@   r   r   r   r=   a   rB   c                 S   rC   r   r*   r@   r   r   r   r=   b   rD   r   rE   rJ   r   r%   r   rM   _   rN   z&ArgumentsBase.__new_from_mlir_values__N)rO   rP   rQ   rS   rW   rM   r   r   r   r   rR   G   s    	rR   r   )0typingr   	functoolsr   dataclassesr   r   r   triton.tools.disasmr   ImportErrorr   cutlass.cutecuter   r   r	   r
   r   cutlass.base_dsl.typingr   cutlass.cutlass_dslr   	Constexprintboolstrfloattyper,   base_dslruntimer   load_cubin_module_dataload_cubin_module_data_ogcompilecute_compile_ogfloat16bfloat16float32int32int64torch2cute_dtype_mapr   r   r   r   rR   r   r   r   r   <module>   s>   	
 