U
    hâËdî˜  ã                
   @   s¢  d dl mZ d dlZd dlZd dlmZ d dlmZ d dl	m
Z
mZ d dlmZ d dlmZ d dlmZmZ d dlmZ d d	lmZ d
dlmZ d dlmZ d dlmZmZmZ d dlm Z m!Z!m"Z" e
ƒ Z#e#j$Z$e#j%Z&e#j'Z'dd„ Z(e&e )e¡dƒdd„ ƒZ*e&e )e¡dƒdd„ ƒZ+e&e )e¡dƒdd„ ƒZ,e&e )e¡dƒdd„ ƒZ-e&e )e¡dƒdd„ ƒZ.e&e d ƒd!d"„ ƒZ/e&e d#ƒd$d%„ ƒZ0e&e d&ƒd'd(„ ƒZ1e$ej2j3ƒd)d*„ ƒZ4e$d+e!ƒd,d-„ ƒZ5e$ej6j7ej8ƒd.d/„ ƒZ9d a:d0d1„ Z;e$ej<j=ej>ej?ƒd2d3„ ƒZ@e$ej<j=ejAej?ƒe$ej<j=ejBej?ƒd4d5„ ƒƒZCe$ejDj=ej>ej?ƒd6d7„ ƒZEe$ejDj=ejAej?ƒe$ejDj=ejBej?ƒd8d9„ ƒƒZFe$ejGƒd:d;„ ƒZHe$ejIƒd<d=„ ƒZJe$ejKƒd>d?„ ƒZLe$ejMƒd@dA„ ƒZNe$ejMejOƒdBdC„ ƒZPe$ejQejOejOejOejOejOƒe$ejQejOejOejRejOejOƒe$ejQejOejOejSejOejOƒe$ejQejOejOejTejOejOƒdDdE„ ƒƒƒƒZUe$ejVejOejOejWƒdFdG„ ƒZXe$ejYejOejOƒe$ejYejOejRƒe$ejYejOejSƒe$ejYejOejTƒdHdI„ ƒƒƒƒZZe$ej[ejOejOƒe$ej[ejOejRƒe$ej[ejOejSƒe$ej[ejOejTƒdJdK„ ƒƒƒƒZ\e$ej]ƒdLdM„ ƒZ^e$ej_ƒdNdO„ ƒZ`e$ejaej?ƒdPdQ„ ƒZbe$ejcej?ej?ej?ƒdRdS„ ƒZddTdU„ ZeeejfejgƒdVdW„ ƒZheejgejfƒdXdY„ ƒZidZd[„ Zjeejfejkƒd\d]„ ƒZleejkejfƒeej>ejfƒd^d_„ ƒƒZmd`da„ Znenejojpdbƒ enejqdbƒ enejrdbƒ enejojsdcƒ enejtdcƒ enejudcƒ enejojvddƒ enejwddƒ enejxddƒ e$ejojyejfƒdedf„ ƒZze$ej{ejfƒdgdh„ ƒZ|e$ejoj}ejfƒdidj„ ƒZ~e$eejfƒdkdl„ ƒZ€e$ejoj�ejfejfejfƒdmdn„ ƒZ‚e$ejƒejfejfƒe$ej„ejfejfƒdodp„ ƒƒZ…dqZ†drds„ Z‡e$ejojˆejfejfƒe‡dtƒƒ e$ej‰ejfejfƒe‡dtƒƒ e$ejojŠejfejfƒe‡duƒƒ e$ej‹ejfejfƒe‡duƒƒ e$ejojŒejfejfƒe‡dvƒƒ e$ej�ejfejfƒe‡dvƒƒ e$ejojŽejfejfƒe‡dwƒƒ e$ej�ejfejfƒe‡dwƒƒ e$ejoj�ejfejfƒe‡dxƒƒ e$ej‘ejfejfƒe‡dxƒƒ e$ejoj’ejfejfƒe‡dyƒƒ e$ej“ejfejfƒe‡dyƒƒ dzd{„ Z”e”ejoj•d|dwƒ e”ejoj–d}dyƒ ej—d~ej˜diZ™e$ejšej—ƒe$ejšej˜ƒd€d�„ ƒƒZ›e$ejœej�ƒd‚dƒ„ ƒZže$ejœejŸƒd„d…„ ƒZ e$ej¡ej?ƒd†d‡„ ƒZ¢e$ej£ejOƒe$ej£ej�ƒdˆd‰„ ƒƒZ¤e$ej£ejRƒe$ej£ejŸƒdŠd‹„ ƒƒZ¥e$ej¦ej?ej?ej?ƒdŒd�„ ƒZ§e$e¨ejSejSƒdŽd�„ ƒZ©e$e¨ejTejSƒe$e¨ejSejTƒe$e¨ejTejTƒd�d‘„ ƒƒƒZªe$e«ejSejSƒd’d“„ ƒZ¬e$e«ejTejSƒe$e«ejSejTƒe$e«ejTejTƒd”d•„ ƒƒƒZ­e$e®ejSƒe$e®ejTƒd–d—„ ƒƒZ¯e$e®ejSejkƒe$e®ejTejkƒd˜d™„ ƒƒZ°dšd›„ Z±ej²dœ Z³dœej² Z´e$ejµejSƒe±e³ƒƒ e$ejµejTƒe±e³ƒƒ e$ej¶ejSƒe±e´ƒƒ e$ej¶ejTƒe±e´ƒƒ d�dž„ Z·dŸd „ Z¸e$ej¹jqej8ejºej?ƒe$ej¹jqej8ejBej?ƒe$ej¹jqej8ejAej?ƒe¸d¡d¢„ ƒƒƒƒZ»e$ej¹jtej8ejºej?ƒe$ej¹jtej8ejBej?ƒe$ej¹jtej8ejAej?ƒe¸d£d¤„ ƒƒƒƒZ¼e$ej¹j½ej8ejºej?ƒe$ej¹j½ej8ejBej?ƒe$ej¹j½ej8ejAej?ƒe¸d¥d¦„ ƒƒƒƒZ¾e$ej¹j¿ej8ejºej?ƒe$ej¹j¿ej8ejBej?ƒe$ej¹j¿ej8ejAej?ƒe¸d§d¨„ ƒƒƒƒZÀd©dª„ ZÁeÁej¹jÂd«ƒ eÁej¹jÃd¬ƒ eÁej¹jÄd­ƒ e$ej¹jÅej8ejºej?ƒe$ej¹jÅej8ejBej?ƒe$ej¹jÅej8ejAej?ƒe¸d®d¯„ ƒƒƒƒZÆe$ej¹j¨ej8ejºej?ƒe$ej¹j¨ej8ejAej?ƒe$ej¹j¨ej8ejBej?ƒe¸d°d±„ ƒƒƒƒZÇe$ej¹j«ej8ejºej?ƒe$ej¹j«ej8ejAej?ƒe$ej¹j«ej8ejBej?ƒe¸d²d³„ ƒƒƒƒZÈe$ej¹jÉej8ejºej?ƒe$ej¹jÉej8ejAej?ƒe$ej¹jÉej8ejBej?ƒe¸d´dµ„ ƒƒƒƒZÊe$ej¹jËej8ejºej?ƒe$ej¹jËej8ejAej?ƒe$ej¹jËej8ejBej?ƒe¸d¶d·„ ƒƒƒƒZÌe$ej¹jÍej8ej?ej?ƒd¸d¹„ ƒZÎe$ej¹jÏej8ejºej?ej?ƒe$ej¹jÏej8ejAej?ej?ƒe$ej¹jÏej8ejBej?ej?ƒdºd»„ ƒƒƒZÐe$ejÑejÒƒd¼d½„ ƒZÓdÃd¿dÀ„ZÔe'e"ƒdÁdÂ„ ƒZÕee Ö¡ e$ƒ dS )Äé    )ÚreduceN)Úir)ÚRegistryÚ
lower_cast)Úparse_dtype)Úmodels)ÚtypesÚcgutils)Úufunc_db)Úregister_ufuncsé   )Únvvm)Úcuda)Ú	nvvmutilsÚstubsÚerrors)Údim3Ú
grid_groupÚCUDADispatcherc                 C   sB   t  | d| ¡}t  | d| ¡}t  | d| ¡}t | |||f¡S )Nz%s.xz%s.yz%s.z)r   Ú	call_sregr	   Zpack_struct)ÚbuilderÚprefixÚxÚyÚz© r   úL/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/cudaimpl.pyÚinitialize_dim3   s    r   Z	threadIdxc                 C   s
   t |dƒS )NÚtid©r   ©Úcontextr   ÚsigÚargsr   r   r   Úcuda_threadIdx    s    r$   ZblockDimc                 C   s
   t |dƒS )NZntidr   r    r   r   r   Úcuda_blockDim%   s    r%   ZblockIdxc                 C   s
   t |dƒS )NZctaidr   r    r   r   r   Úcuda_blockIdx*   s    r&   ZgridDimc                 C   s
   t |dƒS )NZnctaidr   r    r   r   r   Úcuda_gridDim/   s    r'   Úlaneidc                 C   s   t  |d¡S )Nr(   )r   r   r    r   r   r   Úcuda_laneid4   s    r)   r   c                 C   s   |  |d¡S ©Nr   ©Úextract_valuer    r   r   r   Údim3_x9   s    r-   r   c                 C   s   |  |d¡S ©Nr   r+   r    r   r   r   Údim3_y>   s    r/   r   c                 C   s   |  |d¡S )Né   r+   r    r   r   r   Údim3_zC   s    r1   c                 C   s(   |   tjd¡}|j}| t |¡|f¡S r.   )Úget_constantr   Úint32ÚmoduleÚcallr   Z declare_cudaCGGetIntrinsicHandle)r!   r   r"   r#   ÚoneÚlmodr   r   r   Úcg_this_gridH   s    þr8   zGridGroup.syncc                 C   s,   |   tjd¡}|j}| t |¡||f˜¡S r*   )r2   r   r3   r4   r5   r   Zdeclare_cudaCGSynchronize)r!   r   r"   r#   Úflagsr7   r   r   r   Úptx_sync_groupQ   s    þr:   c                 C   s   |d S r*   r   r    r   r   r   Úcuda_const_array_like\   s    r;   c                 C   s   t d7 a d | t ¡S )zÍDue to bug with NVVM invalid internalizing of shared memory in the
    PTX output.  We can't mark shared memory to be internal. We have to
    ensure unique name is generated for shared memory symbol.
    r   z{0}_{1})Ú_unique_smem_idÚformat©Únamer   r   r   Ú_get_unique_smem_idf   s    r@   c              	   C   s8   |j d j}t|j d ƒ}t| ||f|tdƒtjdd�S )Nr   r   Ú_cudapy_smemT©ÚshapeÚdtypeÚsymbol_nameÚ	addrspaceÚcan_dynsized)r#   Úliteral_valuer   Ú_generic_arrayr@   r   ÚADDRSPACE_SHARED©r!   r   r"   r#   ÚlengthrD   r   r   r   Úcuda_shared_array_integerp   s    ýrM   c              	   C   s>   dd„ |j d D ƒ}t|j d ƒ}t| |||tdƒtjdd�S )Nc                 S   s   g | ]
}|j ‘qS r   ©rH   ©Ú.0Úsr   r   r   Ú
<listcomp>}   s     z+cuda_shared_array_tuple.<locals>.<listcomp>r   r   rA   TrB   )r#   r   rI   r@   r   rJ   ©r!   r   r"   r#   rC   rD   r   r   r   Úcuda_shared_array_tuplez   s    
ýrT   c              	   C   s4   |j d j}t|j d ƒ}t| ||f|dtjdd�S )Nr   r   Ú_cudapy_lmemFrB   )r#   rH   r   rI   r   ÚADDRSPACE_LOCALrK   r   r   r   Úcuda_local_array_integer…   s    ýrW   c              	   C   s:   dd„ |j d D ƒ}t|j d ƒ}t| |||dtjdd�S )Nc                 S   s   g | ]
}|j ‘qS r   rN   rO   r   r   r   rR   ’   s     z(ptx_lmem_alloc_array.<locals>.<listcomp>r   r   rU   FrB   )r#   r   rI   r   rV   rS   r   r   r   Úptx_lmem_alloc_array�   s    
ýrX   c                 C   sD   |rt ‚d}|j}t t ¡ d¡}t |||¡}| |d¡ |  ¡ S )Nzllvm.nvvm.membar.ctar   ©	ÚAssertionErrorr4   r   ÚFunctionTypeÚVoidTyper	   Úget_or_insert_functionr5   Úget_dummy_value©r!   r   r"   r#   Úfnamer7   ÚfntyÚsyncr   r   r   Úptx_threadfence_blockš   s    rc   c                 C   sD   |rt ‚d}|j}t t ¡ d¡}t |||¡}| |d¡ |  ¡ S )Nzllvm.nvvm.membar.sysr   rY   r_   r   r   r   Úptx_threadfence_system¥   s    rd   c                 C   sD   |rt ‚d}|j}t t ¡ d¡}t |||¡}| |d¡ |  ¡ S )Nzllvm.nvvm.membar.glr   rY   r_   r   r   r   Úptx_threadfence_device°   s    re   c                 C   s*   |   tjd¡}t tj¡}t| |||gƒS )Nl   ÿÿ )r2   r   r3   ÚnoneÚptx_syncwarp_mask)r!   r   r"   r#   ÚmaskZmask_sigr   r   r   Úptx_syncwarp»   s    ri   c                 C   sD   d}|j }t t ¡ t d¡f¡}t |||¡}| ||¡ |  ¡ S )Nzllvm.nvvm.bar.warp.syncé    )	r4   r   r[   r\   ÚIntTyper	   r]   r5   r^   r_   r   r   r   rg   Â   s    rg   c              
   C   sü  |\}}}}}|j d }	|	tjkr6| |t |	j¡¡}d}
