Ë
    $täiG  ã                  óT  — d dl mZ d dlmZ d dlmZ d dlmc mZ d dl	m
Z
 d dlmZmZmZmZmZmZmZmZmZ d dlmZ d dlmZmZmZmZmZmZmZ d d	lm Z  erd d
l!m"Z" d dlm#Z#  ed¬«       G d„ d«      «       Z$ ed¬«       G d„ d«      «       Z%d!d„Z&	 	 	 	 d"d„Z' ed¬«       G d„ d«      «       Z( e(d dejR                  ¬«       e(ddejT                  ¬«      gZ+	 	 	 	 	 	 	 	 	 	 d#d„Z,ed$d„«       Z- ed¬«       G d„ d«      «       Z.ed$d„«       Z/	 	 	 	 	 	 	 	 	 	 d%d„Z0ed$d „«       Z1y)&é    )Úannotations)Ú	dataclass)ÚTYPE_CHECKINGN)Ú	translate)	Ú	BaseCTypeÚBindingÚCTypeÚExprÚ
NamedCTypeÚopmath_tÚscalar_tÚStructuredImplSignatureÚVectorizedCType)Úwith_native_function)ÚArgumentÚBaseTyÚBaseTypeÚDispatchKeyÚNativeFunctionsGroupÚ
ScalarTypeÚUfuncKey)Ú
OrderedSet)ÚSequence)ÚUfunctorBindingsT)Úfrozenc                  ó\   — e Zd ZU ded<   ded<   ded<   dd„Zdd„Zdd	„Zdd
„Zdd„Zdd„Z	y)ÚUfunctorSignaturer   Úgz
int | NoneÚscalar_tensor_idxÚstrÚnamec                ób   — t        j                  | j                  | j                  t        ¬«      S )N)r   r   )ÚufuncÚufunctor_argumentsr   r   r   ©Úselfs    úb/Volumes/fast/ai/experiments/MLX_z-image/.venv/lib/python3.12/site-packages/torchgen/dest/ufunc.pyÚ	argumentszUfunctorSignature.argumentsI   s&   € Ü×'Ñ'Ø�F‰F d×&<Ñ&<Äxô
ð 	
ó    c                ó”   — | j                  «       j                  D �cg c]   }|j                  |j                  › d�«      ‘Œ" c}S c c}w )NÚ_)r(   ÚctorÚrenamer!   )r&   Úbs     r'   ÚfieldszUfunctorSignature.fieldsN   s<   € à04·±Ó0@×0EÒ0EÓFÑ0E¨1�—‘˜AŸF™F˜8 1˜Õ&Ð0EÑFÐFùÒFs   �%Ac                ó    — t        t        «      S ©N)r   r   r%   s    r'   Úreturns_typezUfunctorSignature.returns_typeR   s   € ô œÓ"Ð"r)   c                óN   — dj                  d„ | j                  «       D «       «      S )NÚ
c              3  óT   K  — | ]   }|j                   › d |j                  › d�–— Œ" y­w)Ú Ú;N)Útyper!   )Ú.0Úfs     r'   Ú	<genexpr>z0UfunctorSignature.decl_fields.<locals>.<genexpr>X   s&   è ø€ ÐF¹°1˜AŸF™F˜8 1 Q§V¡V H¨AÔ.¹ùó   ‚&()Újoinr/   r%   s    r'   Údecl_fieldszUfunctorSignature.decl_fieldsW   s   € Ø�y‰yÑF¸¿¹¼ÓFÓFÐFr)   c                óì   — dj                  d„ | j                  «       j                  D «       «      }dj                  d„ | j                  «       j                  D «       «      }| j                  › d|› d|› d�S )Nú, c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wr1   ©Údecl©r9   Úas     r'   r;   z5UfunctorSignature.inline_defn_ctor.<locals>.<genexpr>[   s   è ø€ ÐEÑ/D¨!˜QŸV™VŸXÑ/Dùó   ‚c              3  óT   K  — | ]   }|j                   › d |j                   › d�–— Œ" y­w)z_(Ú)N©r!   rD   s     r'   r;   z5UfunctorSignature.inline_defn_ctor.<locals>.<genexpr>^   s(   è ø€ ÐSÑ=R¸ §¡˜x r¨!¯&©&¨°Ô3Ñ=Rùr<   Ú(z) : z {})r=   r(   r,   r!   )r&   Úargs_strÚinit_strs      r'   Úinline_defn_ctorz"UfunctorSignature.inline_defn_ctorZ   sc   € Ø—9‘9ÑE¨t¯~©~Ó/?×/DÒ/DÓEÓEˆð —9‘9ÑS¸T¿^¹^Ó=M×=RÒ=RÓSÓSˆØ—)‘)�˜A˜h˜Z t¨H¨:°UÐ;Ð;r)   c                óª   — dj                  d„ | j                  «       j                  D «       «      }| j                  «       j	                  «       › d|› d�S )Nr@   c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wr1   rB   rD   s     r'   r;   z/UfunctorSignature.decl_apply.<locals>.<genexpr>b   s   è ø€ ÐFÑ/E¨!˜QŸV™VŸXÑ/EùrF   z operator()(z) const)r=   r(   Úapplyr2   Úcpp_type)r&   rK   s     r'   Ú
decl_applyzUfunctorSignature.decl_applya   sI   € Ø—9‘9ÑF¨t¯~©~Ó/?×/EÒ/EÓFÓFˆØ×#Ñ#Ó%×.Ñ.Ó0Ð1°¸h¸ZÀwÐOÐOr)   N)Úreturnr   ©rS   zlist[Binding])rS   r	   ©rS   r    )
Ú__name__Ú
__module__Ú__qualname__Ú__annotations__r(   r/   r2   r>   rM   rR   © r)   r'   r   r   C   s4   … àÓØ!Ó!Ø
ƒIó
ó
Gó#ó
Gó<ôPr)   r   c                  ó<   — e Zd ZU ded<   ded<   ded<   d
d„Zdd„Zy	)ÚUfuncSignaturer   r   r    r!   r	   Ú	compute_tc                óX   — t        j                  | j                  | j                  ¬«      S )N)r]   )r#   Úufunc_argumentsr   r]   r%   s    r'   r(   zUfuncSignature.argumentsl   s   € Ü×$Ñ$ T§V¡V°t·~±~ÔFÐFr)   c           
     ó‚   — | j                   › ddj                  d„ t        || j                  «       «      D «       «      › d�S )NrJ   r@   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr1   ©ÚexprrD   s     r'   r;   z&UfuncSignature.call.<locals>.<genexpr>p   s   è ø€ Ð'YÑ8X°1¨¯­Ñ8Xùó   ‚rH   ©r!   r=   r   r(   ©r&   Úctxs     r'   ÚcallzUfuncSignature.callo   s8   € Ø—)‘)�˜A˜dŸi™iÑ'Y¼	À#ÀtÇ~Á~ÓGWÔ8XÓ'YÓYÐZÐZ[Ð\Ð\r)   NrT   )rg   zSequence[Binding | Expr]rS   r    )rV   rW   rX   rY   r(   rh   rZ   r)   r'   r\   r\   f   s   … àÓØ
ƒIØÓóGô]r)   r\   c                ó€   — t        d„ | j                  j                  j                  j                  D «       «      }|dk(  S )Nc              3  óV   K  — | ]!  }|j                   j                  «       sŒd –— Œ# y­w)é   N)r8   Úis_tensor_likerD   s     r'   r;   z<eligible_for_binary_scalar_specialization.<locals>.<genexpr>†   s$   è ø€ ð Ù;ˆa¸q¿v¹v×?TÑ?TÕ?VŒÑ;ùs   ‚)¢)é   )ÚsumÚ
functionalÚfuncr(   Úflat_non_out)r   Únum_tensorss     r'   Ú)eligible_for_binary_scalar_specializationrs   …   s<   € Üñ Ø—<‘<×$Ñ$×.Ñ.×;Ò;óó €Kð ˜!ÑÐr)   c                ó<  — i }g }| j                   j                  }t        j                  dt        j                  dt        j
                  d i}t        | «      r0t        j                  t        j                  t        j
                  g}nIt        j
                  g}t        j                  t        j                  fD ]  }||v sŒt        d|› d�«      ‚ |D �]´  }||v rHt        | ||   ||   j                  ¬«      }||   j                  D ]  }||j                  |i «      |<   Œ ŒPd }	t        «       }
