U
    •Ê¹jˆE  ã                   @   sú  d dl mZ d dlmZmZmZmZmZmZ d dl	m
  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 d dlmZmZmZmZm Z m!Z!m"Z" d d	l#m$Z$ ed
d�G dd„ dƒƒZ%ed
d�G dd„ dƒƒZ&e e'dœdd„Z(e eee!ee"e%f f e)f dœdd„Z*ed
d�G dd„ dƒƒZ+e+d de"j,d�e+dde"j-d�gZ.e e!ee"e%f ee e)dœdd„Z/ee e)dœdd„ƒZ0ed
d�G d d!„ d!ƒƒZ1ee e)dœd"d#„ƒZ2e e!ee"e&f ee e)dœd$d%„Z3ee e)dœd&d'„ƒZ4dS )(é    )Ú	dataclass)ÚDictÚListÚOptionalÚSequenceÚTupleÚUnionN)Ú	translate)	Ú	BaseCTypeÚBindingÚCTypeÚExprÚ
NamedCTypeÚopmath_tÚscalar_tÚStructuredImplSignatureÚVectorizedCType)ÚUfunctorBindings)Úwith_native_function)ÚArgumentÚBaseTyÚBaseTypeÚDispatchKeyÚNativeFunctionsGroupÚ
ScalarTypeÚUfuncKey)Ú
OrderedSetT)Úfrozenc                   @   s‚   e Zd ZU eed< ee ed< eed< edœdd„Z	e
e dœdd„Zedœd	d
„Zedœdd„Zedœdd„Zedœdd„ZdS )ÚUfunctorSignatureÚgÚscalar_tensor_idxÚname©Úreturnc                 C   s   t j| j| jtd�S )N)r    r   )ÚufuncZufunctor_argumentsr   r    r   ©Úself© r'   úR/var/www/html/bwave/python/venv/lib/python3.8/site-packages/torchgen/dest/ufunc.pyÚ	argumentsA   s
      ÿzUfunctorSignature.argumentsc                 C   s   dd„ |   ¡ jD ƒS )Nc                 S   s   g | ]}|  |j› d �¡‘qS )Ú_)Úrenamer!   ©Ú.0Úbr'   r'   r(   Ú
<listcomp>H   s     z,UfunctorSignature.fields.<locals>.<listcomp>)r)   Úctorr%   r'   r'   r(   ÚfieldsF   s    zUfunctorSignature.fieldsc                 C   s   t tƒS ©N)r
   r   r%   r'   r'   r(   Úreturns_typeJ   s    zUfunctorSignature.returns_typec                 C   s   d  dd„ |  ¡ D ƒ¡S )NÚ
c                 s   s"   | ]}|j › d |j› d�V  qdS )ú ú;N)Útyper!   )r-   Úfr'   r'   r(   Ú	<genexpr>P   s     z0UfunctorSignature.decl_fields.<locals>.<genexpr>)Újoinr1   r%   r'   r'   r(   Údecl_fieldsO   s    zUfunctorSignature.decl_fieldsc                 C   sL   d  dd„ |  ¡ jD ƒ¡}d  dd„ |  ¡ jD ƒ¡}| j› d|› d|› d�S )Nú, c                 s   s   | ]}|  ¡ V  qd S r2   ©Údecl©r-   Úar'   r'   r(   r9   S   s     z5UfunctorSignature.inline_defn_ctor.<locals>.<genexpr>c                 s   s"   | ]}|j › d |j › d�V  qdS )z_(ú)N©r!   r?   r'   r'   r(   r9   V   s     ú(z) : z {})r:   r)   r0   r!   )r&   Úargs_strZinit_strr'   r'   r(   Úinline_defn_ctorR   s    z"UfunctorSignature.inline_defn_ctorc                 C   s2   d  dd„ |  ¡ jD ƒ¡}|  ¡  ¡ › d|› d�S )Nr<   c                 s   s   | ]}|  ¡ V  qd S r2   r=   r?   r'   r'   r(   r9   Z   s     z/UfunctorSignature.decl_apply.<locals>.<genexpr>z operator()(z) const)r:   r)   Úapplyr3   Zcpp_type)r&   rD   r'   r'   r(   Ú
decl_applyY   s    zUfunctorSignature.decl_applyN)Ú__name__Ú
__module__Ú__qualname__r   Ú__annotations__r   ÚintÚstrr   r)   r   r   r1   r   r3   r;   rE   rG   r'   r'   r'   r(   r   ;   s   
r   c                   @   sT   e Zd ZU eed< eed< eed< ee dœdd„Z	e
eeef  edœdd	„Zd
S )ÚUfuncSignaturer   r!   Ú	compute_tr"   c                 C   s   t j| j| jd�S )N)rO   )r$   Zufunc_argumentsr   rO   r%   r'   r'   r(   r)   d   s    zUfuncSignature.arguments©Úctxr#   c              	   C   s,   | j › dd dd„ t||  ¡ ƒD ƒ¡› d�S )NrC   r<   c                 s   s   | ]}|j V  qd S r2   ©Úexprr?   r'   r'   r(   r9   h   s     z&UfuncSignature.call.<locals>.<genexpr>rA   ©r!   r:   r	   r)   ©r&   rQ   r'   r'   r(   Úcallg   s    zUfuncSignature.callN)rH   rI   rJ   r   rK   rM   r   r   r   r)   r   r   r   rV   r'   r'   r'   r(   rN   ^   s
   