|j}t t 	t d¡t d¡f¡t d¡t d¡t d¡t d¡t d¡f¡}t
 |||
¡}|	jdkrü| ||||||f¡}|	tjkrú| |d¡}| |d¡}| |t ¡ ¡}t
 |||f¡}nü| |t d¡¡}| ||  tjd¡¡}| |t d¡¡}| ||||||f¡}| ||||||f¡}| |d¡}| |d¡}| |d¡}| |t d¡¡}| |t d¡¡}| ||  tjd¡¡}| ||¡}|	tjk�rè| |t ¡ ¡}t
 |||f¡}|S )a�  
    The NVVM intrinsic for shfl only supports i32, but the cuda intrinsic
    function supports both 32 and 64 bit ints and floats, so for feature parity,
    i64, f32, and f64 are implemented. Floats by way of bitcasting the float to
    an int, then shuffling, then bitcasting back. And 64-bit values by packing
    them into 2 32bit values, shuffling thoose, and then packing back together.
    r0   zllvm.nvvm.shfl.sync.i32rj   r   r   é@   )r#   r   Úreal_domainÚbitcastr   rk   Úbitwidthr4   r[   ÚLiteralStructTyper	   r]   r5   Úfloat32r,   Ú	FloatTypeZmake_anonymous_structÚtruncZlshrr2   Úi8ÚzextZshlÚor_Úfloat64Ú
