o
    c۷i                     @   sB   d dl Z d dlmZ d dlmZ d dlmZ G dd dejZdS )    N)Integer)dsl_user_opc                       sJ   e Zd ZdZe	dddddedef fddZdd	 Zd
d Z	  Z
S )
FastDivmodz6We store the divisor along with the FastDivmodDivisor.N)locipdivisoris_power_of_2c                   s   t  j||||d || _d S )N)r   r   r   )super__init__r   )selfr   r   r   r   	__class__ E/home/ubuntu/vllm_env/lib/python3.10/site-packages/quack/fast_math.pyr
      s   	
zFastDivmod.__init__c                 C   s   | j gt| j S )z.Extract MLIR values for Host->Device transfer.)_divisorcutlassextract_mlir_valuesr   )r   r   r   r   __extract_mlir_values__   s   z"FastDivmod.__extract_mlir_values__c                 C   s0   t t}|d |_t| j|dd |_|S )z/Reconstruct FastDivmodDivisor from MLIR values.r      N)object__new__r   r   r   new_from_mlir_valuesr   )r   valuesnew_objr   r   r   __new_from_mlir_values__   s   

z#FastDivmod.__new_from_mlir_values__)N)__name__
__module____qualname____doc__r   r   boolr
   r   r   __classcell__r   r   r   r   r   	   s    r   )	r   cutlass.cutecutecutlass.base_dsl.typingr   cutlass.cutlass_dslr   FastDivmodDivisorr   r   r   r   r   <module>   s
   