rN   )r   r#   c                 C   s"   t dd„ | jjjjD ƒƒ}|dkS )Nc                 s   s   | ]}|j  ¡ rd V  qdS )é   N)r7   Úis_tensor_liker?   r'   r'   r(   r9   ~   s    
 z<eligible_for_binary_scalar_specialization.<locals>.<genexpr>é   )ÚsumÚ
functionalÚfuncr)   Úflat_non_out)r   Znum_tensorsr'   r'   r(   Ú)eligible_for_binary_scalar_specialization}   s    
ÿr^   c                 C   sà  i }g }| j j}tjdtjdtjd i}t| ƒr@tjtjtjg}n2tjg}tjtjfD ]}||ksTtd|› d�ƒ‚qT|D �]X}||kr¾t| || || j	d�}|| j
D ]}|| |i ¡|< q¦qvd }	tƒ }
tjtjfD ]H}||krâqÔ|	d krö|| j	}	n|	|| j	k�stdƒ‚|
|| j
O }
qÔ|	d k	�s,t‚|› d|	› �}t| || |d�}|
D ]}|| |i ¡|< �qPt| d|	› �ttƒd	�}| ¡ | ¡ j }| d
|j	› d| ¡ › d| ¡ › d| ¡ › d| |¡› d�¡ qv|d |¡fS )NrW   r   zcannot use z on non-binary function)r    r!   z0ScalarOnly and Generic must have same ufunc namer*   úufunc::©r!   rO   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ÚCUDAFunctorr^   ÚAssertionErrorr   r!   Úsupported_dtypesÚ
setdefaultr   Ú
ScalarOnlyÚGenericrN   r
   r   r1   r)   rF   Úappendr;   rE   rG   rV   r:   )r   Úufunctor_sigsÚ	ufunctorsÚloopsZscalar_tensor_idx_lookupÚkeysÚkÚufunctor_sigÚdtypeZ
ufunc_namerg   Úlkr!   Z	ufunc_sigZ	apply_ctxr'   r'   r(   Úcompute_ufunc_cuda_functors„   sŠ       ýý
  ÿ	ÿþ  ÿ  ÿþüûúùÿrt   c                   @   s&   e Zd ZU eed< eed< eed< dS )Ú BinaryScalarSpecializationConfigÚ
scalar_idxÚctor_tensorÚ	ufunc_keyN)rH   rI   rJ   rL   rK   rM   r   r'   r'   r'   r(   ru   ×   s   
ru   r&   )rv   rw   rx   rW   Úother)r   rr   Úinner_loopsÚ
parent_ctxr#   c           
      C   sî   d}|d7 }t D ]–}|j|kr q||j }|jd }t|ƒ}| td|› d�t|jtt	ƒƒd�¡ d 
dd	„ t|| ¡ jƒD ƒ¡}	|d
|› d|j› d|	› d|› d�	7 }q|tj }d 
dd	„ t|| ¡ jƒD ƒ¡}	|d|j› d|	› d�7 }|S )Nz+using opmath_t = at::opmath_type<scalar_t>;zif (false) {}
rW   ziter.scalar_value<opmath_t>(rA   )rS   r7   r<   c                 s   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9     s    z0compute_ufunc_cuda_dtype_body.<locals>.<genexpr>zelse if (iter.is_cpu_scalar(z)) {
  z<scalar_t> ufunctor(z);
  iter.remove_operand(z");
  gpu_kernel(iter, ufunctor);
}c                 s   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9     s    z
else {
  gpu_kernel(iter, z<scalar_t>(z
));
}
    )Ú!BinaryScalarSpecializationConfigsrx   rv   Úlistrk   r   r   rw   r
   r   r:   r	   r)   r0   r!   r   re   )