DoubleType)r!   r   r"   r#   rh   ÚmodeÚvalueÚindexÚclampZ
value_typer`   r7   ra   ÚfuncÚretÚrvÚpredZfvZvalue1Z
value_lshrZvalue2Zret1Zret2Zrv1Zrv2Zrv1_64Zrv2_64Zrv_shlr   r   r   Úptx_shfl_sync_i32Ì   sJ    

 ÿþ

r�   c                 C   s^   d}|j }t t t d¡t d¡f¡t d¡t d¡t d¡f¡}t |||¡}| ||¡S )Nzllvm.nvvm.vote.syncrj   r   )r4   r   r[   rp   rk   r	   r]   r5   )r!   r   r"   r#   r`   r7   ra   r}   r   r   r   Úptx_vote_sync  s    ÿþr‚   c                 C   s†   |\}}|j d j}|j d tjkr6| |t |¡¡}d |¡}|j}t 	t d¡t d¡t |¡f¡}	t
 ||	|¡}
| |
||f¡S )Nr   zllvm.nvvm.match.any.sync.i{}rj   )r#   ro   r   rm   rn   r   rk   r=   r4   r[   r	   r]   r5   ©r!   r   r"   r#   rh   rz   Úwidthr`   r7   ra   r}   r   r   r   Úptx_match_any_sync  s    
"r…   c                 C   s–   |\}}|j d j}|j d tjkr6| |t |¡¡}d |¡}|j}t 	t 
t d¡t d¡f¡t d¡t |¡f¡}	t ||	|¡}
| |
||f¡S )Nr   zllvm.nvvm.match.all.sync.i{}rj   )r#   ro   r   rm   rn   r   rk   r=   r4   r[   rp   r	   r]   r5   rƒ   r   r   r   Úptx_match_all_sync  s    
ÿþr†   c                 C   s,   t jt  t  d¡g ¡dddd�}| |g ¡S )Nrj   zactivemask.b32 $0;ú=rT©Zside_effect©r   Ú	InlineAsmr[   rk   r5   ©r!   r   r"   r#   Ú
activemaskr   r   r   Úptx_activemask/  s      ÿr�   c                 C   s,   t jt  t  d¡g ¡dddd�}| |g ¡S )Nrj   zmov.u32 $0, %lanemask_lt;r‡   Trˆ   r‰   r‹   r   r   r   Úptx_lanemask_lt6  s     þrŽ   c                 C   s   |  |d ¡S r*   )Zctpopr    r   r   r   Úptx_popc>  s    r�   c                 C   s
   |j |Ž S ©N)Úfmar    r   r   r   Úptx_fmaC  s    r’   c                 C   sD   dddœ}z