t        j                  t        j                  fD ]H  }||vrŒ|	€||   j                  }	n|	||   j                  k7  rt        d«      ‚|
||   j                  z  }
ŒJ |	€t        d«      ‚|› d|	› �}t        | ||   |¬«      }|
D ]  }||j                  |i «      |<   Œ t        | d	|	› �t!        t"        «      ¬
«      }|j%                  «       |j'                  «       j(                  z   }|j+                  d|j                  › d|j-                  «       › d|j/                  «       › d|j1                  «       › d|j3                  |«      › d�«       �Œ· |dj5                  |«      fS )Nrk   r   zcannot use z on non-binary function)r   r!   z0ScalarOnly and Generic must have same ufunc namezufunc_name must be non-Noner+   úufunc::©r!   r]   z%
template <typename scalar_t>
struct z3 {
  using opmath_t = at::opmath_type<scalar_t>;
  z
  z
  __device__ z {
    return z	;
  }
};
r4   )ÚoutÚufunc_inner_loopr   ÚCUDAFunctorOnSelfÚCUDAFunctorOnOtherÚCUDAFunctorrs   ÚAssertionErrorr   r!   Úsupported_dtypesÚ
setdefaultr   Ú
ScalarOnlyÚGenericr\   r   r   r/   r(   rP   Úappendr>   rM   rR   rh   r=   )r   Úufunctor_sigsÚ	ufunctorsÚloopsÚscalar_tensor_idx_lookupÚkeysÚkÚufunctor_sigÚdtypeÚ
ufunc_namer}   Úlkr!   Ú	ufunc_sigÚ	apply_ctxs                  r'   Úcompute_ufunc_cuda_functorsrŽ   Œ   sÒ  € ð JL€MØ€IØ�E‰E×"Ñ"€Eä×"Ñ" AÜ×#Ñ# QÜ×Ñ˜dð Ðô
 1°Ô3ä×&Ñ&Ü×'Ñ'Ü× Ñ ð