r   rr   rz   r{   ÚbodyÚconfigrq   rv   rQ   Zufunctor_ctor_exprs_strr'   r'   r(   Úcompute_ufunc_cuda_dtype_bodyì   sH    



þÿ
ÿÿþþÿ

ÿþþ
r€   c           	      C   sÐ   t | ƒ\}}t| t | tj¡ƒ}g }| ¡ D ].\}}| d|› dt| ||| 	¡ ƒ› d�¡ q,d 
|¡}t| ƒ}d|› d| ¡ › d| ¡ › d| ¡ › d|j› d	|› d
|j› d|j› d| ¡ › d| | 	¡ ¡› d�S )Nú"
AT_DISPATCH_CASE(at::ScalarType::ú,
  [&]() {
    ú
  }
)
r4   z

ú;
ú;

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

ú {
  ú;
}
)rt   r   r$   Úkernel_namer   ÚCUDAÚitemsrk   r€   r)   r:   ÚStubSignatureÚ	type_defnÚdispatch_declÚkernel_defnr!   ÚdefnÚdirect_call)	r   rl   rm   ÚsigÚdtype_casesrr   Úinner_ufunc_sigsÚdtype_cases_strÚstub_sigr'   r'   r(   Úcompute_ufunc_cuda  sD    ÿýÿ

ÿýüúùøõõóòr™   c                   @   sÌ   e Zd ZU eed< eedœdd„ƒZeedœdd„ƒZeedœdd„ƒZ	e
e dœd	d
„Zedœdd„Zedœdd„Zedœdd„Zedœdd„Zedœdd„Zee edœdd„Zee edœdd„ZdS )rŽ   r   r"   c                 C   s   t | jjjjjƒ› d�S )NZ_stub©rM   r   r[   r\   r!   r%   r'   r'   r(   r!   P  s    zStubSignature.namec                 C   s   t | jjjjjƒ› d�S )NZ_kernelrš   r%   r'   r'   r(   r‹   T  s    zStubSignature.kernel_namec                 C   s   t | jjjjjƒ› d�S )NÚ_fnrš   r%   r'   r'   r(   Ú	type_nameX  s    zStubSignature.type_namec                 C   s   t  | j¡S r2   )r$   Zstub_argumentsr   r%   r'   r'   r(   r)   \  s    zStubSignature.argumentsc                 C   s$   |   ¡ }dd dd„ |D ƒ¡› d�S )Nzvoid(*)(TensorIteratorBase&, r<   c                 s   s   | ]}|j V  qd S r2   )r7   r?   r'   r'   r(   r9   a  s     z%StubSignature.type.<locals>.<genexpr>rA   )r)   r:   )r&   Zcpp_argsr'   r'   r(   r7   _  s    zStubSignature.typec                 C   s   d| j › d| j› d�S )NzDECLARE_DISPATCH(r<   rA   )rœ   r!   r%   r'   r'   r(   r�   c  s    zStubSignature.dispatch_declc                 C   s   d| j › d�S )NzDEFINE_DISPATCH(rA   rB   r%   r'   r'   r(   Údispatch_defnf  s    zStubSignature.dispatch_defnc                 C   s(   d| j › dd dd„ |  ¡ D ƒ¡› d�S )Nzvoid z(TensorIteratorBase& iter, r<   c                 s   s   | ]}|  ¡ V  qd S r2   )r’   r?   r'   r'   r(   r9   j  s     z,StubSignature.kernel_defn.<locals>.<genexpr>rA   )r‹   r:   r)   r%   r'   r'   r(   r‘   i  s    zStubSignature.kernel_defnc                 C   s   d| j › d|  ¡ › �S )Nzusing ú = )rœ   r7   r%   r'   r'   r(   r�   l  s    zStubSignature.type_defnrP   c              	   C   s,   | j › dd dd„ t||  ¡ ƒD ƒ¡› d�S )Nz(device_type(), *this, r<   c                 s   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9   q  s     z%StubSignature.call.<locals>.<genexpr>rA   rT   rU   r'   r'   r(   rV   p  s    zStubSignature.callc              	   C   s,   | j › dd dd„ t||  ¡ ƒD ƒ¡› d�S )Nz(*this, r<   c                 s   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9   u  s     z,StubSignature.direct_call.<locals>.<genexpr>rA   )r‹   r:   r	   r)   rU   r'   r'   r(   r“   t  s    zStubSignature.direct_callN)rH   rI   rJ   r   rK   ÚpropertyrM   r!   r‹   rœ   r   r   r)   r7   r�   r�   r‘   r�   r   rV   r“   r'   r'   r'   r(   rŽ   L  s   
rŽ   c                 C   sZ   t | ƒ}t| t | tj¡ƒ}d| ¡ › d| ¡ › d| ¡ › d| 	¡ › d| 
| ¡ ¡› d�S )Nr4   r„   r…   r‰   rŠ   )rŽ   r   r$   r‹   r   ZCPUr�   r�   r�   r’   rV   r)   )r   r˜   r”   r'   r'   r(   Úcompute_ufunc_cpux  s    ÿþýûúr    c                    s¦  t j|ks t|› d| ¡ › �ƒ‚| ¡ t jt jhks8t‚|t j }d }t j|krZ|t j }g }g ‰ |D ]f}t|jtƒrŠ|jjt	t
jƒkrŠqf| d|j› d|j› d�¡ ˆ  td|j› �t|jjttƒƒƒ¡ qf|d k	�rL|D ]n}t|jtƒ�r|jjt	t
jƒk�rqÜ| d|j› d|j› d�¡ ˆ  td	|j› �t|jjtttƒƒƒƒ¡ qÜg }g }	| jjjjD ]~}
|
j ¡ �st�q`|
jt	t
jƒk�sŠt‚| t|
jt|
jttƒƒ|
d
�¡ |d k	�r`|	 t|
jt|
jtttƒƒƒ|
d
�¡ �q`tt ttttf  dœ‡ fdd„}d |¡}|d k	�rpd|› dd dd„ |D ƒ¡› d| ||ƒ¡› dd dd„ |	D ƒ¡› d| ||	ƒ¡› d�S d|› dd dd„ |D ƒ¡› d| ||ƒ¡› d�S d S )Nr<   zauto _s_rž   z.to<scalar_t>();Z_s_zauto _v_z$ = at::vec::Vectorized<scalar_t>(_s_z);Z_v_)r!   ÚnctypeÚargument)r.   r#   c                    s   g }|  ˆ ¡ |  | ¡ |S r2   )Úextend)r.   Úr©rQ   r'   r(   Úwith_ctxÍ  s    

z.compute_ufunc_cpu_dtype_body.<locals>.with_ctxr4   z
cpu_kernel_vec(iter,
  [=](c                 s   s   | ]}|  ¡ V  qd S r2   r=   r,   r'   r'   r(   r9   Ø  s     z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>z) { return z; },
  [=](c                 s   s   | ]}|  ¡ V  qd S r2   r=   r,   r'   r'   r(   r9   Ù  s     z; }
);
z
cpu_kernel(iter,
  [=](c                 s   s   | ]}|  ¡ V  qd S r2   r=   r,   r'   r'   r(   r9   à  s     ) r   Ú	CPUScalarrf   ro   Ú	CPUVectorÚ
isinstancer¢   r   r7   r   r   ZScalarrk   r!   r   r   r¡   r
   r   r   r[   r\   r)   r]   rX   ZTensorr   r   r   r   r:   rV   )r   rr   rz   r{   Zscalar_loopZvec_loopr~   r.   Zscalar_bindingsZvec_bindingsr@   r¦   Zbody_strr'   r¥   r(   Úcompute_ufunc_cpu_dtype_bodyˆ  s�     


	ÿ(
ÿÿ
þÿ	ýÿ
ýÿ$

ÿýýüüÿýýrª   c                 C   sx  t | ƒ}| jj}i }tjtjfD ]Î}g }||kr:| |¡ tj|krZ|tjkrZ| tj¡ tj|krp| tj¡ |D ]x}|| j	D ]h}|tjkršt
tƒ}n|tjkr²tt
tƒƒ}ntƒ ‚| |i ¡}	||	kr‚t| d|| j› �|d�|	|< q‚qtq g }
| ¡ D ].\}}	|
 d|› dt| ||	| ¡ ƒ› d�¡ qüd |
¡}d| ¡ › d|j› d	|› d
| ¡ › d| ¡ › d|j› d|j› d�S )Nr_   r`   r�   r‚   rƒ   r4   z
namespace {

r†   r‡   z#
  );
}

} // anonymous namespace

r„   z;
REGISTER_DISPATCH(rˆ   z);
)rŽ   ra   rb   r   r§   r¨   rk   ri   rj   rg   r
   r   r   rf   rh   rN   r!   r�   rª   r)   r:   r‘   r�   r�   r‹   )r   r˜   rn   Z
ufunc_sigsrp   Zlksrs   rr   rO   r–   r•   r—   r'   r'   r(   Úcompute_ufunc_cpu_kernelå  sd    




  ÿÿýÿ

ýüûõôüór«   )5Údataclassesr   Útypingr   r   r   r   r   r   Ztorchgen.api.ufuncÚapir$   Ztorchgen.api.translater	   Ztorchgen.api.typesr
   r   r   r   r   r   r   r   r   r   Ztorchgen.contextr   Ztorchgen.modelr   r   r   r   r   r   r   Ztorchgen.utilsr   r   rN   Úboolr^   rM   rt   ru   rd   rc   r|   r€   r™   rŽ   r    rª   r«   r'   r'   r'   r(   Ú<module>   sb    ,$	"þSýýú
û/0+
û]