||  W S  t k
r>   d| › d�}t |¡‚Y nX d S )N)Zf32Úf)Zf64Úd)rj   rl   z$Conversion between float16 and floatú unsupported©ÚKeyErrorr   ZCudaLoweringError©ro   ÚtypemapÚmsgr   r   r   Úfloat16_float_ty_constraintH  s    

r›   c           	      C   sd   |j |j kr|S t|j ƒ\}}t |  |¡t d¡g¡}t |d|› d�d|› d�¡}| ||g¡S )Né   zcvt.ú.f16 $0, $1;ú=ú,h)ro   r›   r   r[   Úget_value_typerk   rŠ   r5   ©	r!   r   ÚfromtyÚtotyÚvalÚtyÚ
constraintra   Úasmr   r   r   Úfloat16_to_float_castR  s    r¨   c           	      C   sb   |j |j kr|S t|j ƒ\}}t t d¡|  |¡g¡}t |d|› d�d|› �¡}| ||g¡S )Nrœ   úcvt.rn.f16.ú $0, $1;ú=h,)ro   r›   r   r[   rk   r    rŠ   r5   r¡   r   r   r   Úfloat_to_float16_cast^  s    r¬   c                 C   sH   dddddœ}z
||  W S  t k
rB   d| › d�}t |¡‚Y nX d S )NÚcÚhÚrÚl)é   rœ   rj   rl   z"Conversion between float16 and intr•   r–   r˜   r   r   r   Úfloat16_int_constraintj  s    
r²   c           
      C   sf   |j }t|ƒ}|jrdnd}t |  |¡t d¡g¡}t |d|› |› d�d|› d�¡}	| |	|g¡S )NrQ   Úurœ   zcvt.rni.r�   rž   rŸ   )	ro   r²   Úsignedr   r[   r    rk   rŠ   r5   ©
r!   r   r¢   r£   r¤   ro   r¦   Z
signednessra   r§   r   r   r   Úfloat16_to_integer_castt  s    
þr¶   c           
      C   sd   |j }t|ƒ}|jrdnd}t t d¡|  |¡g¡}t |d|› |› d�d|› �¡}	| |	|g¡S )NrQ   r³   rœ   r©   rª   r«   )	ro   r²   r´   r   r[   rk   r    rŠ   r5   rµ   r   r   r   Úinteger_to_float16_cast�  s    
ÿþr·   c                    s    t | tjtjƒ‡ fdd„ƒ}d S )Nc                    sB   t  t  d¡t  d¡t  d¡g¡}t  |ˆ › d�d¡}| ||¡S )Nrœ   z.f16 $0,$1,$2;ú=h,h,h©r   r[   rk   rŠ   r5   ©r!   r   r"   r#   ra   r§   ©Úopr   r   Úptx_fp16_binary‘  s
    ÿz*lower_fp16_binary.<locals>.ptx_fp16_binary©Úlowerr   Úfloat16)Úfnr¼   r½   r   r»   r   Úlower_fp16_binary�  s    rÂ   ÚaddÚsubÚmulc                 C   s4   t  t  d¡t  d¡g¡}t  |dd¡}| ||¡S )Nrœ   zneg.f16 $0, $1;ú=h,hr¹   rº   r   r   r   Úptx_fp16_hneg¤  s    rÇ   c                 C   s   t | |||ƒS r�   )rÇ   r    r   r   r   Úoperator_hneg«  s    rÈ   c                 C   s4   t  t  d¡t  d¡g¡}t  |dd¡}| ||¡S )Nrœ   zabs.f16 $0, $1;rÆ   r¹   rº   r   r   r   Úptx_fp16_habs°  s    rÉ   c                 C   s   t | |||ƒS r�   )rÉ   r    r   r   r   Úoperator_habs·  s    rÊ   c                 C   sH   t  d¡t  d¡t  d¡g}t  t  d¡|¡}t  |dd¡}| ||¡S )Nrœ   zfma.rn.f16 $0,$1,$2,$3;z=h,h,h,h)r   rk   r[   rŠ   r5   )r!   r   r"   r#   Zargtysra   r§   r   r   r   Úptx_hfma¼  s    rË   c                 C   s   dd„ }|   ||||¡S )Nc                 S   s   t j | |¡S r�   )r   Úfp16Zhdiv)r   r   r   r   r   Úfp16_divÇ  s    zfp16_div_impl.<locals>.fp16_div©Zcompile_internal)r!   r   r"   r#   rÍ   r   r   r   Úfp16_div_implÄ  s    rÏ   z’{{
          .reg .pred __$$f16_cmp_tmp;
          setp.{op}.f16 __$$f16_cmp_tmp, $1, $2;
          selp.u16 $0, 1, 0, __$$f16_cmp_tmp;
        }}c                    s   ‡ fdd„}|S )Nc           	         sr   t  t  d¡t  d¡t  d¡g¡}t  |tjˆ d�d¡}| ||¡}|  tj	d¡}| 