‰ô ×$Ñ$Ð%ˆÜ×,Ñ,¬h×.IÑ.IÓJˆAØ�EŠzÜ$ {°1°#Ð5LÐ%MÓNÐNð Kô ˆð �‰:Ü,ØÐ%=¸aÑ%@ÀuÈQÁxÇ}Á}ôˆLð ˜q™×2Ô2�Ø9E�×(Ñ(¨°Ó3°AÒ6ð 3àð ˆ
Ü3=³<ÐÜ×&Ñ&¬×(8Ñ(8Ó9ˆBØ˜‰ØØÐ!Ø" 2™YŸ^™^‘
ð   r¡§¡Ò/Ü(ØJóð ð   b¡	× :Ñ :Ñ:Ñð :ð ÐÜ Ð!>Ó?Ð?à��A�j�\Ð"ˆÜ(ØÐ!9¸!Ñ!<À4ô
ˆó &ˆEØ5AˆM×$Ñ$ U¨BÓ/°Ò2ð &ô #Ø�g˜j˜\Ð*´iÄÓ6Iô
ˆ	ð !×'Ñ'Ó)¨L×,BÑ,BÓ,D×,JÑ,JÑJˆ	Ø×Ñðà×ÑÐð à×ÑÓÐð Ø× Ñ Ó"Ð#ð $Ø×%Ñ%Ó'Ð(ð )Ø�N‰N˜9Ó%Ð&ð 'ð
ö	
ð] ðx ˜$Ÿ)™) IÓ.Ð.Ð.r)   c                  ó,   — e Zd ZU ded<   ded<   ded<   y)Ú BinaryScalarSpecializationConfigÚintÚ
scalar_idxr    Úctor_tensorr   Ú	ufunc_keyN)rV   rW   rX   rY   rZ   r)   r'   r�   r�   â   s   … àƒOØÓØÔr)   r�   r&   )r’   r“   r”   rk   Úotherc                óŽ  — d}|dz  }t         D ]Ñ  }|j                  |vrŒ||j                     }|j                  dz   }t        |«      }|j	                  t        d|› d�t        |j                  t        t        «      «      ¬«      «       dj                  d„ t        ||j                  «       j                  «      D «       «      }	|d	|› d
|j                  › d|	› d|› d�	z  }ŒÓ |t        j                      }dj                  d„ t        ||j                  «       j                  «      D «       «      }	|d|j                  › d|	› d�z  }|S )Nz+using opmath_t = at::opmath_type<scalar_t>;zif (false) {}
rk   ziter.scalar_value<opmath_t>(rH   )rc   r8   r@   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr1   rb   rD   s     r'   r;   z0compute_ufunc_cuda_dtype_body.<locals>.<genexpr>  s   è ø€ ð ,
ÙI�qˆA�F�FÑIùrd   zelse if (iter.is_cpu_scalar(z)) {
  z<scalar_t> ufunctor(z);
  iter.remove_operand(z");
  gpu_kernel(iter, ufunctor);
}c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr1   rb   rD   s     r'   r;   z0compute_ufunc_cuda_dtype_body.<locals>.<genexpr>  s   è ø€ ð (ÙL�1ˆ��ÑLùrd   z
else {
  gpu_kernel(iter, z<scalar_t>(z
));
}
    )Ú!BinaryScalarSpecializationConfigsr”   r’   Úlistr�   r
   r   r“   r   r   r=   r   r(   r,   r!   r   r{   )
r   r‰   Úinner_loopsÚ
parent_ctxÚbodyÚconfigrˆ   r’   rg   Úufunctor_ctor_exprs_strs
             r'   Úcompute_ufunc_cuda_dtype_bodyr    ÷   s„  € ð 9€DØÐÑ€Dß3ˆØ×Ñ ;Ñ.ØØ" 6×#3Ñ#3Ñ4ˆØ×&Ñ&¨Ñ*ˆ
ô %)¨Ó$4ˆØ�
‰
ÜØ3°J°<¸qÐAÜ × 2Ñ 2´I¼hÓ4GÓHôô	
ð #'§)¡)ñ ,
Ü% c¨<×+AÑ+AÓ+C×+HÑ+HÔIó,
ó #
Ðð 	ð Ø'˜Lð )Ø×ÑÐÐ)Ð*AÐ)Bð CØ!�lð #ðñ 	‰ð) 4ð6 œx×3Ñ3Ñ4€LØ"Ÿi™iñ (Ü! *¨l×.DÑ.DÓ.F×.KÑ.KÔLó(ó Ðð 	ð à ×%Ñ%Ð& kÐ2IÐ1Jð Kðñ €Dð
 €Kr)   c                ór  — t        | «      \  }}t        | t        j                  | t        j
                  «      «      }g }|j                  «       D ]7  \  }}|j                  d|› dt        | |||j                  «       «      › d�«       Œ9 dj                  |«      }t        | «      }d|› d|j                  «       › d|j                  «       › d|j                  «       › d|j                  › d|› d	|j                  › d