|t  d¡¡}| d||¡S )Nrœ   r»   r¸   r   z!=)r   r[   rk   rŠ   Ú	_fp16_cmpr=   r5   r2   r   Zint16rn   Zicmp_unsigned)	r!   r   r"   r#   ra   r§   ÚresultÚzeroZ
int_resultr»   r   r   Úptx_fp16_comparisonÕ  s    "z*_gen_fp16_cmp.<locals>.ptx_fp16_comparisonr   )r¼   rÓ   r   r»   r   Ú_gen_fp16_cmpÔ  s    rÔ   ÚeqÚneÚgeÚgtÚleÚltc                    s    t | tjtjƒ‡ fdd„ƒ}d S )Nc                    s(   t ˆ ƒ| |||ƒ}| ||d |d ¡S )Nr   r   )rÔ   Úselect)r!   r   r"   r#   Úchoicer»   r   r   Úptx_fp16_minmaxï  s    z*lower_fp16_minmax.<locals>.ptx_fp16_minmaxr¾   )rÁ   r`   r¼   rÝ   r   r»   r   Úlower_fp16_minmaxî  s    rÞ   ÚmaxÚminZ
__nv_cbrtfZ	__nv_cbrtc           
      C   sF   |j }t| }|  |¡}|j}t ||g¡}t |||¡}	| |	|¡S r�   )	Úreturn_typeÚ
cbrt_funcsr    r4   r   r[   r	   r]   r5   )
r!   r   r"   r#   r¥   r`   Zftyr7   ra   rÁ   r   r   r   Úptx_cbrt  s    
rã   c              	   C   s2   t  |jt t d¡t d¡f¡d¡}| ||¡S )Nrj   Z	__nv_brev©r	   r]   r4   r   r[   rk   r5   ©r!   r   r"   r#   rÁ   r   r   r   Úptx_brev_u4  s    ýræ   c              	   C   s2   t  |jt t d¡t d¡f¡d¡}| ||¡S )Nrl   Z__nv_brevllrä   rå   r   r   r   Úptx_brev_u8  s    ýrç   c                 C   s   |  |d |  tjd¡¡S r*   )Zctlzr2   r   Úbooleanr    r   r   r   Úptx_clz'  s    þré   c              	   C   s2   t  |jt t d¡t d¡f¡d¡}| ||¡S )Nrj   Z__nv_ffsrä   rå   r   r   r   Ú
ptx_ffs_32.  s    ýrê   c              	   C   s2   t  |jt t d¡t d¡f¡d¡}| ||¡S )Nrj   rl   Z
__nv_ffsllrä   rå   r   r   r   Ú
ptx_ffs_648  s    ýrë   c                 C   s   |\}}}|  |||¡S r�   )rÛ   )r!   r   r"   r#   ÚtestÚaÚbr   r   r   Úptx_selpB  s    
rï   c              	   C   s4   t  |jt t ¡ t ¡ t ¡ f¡d¡}| ||¡S )NZ
__nv_fmaxf©r	   r]   r4   r   r[   rr   r5   rå   r   r   r   Ú
ptx_max_f4H  s    þûrñ   c              
   C   sh   t  |jt t ¡ t ¡ t ¡ f¡d¡}| ||  ||d |jd t	j
¡|  ||d |jd t	j
¡g¡S )NZ	__nv_fmaxr   r   ©r	   r]   r4   r   r[   rx   r5   Úcastr#   r   Údoublerå   r   r   r   Ú
ptx_max_f8S  s    þûþrõ   c              	   C   s4   t  |jt t ¡ t ¡ t ¡ f¡d¡}| ||¡S )NZ
__nv_fminfrð   rå   r   r   r   Ú
ptx_min_f4d  s    þûrö   c              
   C   sh   t  |jt t ¡ t ¡ t ¡ f¡d¡}| ||  ||d |jd t	j
¡|  ||d |jd t	j
¡g¡S )NZ	__nv_fminr   r   rò   rå   r   r   r   Ú
ptx_min_f8o  s    þûþr÷   c              	   C   sJ   t  |jt t d¡t ¡ f¡d¡}| ||  ||d |j	d t
j¡g¡S )Nrl   Z__nv_llrintr   )r	   r]   r4   r   r[   rk   rx   r5   ró   r#   r   rô   rå   r   r   r   Ú	ptx_round€  s    þûÿrø   c                 C   s   dd„ }|   ||||¡S )Nc                 S   sÀ   t  | ¡st  | ¡r| S |dkrb|dkr:d|d  }d}nd| }d}| | | }t  |¡rt| S nd|  }| | }t|ƒ}t  || ¡dkrždt|d ƒ }|dkr´|| | }n||9 }|S )Nr   é   g      $@g’ÕMÏð€Dg      ð?g      à?g       @)ÚmathÚisinfÚisnanÚroundÚfabs)r   ÚndigitsZpow1Zpow2r   r   r   r   r   Úround_ndigits•  s(    

z$round_to_impl.<locals>.round_ndigitsrÎ   )r!   r   r"   r#   r   r   r   r   Úround_to_impl’  s    !r  c                    s   ‡ fdd„}|S )Nc                    s$   |j \}|  |ˆ ¡}| ||d ¡S r*   )r#   r2   Zfmul)r!   r   r"   r#   ZargtyÚfactor©Úconstr   r   Úimplº  s    zgen_deg_rad.<locals>.implr   )r  r  r   r  r   Úgen_deg_rad¹  s    r  g     €f@c                    s˜   |t jkr t j|dd�}|g}ntjˆ |t|ƒd�}‡ ‡fdd„t||ƒD ƒ}|j}||krltd||f ƒ‚|j	t|ƒkr�td|j	t|ƒf ƒ‚||fS )z4
    Convert integer indices into tuple of intp
    r   )rD   Úcount)r  c                    s"   g | ]\}}ˆ  ˆ ||tj¡‘qS r   )ró   r   Úintp)rP   ÚtÚi©r   r!   r   r   rR   Ò  s   ÿz&_normalize_indices.<locals>.<listcomp>zexpect %s but got %sz#indexing %d-D array with %d-D index)