|j                  › d|j!                  «       › d|j#                  |j                  «       «      › d�S )Nú"
AT_DISPATCH_CASE(at::ScalarType::ú,
  [&]() {
    ú
  }
)
r4   z

ú;
ú. {
  AT_DISPATCH_SWITCH(iter.common_dtype(), "ú",
    z
  );
}
REGISTER_DISPATCH(ú, &z)

ú {
  ú;
}
)rŽ   r   r#   Úkernel_namer   ÚCUDAÚitemsr�   r    r(   r=   ÚStubSignatureÚ	type_defnÚdispatch_declÚkernel_defnr!   ÚdefnÚdirect_call)	r   r‚   rƒ   ÚsigÚdtype_casesr‰   Úinner_ufunc_sigsÚdtype_cases_strÚstub_sigs	            r'   Úcompute_ufunc_cudar¹   &  sc  € ô  ;¸1Ó=Ñ€M�9ô " !¤U×%6Ñ%6°q¼+×:JÑ:JÓ%KÓ
L€CØ€KØ#0×#6Ñ#6Ö#8ÑˆÐØ×Ñð"Ø"' ð )ä" 1 eÐ-=¸s¿}¹}»ÓOÐPð Qðõ	
ð $9ð —i‘i Ó,€Oä˜QÓ€HðØ
€ð à	×ÑÓÐ ð Ø	×ÑÓÐ ð à	×ÑÓÐ ð ,Ø,/¯H©H¨:ð 6ØÐð ð —=‘=�/  X×%9Ñ%9Ð$:ð ;à‡�ƒ€ð Ø×Ñ˜Ÿ™›Ó(Ð)ð *ðð r)   c                  óŽ   — e Zd ZU ded<   edd„«       Zedd„«       Zedd„«       Zdd„Zdd„Z	dd„Z
dd	„Zdd
„Zdd„Zdd„Zdd„Zy)r®   r   r   c                ó‚   — t        | j                  j                  j                  j                  j                  «      › d�S )NÚ_stub©r    r   ro   rp   r!   r%   s    r'   r!   zStubSignature.name[  s1   € ä�d—f‘f×'Ñ'×,Ñ,×1Ñ1×6Ñ6Ó7Ð8¸Ð>Ð>r)   c                ó‚   — t        | j                  j                  j                  j                  j                  «      › d�S )NÚ_kernelr½   r%   s    r'   r«   zStubSignature.kernel_name_  s1   € ä�d—f‘f×'Ñ'×,Ñ,×1Ñ1×6Ñ6Ó7Ð8¸Ð@Ð@r)   c                ó‚   — t        | j                  j                  j                  j                  j                  «      › d�S )NÚ_fnr½   r%   s    r'   Ú	type_namezStubSignature.type_namec  s1   € ä�d—f‘f×'Ñ'×,Ñ,×1Ñ1×6Ñ6Ó7Ð8¸Ð<Ð<r)   c                ó@   — t        j                  | j                  «      S r1   )r#   Ústub_argumentsr   r%   s    r'   r(   zStubSignature.argumentsg  s   € Ü×#Ñ# D§F¡FÓ+Ð+r)   c                óZ   — | j                  «       }ddj                  d„ |D «       «      › d�S )Nzvoid(*)(TensorIteratorBase&, r@   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr1   )r8   rD   s     r'   r;   z%StubSignature.type.<locals>.<genexpr>l  s   è ø€ Ð8RÉÀA¸¿½Éùrd   rH   )r(   r=   )r&   Úcpp_argss     r'   r8   zStubSignature.typej  s.   € Ø—>‘>Ó#ˆØ.¨t¯y©yÑ8RÉÓ8RÓ/RÐ.SÐSTÐUÐUr)   c                ó<   — d| j                   › d| j                  › d�S )NzDECLARE_DISPATCH(r@   rH   )rÂ   r!   r%   s    r'   r°   zStubSignature.dispatch_decln  s    € Ø" 4§>¡>Ð"2°"°T·Y±Y°K¸qÐAÐAr)   c                ó"   — d| j                   › d�S )NzDEFINE_DISPATCH(rH   rI   r%   s    r'   Údispatch_defnzStubSignature.dispatch_defnq  s   € Ø! $§)¡) ¨AÐ.Ð.r)   c                óp   — d| j                   › ddj                  d„ | j                  «       D «       «      › d�S )Nzvoid z(TensorIteratorBase& iter, r@   c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wr1   )r²   rD   s     r'   r;   z,StubSignature.kernel_defn.<locals>.<genexpr>u  s   è ø€ ÐMqÑ`pÐ[\ÈaÏfÉfÏhÑ`pùrF   rH   )r«   r=   r(   r%   s    r'   r±   zStubSignature.kernel_defnt  s:   € Ø�t×'Ñ'Ð(Ð(CÀDÇIÁIÑMqÐ`d×`nÑ`nÔ`pÓMqÓDqÐCrÐrsÐtÐtr)   c                óB   — d| j                   › d| j                  «       › �S )Nzusing ú = )rÂ   r8   r%   s    r'   r¯   zStubSignature.type_defnw  s    € Ø˜Ÿ™Ð' s¨4¯9©9«;¨-Ð8Ð8r)   c           
     ó‚   — | j                   › ddj                  d„ t        || j                  «       «      D «       «      › d�S )Nz(device_type(), *this, r@   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr1   rb   rD   s     r'   r;   z%StubSignature.call.<locals>.<genexpr>|  s   è ø€ Ð=oÑNnÈ¸a¿f½fÑNnùrd   rH   re   rf   s     r'   rh   zStubSignature.call{  s=   € Ø—)‘)�Ð3°D·I±IÑ=oÌiÐX[Ð]a×]kÑ]kÓ]mÔNnÓ=oÓ4oÐ3pÐpqÐrÐrr)   c           
     ó‚   — | j                   › ddj                  d„ t        || j                  «       «      D «       «      › d�S )Nz(*this, r@   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr1   rb   rD   s     r'   r;   z,StubSignature.direct_call.<locals>.<genexpr>€  s   è ø€ Ð5gÑFfÀ°a·fµfÑFfùrd   rH   )r«   r=   r   r(   rf   s     r'   r³   zStubSignature.direct_call  s?   € Ø×"Ñ"Ð# 8¨D¯I©IÑ5gÄiÐPSÐUY×UcÑUcÓUeÔFfÓ5gÓ,gÐ+hÐhiÐjÐjr)   NrU   rT   )rg   úSequence[Binding]rS   r    )rV   rW   rX   rY   Úpropertyr!   r«   rÂ   r(   r8   r°   rÊ   r±   r¯   rh   r³   rZ   r)   r'   r®   r®   W  sq   … àÓàò?ó ð?ð òAó ðAð ò=ó ð=ó,óVóBó/óuó9ósôkr)   r®   c                óB  — t        | «      }t        | t        j                  | t        j
                  «      «      }d|j                  «       › d|j                  «       › d|j                  «       › d|j                  «       › d|j                  |j                  «       «      › d�S )Nr4   r¥   z;