r   Zinteger_domainÚUniTupler	   Zunpack_tupleÚlenÚziprD   Ú	TypeErrorÚndim)r!   r   ÚindtyÚindsÚarytyÚvaltyÚindicesrD   r   r  r   Ú_normalize_indicesÉ  s    
ÿÿr  c                    s   ‡ fdd„}|S )Nc                    sj   |j \}}}|\}}}	|j}
t| |||||ƒ\}}|  |¡| ||ƒ}tj| ||||dd�}ˆ | ||
||	ƒS )NT©Z
wraparound)r#   rD   r  Ú
make_arrayr	   Úget_item_pointer)r!   r   r"   r#   r  r  r  Úaryr  r¤   rD   r  ÚlaryÚptr©Údispatch_fnr   r   Úimpá  s    

 ÿÿz_atomic_dispatcher.<locals>.impr   )r  r  r   r  r   Ú_atomic_dispatcherà  s    r   c                 C   s`   |t jkr&|j}| t |¡||f¡S |t jkrL|j}| t |¡||f¡S | d||d¡S d S )NrÃ   Ú	monotonic)	r   rq   r4   r5   r   Zdeclare_atomic_add_float32rw   Zdeclare_atomic_add_float64Ú
atomic_rmw©r!   r   rD   r  r¤   r7   r   r   r   Úptx_atomic_add_tupleò  s    
ÿ
ÿr$  c                 C   s`   |t jkr&|j}| t |¡||f¡S |t jkrL|j}| t |¡||f¡S | d||d¡S d S )NrÄ   r!  )	r   rq   r4   r5   r   Zdeclare_atomic_sub_float32rw   Zdeclare_atomic_sub_float64r"  r#  r   r   r   Úptx_atomic_sub  s    
ÿ
ÿr%  c                 C   sP   |t jjkr<|j}|j}ttd|› �ƒ}| ||ƒ||f¡S td|› d�ƒ‚d S )NZdeclare_atomic_inc_intzUnimplemented atomic inc with ú array©	r   ÚcudadeclZunsigned_int_numba_typesro   r4   Úgetattrr   r5   r  ©r!   r   rD   r  r¤   Úbwr7   rÁ   r   r   r   Úptx_atomic_inc  s    r,  c                 C   sP   |t jjkr<|j}|j}ttd|› �ƒ}| ||ƒ||f¡S td|› d�ƒ‚d S )NZdeclare_atomic_dec_intzUnimplemented atomic dec with r&  r'  r*  r   r   r   Úptx_atomic_dec"  s    r-  c                    s@   t ‡ fdd„ƒ}tjtjtjfD ]}t| tj|tjƒ|ƒ q d S )Nc                    s6   |t jjkr| ˆ ||d¡S tdˆ › d|› d�ƒ‚d S )Nr!  zUnimplemented atomic z with r&  ©r   r(  Úinteger_numba_typesr"  r  ©r!   r   rD   r  r¤   r»   r   r   Úimpl_ptx_atomic1  s    z+ptx_atomic_bitwise.<locals>.impl_ptx_atomic)r   r   r  r  ÚTupler¿   ÚArrayÚAny)Zstubr¼   r1  r¥   r   r»   r   Úptx_atomic_bitwise0  s    r5  ÚandÚorÚxorc                 C   s0   |t jjkr| d||d¡S td|› d�ƒ‚d S )NZxchgr!  zUnimplemented atomic exch with r&  r.  r0  r   r   r   Úptx_atomic_exchA  s    r9  c                 C   sš   |j }|tjkr&| t |¡||f¡S |tjkrF| t |¡||f¡S |tjtj	fkrh|j
d||dd�S |tjtjfkrŠ|j
d||dd�S td| ƒ‚d S ©Nrß   r!  ©ZorderingZumaxz&Unimplemented atomic max with %s array)r4   r   rw   r5   r   Zdeclare_atomic_max_float64rq   Zdeclare_atomic_max_float32r3   Úint64r"  Úuint32Úuint64r  r#  r   r   r   Úptx_atomic_maxL  s    
ÿ
ÿr?  c                 C   sš   |j }|tjkr&| t |¡||f¡S |tjkrF| t |¡||f¡S |tjtj	fkrh|j
d||dd�S |tjtjfkrŠ|j
d||dd�S td| ƒ‚d S ©Nrà   r!  r;  Zuminz&Unimplemented atomic min with %s array)r4   r   rw   r5   r   Zdeclare_atomic_min_float64rq   Zdeclare_atomic_min_float32r3   r<  r"  r=  r>  r  r#  r   r   r   Úptx_atomic_min`  s    
ÿ
ÿrA  c                 C   sš   |j }|tjkr&| t |¡||f¡S |tjkrF| t |¡||f¡S |tjtj	fkrh|j
d||dd�S |tjtjfkrŠ|j
d||dd�S td| ƒ‚d S r:  )r4   r   rw   r5   r   Zdeclare_atomic_nanmax_float64rq   Zdeclare_atomic_nanmax_float32r3   r<  r"  r=  r>  r  r#  r   r   r   Úptx_atomic_nanmaxt  s    
ÿ
ÿrB  c                 C   sš   |j }|tjkr&| t |¡||f¡S |tjkrF| t |¡||f¡S |tjtj	fkrh|j
d||dd�S |tjtjfkrŠ|j
d||dd�S td| ƒ‚d S r@  )r4   r   rw   r5   r   Zdeclare_atomic_nanmin_float64rq   Zdeclare_atomic_nanmin_float32r3   r<  r"  r=  r>  r  r#  r   r   r   Úptx_atomic_nanminˆ  s    
ÿ
ÿrC  c                 C   sT   |  |jd tj|jd |jd ¡}|d |  tjd¡|d |d f}t| |||ƒS )Nr   r   r0   )rá   r#   r   r  r2   Úptx_atomic_casr    r   r   r   Úptx_atomic_compare_and_swapœ  s    $"rE  c                 C   sš   |j \}}}}|\}}	}
}t| |||	||ƒ\}}|  |¡| ||ƒ}tj| ||||dd�}|jtjjkrˆ|j	}|jj
}t |||||
|¡S td|j ƒ‚d S )NTr  z&Unimplemented atomic cas with %s array)r#   r  r  r	   r  rD   r   r(  r/  r4   ro   r   Zatomic_cmpxchgr  )r!   r   r"   r#   r  r  Zoldtyr  r  r  Úoldr¤   r  r  r  r7   ro   r   r   r   rD  £  s    ÿÿrD  c                 C   s@   t jt  t  ¡ t  d¡g¡dddd�}|d }| ||g¡ d S )Nrj   znanosleep.u32 $0;r¯   Trˆ   r   )r   rŠ   r[   r\   rk   r5   )r!   r   r"   r#   Ú	nanosleepÚnsr   r   r   Úptx_nanosleep»  s      ÿrI  Fc               	      sf  t tj|dƒ}|dko$|o$t|ƒdk}|dkr:|s:tdƒ‚ˆ j| }	t|tjtj	fƒpjt|	t
jƒpj|tjk}
|tjkr†|
s†td| ƒ‚ˆ  |¡}t ||¡}|tjkr¸tj|||d�}nh|j}t ||||¡}ˆ  |¡}d|d  ¡ > |_|röd|_nt |tj¡|_| |t  t !d¡¡d¡}t" #t $¡ j%¡}ˆ  |¡}| &|¡}|}g }t't(|ƒƒD ]\}}| )|¡ ||9 }�qXd	d
„ t(|ƒD ƒ}‡ fdd
„|D ƒ}|�røtj*t +t !d¡g ¡dddd�}| ,| -|g ¡t !d¡¡}ˆ  .tj/|¡}| 0||¡g}n‡ fdd
„|D ƒ}t|ƒ}tj1||dd�}ˆ  2|¡ˆ |ƒ}ˆ j3|| 4||j5j6¡||ˆ  .tj/|¡d d� | 7¡ S )Nr   r   zarray length <= 0zunsupported type: %sr>   Zexternalr±   Zgenericc                 S   s   g | ]}|‘qS r   r   rO   r   r   r   rR     s     z"_generic_array.<locals>.<listcomp>c                    s   g | ]}ˆ   tj|¡‘qS r   ©r2   r   r  rO   ©r!   r   r   rR     s     rj   zmov.u32 $0, %dynamic_smem_size;r‡   Trˆ   rl   c                    s   g | ]}ˆ   tj|¡‘qS r   rJ  rO   rK  r   r   rR     s     ÚC)rD   r  Zlayout)ÚdatarC   ÚstridesÚitemsizeZmeminfo)8r   ÚoperatorrÅ   r  Ú
ValueErrorZdata_model_managerÚ
isinstancer   ZRecordÚBooleanr   ZStructModelrÀ   Znumber_domainr  Zget_data_typer   Ú	ArrayTyper   rV   r	   Zalloca_oncer4   Zadd_global_variableZget_abi_sizeofÚ
bit_lengthÚalignÚlinkageÚConstantÚ	UndefinedZinitializerZaddrspacecastZPointerTyperk   ÚllZcreate_target_dataZNVVMZdata_layoutZget_abi_sizeÚ	enumerateÚreversedÚappendrŠ   r[   ru   r5   r2   r  Zudivr3  r  Zpopulate_arrayrn   rM  ÚtypeZ	_getvalue) r!   r   rC   rD   rE   rF   rG   Z	elemcountZdynamic_smemZ
data_modelZother_supported_typeZlldtypeZlarytyZdataptrr7   ZgvmemrV  Z
targetdatarO  Z
laststrideZrstridesr
  ZlastsizerN  ZkstridesZget_dynshared_sizeZdynsmem_sizeZ	kitemsizeZkshaper  r  r  r   rK  r   rI   Æ  sz    

ÿý


ÿ
ÿ


 þÿûrI   c                 C   s   |   ¡ S r�   )r^   )r!   r   r¥   Zpyvalr   r   r   Úcuda_dispatcher_const*  s    r_  )F)×Ú	functoolsr   rP  rú   Zllvmliter   Zllvmlite.bindingZbindingrZ  Znumba.core.imputilsr   r   Znumba.core.typing.npydeclr   Znumba.core.datamodelr   Z
numba.corer   r	   Znumba.npr
   Znumba.np.npyimplr   Zcudadrvr   Znumbar   Z
numba.cudar   r   r   Znumba.cuda.typesr   r   r   Úregistryr¿   Zlower_getattrZ
lower_attrZlower_constantr   ÚModuler$   r%   r&   r'   r)   r-   r/   r1   ZcgZ	this_gridr8   r:   r  Z
array_liker3  r;   r<   r@   ZsharedÚarrayZIntegerLiteralr4  rM   r2  r  rT   ÚlocalrW   rX   Zthreadfence_blockrc   Zthreadfence_systemrd   Zthreadfencere   Zsyncwarpri   Úi4rg   Zshfl_sync_intrinsicrt   Zf4Zf8r�   Zvote_sync_intrinsicrè   r‚   Zmatch_any_syncr…   Zmatch_all_syncr†   rŒ   r�   Zlanemask_ltrŽ   Zpopcr�   r‘   r’   r›   rÀ   ÚFloatr¨   r¬   r²   ÚIntegerr¶   r·   rÂ   rÌ   ZhaddrÃ   ÚiaddZhsubrÄ   ÚisubZhmulrÅ   ÚimulZhnegrÇ   ÚnegrÈ   ZhabsrÉ   ÚabsrÊ   ZhfmarË   ÚtruedivÚitruedivrÏ   rÐ   rÔ   ZheqrÕ   ÚhnerÖ   Zhger×   ZhgtrØ   ÚhlerÙ   ZhltrÚ   rÞ   ZhmaxZhminrq   rw   râ   Zcbrtrã   ZbrevZu4ræ   Úu8rç   Zclzré   Zffsrê   rë   Zselprï   rß   rñ   rõ   rà   rö   r÷   rý   rø   r  r  ÚpiZ_deg2radZ_rad2degÚradiansÚdegreesr  r   Zatomicr  r$  r%  Úincr,  Údecr-  r5  Úand_rv   r8  Zexchr9  r?  rA  ZnanmaxrB  ZnanminrC  Zcompare_and_swaprE  ZcasrD  rG  r=  rI  rI   r_  Z
get_ufuncsr   r   r   r   Ú<module>   s  














		
		







	ÿÿÿÿ.










	





  þ










%





 ÿ
d