r©   rª   )r®   r   r#   r«   r   ÚCPUr¯   r°   rÊ   r²   rh   r(   )r   r¸   r´   s      r'   Úcompute_ufunc_cpur×   ƒ  sž   € ä˜QÓ€HÜ
! !¤U×%6Ñ%6°q¼+¿/¹/Ó%JÓ
K€CðØ	×ÑÓÐ ð Ø	×ÑÓÐ ð Ø	×ÑÓÐ ð à‡�ƒ€ð Ø‡=�=�—‘“Ó!Ð"ð #ðð r)   c                ó”  ‡— t         j                  |vrt        |› d|j                  «       › �«      ‚|j                  «       t         j                  t         j                  hk  st        d|j                  «       › �«      ‚|t         j                     }d }t         j                  |v r|t         j                     }g }g Š|D ]Ì  }t        |j                  t        «      r1|j                  j                  t        t        j                  «      k7  rŒN|j                  d|j                  › d|j                  › d�«       ‰j                  t        d|j                  › �t        |j                   j                  t#        t$        «      «      «      «       ŒÎ |�Ú|D ]Õ  }t        |j                  t        «      r1|j                  j                  t        t        j                  «      k7  rŒN|j                  d|j                  › d|j                  › d	�«       ‰j                  t        d
|j                  › �t        |j                   j                  t'        t#        t$        «      «      «      «      «       Œ× g }g }	| j(                  j*                  j,                  j.                  D ]÷  }
|
j                  j1                  «       sŒ|
j                  t        t        j2                  «      k7  rt        d|
j                  › �«      ‚|j                  t5        |
j                  t        |
j                  t#        t$        «      «      |
¬«      «       |€Œ§|	j                  t5        |
j                  t        |
j                  t'        t#        t$        «      «      «      |
¬«      «       Œù dˆfd„}dj7                  |«      }|�hd|› ddj7                  d„ |D «       «      › d|j9                   ||«      «      › ddj7                  d„ |	D «       «      › d|j9                   ||	«      «      › d�S d|› ddj7                  d„ |D «       «      › d|j9                   ||«      «      › d�S )Nr@   z<inner_loops keys must be subset of CPUScalar/CPUVector, got zauto _s_rÎ   z.to<scalar_t>();Ú_s_zauto _v_z$ = at::vec::Vectorized<scalar_t>(_s_z);Ú_v_zExpected Tensor type, got )r!   ÚnctypeÚargumentc                óP   •— g }|j                  ‰«       |j                  | «       |S r1   )Úextend)r.   Úrrg   s     €r'   Úwith_ctxz.compute_ufunc_cpu_dtype_body.<locals>.with_ctxÝ  s#   ø€ Ø"$ˆØ	�‰�ŒØ	�‰�ŒØˆr)   r4   z
cpu_kernel_vec(iter,
  [=](c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wr1   rB   ©r9   r.   s     r'   r;   z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>è  ó   è ø€ Ð3¡?˜a�—‘—¡?ùrF   z) { return z; },
  [=](c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wr1   rB   râ   s     r'   r;   z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>é  s   è ø€ Ð0¡<˜a�—‘—¡<ùrF   z; }
);
z
cpu_kernel(iter,
  [=](c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wr1   rB   râ   s     r'   r;   z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>ð  rã   rF   )r.   rÓ   rS   zlist[Expr | Binding])r   Ú	CPUScalarr|   r†   Ú	CPUVectorÚ
isinstancerÜ   r   r8   r   r   ÚScalarr�   r!   r
   r   rÛ   r   r   r   ro   rp   r(   rq   rl   ÚTensorr   r=   rh   )r   r‰   r›   rœ   Úscalar_loopÚvec_loopr�   r.   Úscalar_bindingsÚvec_bindingsrE   rà   Úbody_strrg   s                @r'   Úcompute_ufunc_cpu_dtype_bodyrð   “  s³  ø€ ô ×Ñ Ñ,Ü ˜w b¨×)9Ñ)9Ó);Ð(<Ð=Ó>Ð>Ø×ÑÓ¤(×"4Ñ"4´h×6HÑ6HÐ!IÒIÜØJÈ;×K[ÑK[ÓK]ÐJ^Ð_ó
ð 	
ð œh×0Ñ0Ñ1€KØ€HÜ×Ñ˜[Ñ(Øœx×1Ñ1Ñ2ˆð €DØ
€CÛˆÜ�a—j‘j¤(Ô+°·
±
·±Ä8Ü�M‰MóD
ò 1
ð Ø�‰�h˜qŸv™v˜h c¨!¯&©&¨Ð1AÐBÔCØ�
‰
”4˜#˜aŸf™f˜X˜¬
°1·8±8·=±=Ä)ÌHÓBUÓ(VÓWÕXð ð ÐÛˆAÜ˜!Ÿ*™*¤hÔ/°A·J±J·O±OÄxÜ—‘óHò 5ð Ø�K‰KØ˜1Ÿ6™6˜(Ð"FÀqÇvÁvÀhÈbÐQôð �J‰JÜØ˜!Ÿ&™&˜�NÜ˜qŸx™xŸ}™}¬o¼iÌÓ>QÓ.RÓSóõð ð" €OØ€LØ�\‰\×Ñ×(Ñ(×5Ô5ˆØ�v‰v×$Ñ$Ô&ØØ�6‰6”XœfŸm™mÓ,Ò,Ü Ð#=¸a¿f¹f¸XÐ!FÓGÐGØ×ÑÜØ—V‘VÜ! !§&¡&¬)´HÓ*=Ó>Øôô	
ð ÑØ×ÑÜØŸ™Ü% a§f¡f¬o¼iÌÓ>QÓ.RÓSØôõð 6õ*ð �y‰y˜‹€HØÐðØ	€
ð à‡y�yÑ3¡?Ó3Ó3Ð4°LÀ×AQÑAQÑRZÐ[jÓRkÓAlÐ@mð nØ‡y�yÑ0¡<Ó0Ó0Ð1°¸h¿m¹mÉHÐUaÓLbÓ>cÐ=dð eð	ð 	ðØ	€
ð à‡y�yÑ3¡?Ó3Ó3Ð4°LÀ×AQÑAQÑRZÐ[jÓRkÓAlÐ@mð nðð 	r)   c                ór  — t        | «      }| j                  j                  }i }t        j                  t        j
                  fD �]4  }g }||v r|j                  |«       t        j                  |v r1|t        j                  u r|j                  t        j                  «       t        j                  |v r|j                  t        j                  «       |D ]¡  }||   j                  D ]�  }|t        j                  u rt        t        «      }n1|t        j
                  u rt        t        t        «      «      }nt        ‚|j                  |i «      }	||	vsŒmt        | d||   j                   › �|¬«      |	|<   Œ� Œ£ �Œ7 g }
|j#                  «       D ]7  \  }}	|
j                  d|› dt%        | ||	|j'                  «       «      › d�«       Œ9 dj)                  |
«      }d|j+                  «       › d|j                   › d	|› d
|j-                  «       › d|j/                  «       › d|j                   › d|j0                  › d�S )Nru   rv   r¢   r£   r¤   r4   z
namespace {

r¦   r§   z#
  );
}

} // anonymous namespace

r¥   z
REGISTER_DISPATCH(r¨   z)
)r®   rw   rx   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„   Ú
ufunc_sigsr‡   Úlksr‹   r‰   r]   r¶   rµ   r·   s               r'   Úcompute_ufunc_cpu_kernelrô   õ  s3  € ä˜QÓ€Hð �E‰E×"Ñ"€EØCE€JÜ× Ñ ¤(×"4Ñ"4Ô5ˆØˆà�‰:Ø�J‰J�qŒMÜ×Ñ %Ñ'¨A´×1CÑ1CÑ,CØ�J‰J”x×*Ñ*Ô+Ü×Ñ˜uÑ$Ø�J‰J”x×'Ñ'Ô(ãˆBØ˜r™×3Ô3�àœ×*Ñ*Ñ*Ü )¬(Ó 3‘IØœ(×,Ñ,Ñ,Ü /´	¼(Ó0CÓ D‘Iä(Ð(Ø#-×#8Ñ#8¸ÀÓ#CÐ ØÐ,Ò,Ü*8Ø '¨%°©)¯.©.Ð)9Ð :Àiô+Ð$ QÒ'ñ 4ò ð 6ð4 €KØ#-×#3Ñ#3Ö#5ÑˆÐØ×Ñð"Ø"' ð )ä! ! UÐ,<¸h×>PÑ>PÓ>RÓSÐTð Uðõ	
ð $6ð —i‘i Ó,€Oðð 
×ÑÓÐ ð ,Ø,4¯M©M¨?ð ;ØÐð ð 
×ÑÓÐ ð Ø	×ÑÓÐ ð Ø—=‘=�/  X×%9Ñ%9Ð$:ð ;ðð r)   )r   r   rS   Úbool)r   r   rS   z?tuple[dict[ScalarType, dict[UfuncKey, UfunctorSignature]], str])
r   r   r‰   r   r›   z!dict[UfuncKey, UfunctorSignature]rœ   rÓ   rS   r    )r   r   rS   r    )
r   r   r‰   r   r›   zdict[UfuncKey, UfuncSignature]rœ   rÓ   rS   r    )2Ú
__future__r   Údataclassesr   Útypingr   Útorchgen.api.ufuncÚapir#   Útorchgen.api.translater   Útorchgen.api.typesr   r   r	   r
   r   r   r   r   r   Útorchgen.contextr   Útorchgen.modelr   r   r   r   r   r   r   Útorchgen.utilsr   Úcollections.abcr   r   r   r\   rs   rŽ   r�   rz   ry   r™   r    r¹   r®   r×   rð   rô   rZ   r)   r'   Ú<module>r     sÑ  ðÝ "å !Ý  ç "Ð "Ý ,÷
÷ 
õ 
õ 2÷÷ ñ õ &ñ Ý(å3ñ@ �$Ô÷Pð Pó ðPñD �$Ô÷	]ð 	]ó ð	]ó<ðS/ØðS/àDóS/ñl �$Ô÷ð ó ðñ %ØØØ×-Ñ-ôñ
 %ØØØ×,Ñ,ôð%Ð !ð,Øð,àð,ð 3ð,ð "ð	,ð
 	ó,ð^ ò&ó ð&ñ` �$Ô÷(kð (kó ð(kðV òó ðð_Øð_àð_ð 0ð_ð "ð	_ð
 	ó_ðD ò;ó ñ;r)   