o
    8ήcE                  
   @   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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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&eddd  Z,e#ed!d"d# Z-e#ed$d%d& Z.e#ed'd(d) Z/e!ej0j1d*d+ Z2e!d,ed-d. Z3e!ej4ej5d/d0 Z6d1d2 Z7e!ej8ej5d3d4 Z9e!ej:j;ej<d5d6 Z=d a>d7d8 Z?e!ej@jAejBejCd9d: ZDe!ej@jAejEejCe!ej@jAejFejCd;d< ZGe!ejHjAejBejCd=d> ZIe!ejHjAejEejCe!ejHjAejFejCd?d@ ZJe!ejKdAdB ZLe!ejMejNdCdD ZOe!ejPejNdEdF ZQe!ejRejNdGdH ZSe!ejTdIdJ ZUe!ejVdKdL ZWe!ejXdMdN ZYe!ejZdOdP Z[e!ejZejNdQdR Z\e!ej]ejNejNejNejNejNe!ej]ejNejNej^ejNejNe!ej]ejNejNej_ejNejNe!ej]ejNejNej`ejNejNdSdT Zae!ejbejNejNejcdUdV Zde!ejeejNejNe!ejeejNej^e!ejeejNej_e!ejeejNej`dWdX Zfe!ejgejNejNe!ejgejNej^e!ejgejNej_e!ejgejNej`dYdZ Zhe!ejid[d\ Zje!ejkd]d^ Zle!ejmejCd_d` Zne!ejoejCejCejCdadb Zpdcdd Zqeejrejsdedf Zteejsejrdgdh Zudidj Zveejrejwdkdl ZxeejwejreejBejrdmdn Zydodp Zzezej{j|dq ezej{j}dr ezej{j~ds e!ej{jejrdtdu Ze!ej{jejrdvdw Ze!ej{jejrejrejrdxdy ZdzZd{d| Ze!ej{jejrejred} e!ej{jejrejred~ e!ej{jejrejred e!ej{jejrejred e!ej{jejrejred e!ej{jejrejred dd Zeej{jdd eej{jdd ejdejdiZe!ejeje!ejejdd Ze!ejejdd Ze!ejejdd Ze!ejejCdd Ze!ejejNe!ejejdd Ze!ejej^e!ejejdd Ze!ejejCejCejCdd Ze!eej_ej_dd Ze!eej`ej_e!eej_ej`e!eej`ej`dd Ze!eej_ej_dd Ze!eej`ej_e!eej_ej`e!eej`ej`dd Ze!eej_e!eej`dd Ze!eej_ejwe!eej`ejwdd Zdd Zejd Zdej Ze!ejej_ee e!ejej`ee e!ejej_ee e!ejej`ee dd Zdd Ze!ejjej<ejejCe!ejjej<ejFejCe!ejjej<ejEejCedd Ze!ejjej<ejejCe!ejjej<ejFejCe!ejjej<ejEejCedd Ze!ejjej<ejejCe!ejjej<ejFejCe!ejjej<ejEejCedd Ze!ejjej<ejejCe!ejjej<ejFejCe!ejjej<ejEejCedd Zdd Zeejjd eejjd eejjd e!ejjej<ejejCe!ejjej<ejFejCe!ejjej<ejEejCedd Ze!ejjej<ejejCe!ejjej<ejEejCe!ejjej<ejFejCedd Ze!ejjej<ejejCe!ejjej<ejEejCe!ejjej<ejFejCedd Ze!ejjej<ejejCe!ejjej<ejEejCe!ejjej<ejFejCedd Ze!ejjej<ejejCe!ejjej<ejEejCe!ejjej<ejFejCedd Ze!ejjej<ejCejCdd Ze!ejej˃ddĄ Z	dddǄZe$eddɄ ZdS )    )reduceN)ir)Registry
lower_cast)parse_dtype	signature)models)typescgutils   )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
   pack_struct)builderprefixxyz r   :/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/cudaimpl.pyinitialize_dim3   s   r   	threadIdxc                 C   
   t |dS )Ntidr   contextr   sigargsr   r   r   cuda_threadIdx      
r&   blockDimc                 C   r   )Nntidr!   r"   r   r   r   cuda_blockDim$   r'   r*   blockIdxc                 C   r   )Nctaidr!   r"   r   r   r   cuda_blockIdx)   r'   r-   gridDimc                 C   r   )Nnctaidr!   r"   r   r   r   cuda_gridDim.   r'   r0   laneidc                 C      t |dS )Nr1   r   r   r"   r   r   r   cuda_laneid3      r4   warpsizec                 C   r2   )Nr6   r3   r"   r   r   r   cuda_warpsize8   r5   r7   r   c                 C      | |dS Nr   extract_valuer"   r   r   r   dim3_x=   r5   r<   r   c                 C   r8   Nr   r:   r"   r   r   r   dim3_yB   r5   r>   r   c                 C   r8   )N   r:   r"   r   r   r   dim3_zG   r5   r@   c                 C   s(   |  tjd}|j}|t||fS r=   )get_constantr	   int32modulecallr    declare_cudaCGGetIntrinsicHandle)r#   r   r$   r%   onelmodr   r   r   cg_this_gridL   s   rH   zGridGroup.syncc                 C   s0   |  tjd}|j}|t|g ||R S r9   )rA   r	   rB   rC   rD   r   declare_cudaCGSynchronize)r#   r   r$   r%   flagsrG   r   r   r   ptx_sync_groupU   s   rK   c                 C   sR   |j }|tjkrtj|ddS t|tjr#tj||jd}t	||S t
d| )Nr   )dimz(Unexpected return type %s from cuda.grid)return_typer	   rB   r   get_global_id
isinstanceUniTuplecountr
   
pack_array
ValueError)r#   r   r$   r%   restypeidsr   r   r   	cuda_grid`   s   
rV   c                 C   s0   t | d| }t | d| }| ||S )Nzntid.znctaid.)r   r   mul)r   rL   r)   r/   r   r   r   _nthreads_for_diml   s   rX   c                 C   s   |j }t|d}|tjkr|S t|tjr:t|d}|jdkr't|||fS |jdkr:t|d}t||||fS t	d| )Nr   r   r?      r   z*Unexpected return type %s of cuda.gridsize)
rM   rX   r	   rB   rO   rP   rQ   r
   rR   rS   )r#   r   r$   r%   rT   nxnynzr   r   r   cuda_gridsizer   s   





r]   c                 C   s   |d S r9   r   r"   r   r   r   cuda_const_array_like   s   r^   c                 C   s   t d7 a d| t S )zDue 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_id   s   rc   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_arrayrc   r   ADDRSPACE_SHAREDr#   r   r$   r%   lengthrg   r   r   r   cuda_shared_array_integer   s   rp   c              	   C   s>   dd |j d D }t|j d }t| |||tdtjddS )Nc                 S      g | ]}|j qS r   rk   .0sr   r   r   
<listcomp>       z+cuda_shared_array_tuple.<locals>.<listcomp>r   r   rd   Tre   )r%   r   rl   rc   r   rm   r#   r   r$   r%   rf   rg   r   r   r   cuda_shared_array_tuple   s   
ry   c              	   C   s4   |j d j}t|j d }t| ||f|dtjddS )Nr   r   _cudapy_lmemFre   )r%   rk   r   rl   r   ADDRSPACE_LOCALrn   r   r   r   cuda_local_array_integer   s   r|   c              	   C   s:   dd |j d D }t|j d }t| |||dtjddS )Nc                 S   rq   r   rr   rs   r   r   r   rv      rw   z(ptx_lmem_alloc_array.<locals>.<listcomp>r   r   rz   Fre   )r%   r   rl   r   r{   rx   r   r   r   ptx_lmem_alloc_array   s   
r}   c                 C   D   |rJ d}|j }tt d}t|||}||d |  S )Nzllvm.nvvm.barrier0r   rC   r   FunctionTypeVoidTyper
   get_or_insert_functionrD   get_dummy_valuer#   r   r$   r%   fnamerG   fntysyncr   r   r   ptx_syncthreads      r   c                 C   >   d}|j }ttdtdf}t|||}|||S )Nzllvm.nvvm.barrier0.popc    rC   r   r   IntTyper
   r   rD   r   r   r   r   ptx_syncthreads_count   
   r   c                 C   r   )Nzllvm.nvvm.barrier0.andr   r   r   r   r   r   ptx_syncthreads_and   r   r   c                 C   r   )Nzllvm.nvvm.barrier0.orr   r   r   r   r   r   ptx_syncthreads_or   r   r   c                 C   r~   )Nzllvm.nvvm.membar.ctar   r   r   r   r   r   ptx_threadfence_block   r   r   c                 C   r~   )Nzllvm.nvvm.membar.sysr   r   r   r   r   r   ptx_threadfence_system   r   r   c                 C   r~   )Nzllvm.nvvm.membar.glr   r   r   r   r   r   ptx_threadfence_device  r   r   c                 C   s,   |  tjd}ttjtj}t| |||gS )Nl    )rA   r	   rB   r   noneptx_syncwarp_mask)r#   r   r$   r%   maskmask_sigr   r   r   ptx_syncwarp  s   r   c                 C   sD   d}|j }tt tdf}t|||}||| |  S )Nzllvm.nvvm.bar.warp.syncr   )	rC   r   r   r   r   r
   r   rD   r   r   r   r   r   r     s   r   c              
   C   s  |\}}}}}|j d }	|	tjv r||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}|S ||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.
    r?   zllvm.nvvm.shfl.sync.i32r   r   r   @   )r%   r	   real_domainbitcastr   r   bitwidthrC   r   LiteralStructTyper
   r   rD   float32r;   	FloatTypemake_anonymous_structtrunclshrrA   i8zextshlor_float64
DoubleType)r#   r   r$   r%   r   modevalueindexclamp
value_typer   rG   r   funcretrvpredfvvalue1
value_lshrvalue2ret1ret2rv1rv2rv1_64rv2_64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.syncr   r   )rC   r   r   r   r   r
   r   rD   )r#   r   r$   r%   r   rG   r   r   r   r   r   ptx_vote_syncT  s   r   c                 C   s   |\}}|j d j}|j d tjv r||t|}d|}|j}t	tdtdt|f}	t
||	|}
||
||fS )Nr   zllvm.nvvm.match.any.sync.i{}r   )r%   r   r	   r   r   r   r   r`   rC   r   r
   r   rD   r#   r   r$   r%   r   r   widthr   rG   r   r   r   r   r   ptx_match_any_sync_  s   
"r   c                 C   s   |\}}|j d j}|j d tjv r||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{}r   )r%   r   r	   r   r   r   r   r`   rC   r   r   r
   r   rD   r   r   r   r   ptx_match_all_synco  s   
r   c                 C   ,   t jt t dg dddd}||g S )Nr   zactivemask.b32 $0;=rTside_effectr   	InlineAsmr   r   rD   r#   r   r$   r%   
activemaskr   r   r   ptx_activemask  s   r   c                 C   r   )Nr   zmov.u32 $0, %lanemask_lt;r   Tr   r   r   r   r   r   ptx_lanemask_lt  s
   r   c                 C   s   | |d S r9   )ctpopr"   r   r   r   ptx_popc  s   r   c                 C   s
   |j | S N)fmar"   r   r   r   ptx_fma  r'   r   c                 C   s:   ddd}z||  W S  t y   d|  d}t|w )N)f32f)f64d)r   r   z$Conversion between float16 and float unsupportedKeyErrorr   CudaLoweringErrorr   typemapmsgr   r   r   float16_float_ty_constraint  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)r   r   r   r   get_value_typer   r   rD   	r#   r   fromtytotyvalty
constraintr   asmr   r   r   float16_to_float_cast  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,)r   r   r   r   r   r   r   rD   r   r   r   r   float_to_float16_cast  s   r   c                 C   s>   ddddd}z||  W S  t y   d|  d}t|w )Nchrl)   r   r   r   z"Conversion between float16 and intr   r   r   r   r   r   float16_int_constraint  s   

r   c           
      C   sf   |j }t|}|jrdnd}t| |tdg}t|d| | dd| d}	||	|gS )Nru   ur   zcvt.rni.r   r   r   )	r   r   signedr   r   r   r   r   rD   
r#   r   r   r   r   r   r   
signednessr   r   r   r   r   float16_to_integer_cast  s   
r   c           
      C   sd   |j }t|}|jrdnd}ttd| |g}t|d| | dd| }	||	|gS )Nru   r   r   r   r   r   )	r   r   r   r   r   r   r   r   rD   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   r   r   rD   r#   r   r$   r%   r   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     r  addsubrW   c                 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   sL   t j dk r
d}nd}ttdtdg}t||d}|||S )N)
   r?   zand.b16 $0, $1, 0x7FFF;zabs.f16 $0, $1;r   r
  )r   runtimeget_versionr   r   r   r   rD   )r#   r   r$   r%   instr   r   r   r   r   ptx_fp16_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   r   r   r   rD   )r#   r   r$   r%   argtysr   r   r   r   r   ptx_hfma  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                        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   r   r   	_fp16_cmpr`   rD   rA   r	   int16r   icmp_unsigned)	r#   r   r$   r%   r   r   resultzero
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+  r  r%  maxmin
__nv_cbrtf	__nv_cbrtc           
      C   sF   |j }t| }| |}|j}t||g}t|||}	||	|S r   )	rM   
cbrt_funcsr   rC   r   r   r
   r   rD   )
r#   r   r$   r%   r   r   ftyrG   r   r  r   r   r   ptx_cbrt@  s   
r,  c              	   C   2   t |jttdtdfd}|||S )Nr   	__nv_brevr
   r   rC   r   r   r   rD   r#   r   r$   r%   r  r   r   r   ptx_brev_u4L     r1  c              	   C   r-  )Nr   __nv_brevllr/  r0  r   r   r   ptx_brev_u8X  r2  r4  c                 C   s   | |d | tjdS r9   )ctlzrA   r	   booleanr"   r   r   r   ptx_clzd  s   r7  c              	   C   r-  )Nr   __nv_ffsr/  r0  r   r   r   
ptx_ffs_32k     r9  c              	   C   s2   t |jttdtdfd}|||S )Nr   r   
__nv_ffsllr/  r0  r   r   r   
ptx_ffs_64u  r:  r<  c                 C   s   |\}}}| |||S r   )r"  )r#   r   r$   r%   testabr   r   r   ptx_selp  s   
r@  c              	   C   4   t |jtt t t fd}|||S )N
__nv_fmaxfr
   r   rC   r   r   r   rD   r0  r   r   r   
ptx_max_f4     rD  c              
   C   h   t |jtt t t fd}||| ||d |jd t	j
| ||d |jd t	j
gS )N	__nv_fmaxr   r   r
   r   rC   r   r   r   rD   castr%   r	   doubler0  r   r   r   
ptx_max_f8     rK  c              	   C   rA  )N
__nv_fminfrC  r0  r   r   r   
ptx_min_f4  rE  rN  c              
   C   rF  )N	__nv_fminr   r   rH  r0  r   r   r   
ptx_min_f8  rL  rP  c              	   C   sJ   t |jttdt fd}||| ||d |j	d t
jgS )Nr   __nv_llrintr   )r
   r   rC   r   r   r   r   rD   rI  r%   r	   rJ  r0  r   r   r   	ptx_round  s   rR  c                 C   s   dd }|  ||||S )Nc                 S   s   t | s
t | r| S |dkr1|dkrd|d  }d}nd| }d}| | | }t |r0| S n	d|  }| | }t|}t || dkrOdt|d  }|dkr[|| | }|S ||9 }|S )Nr      g      $@gMDg      ?g      ?g       @)mathisinfisnanroundfabs)r   ndigitspow1pow2r   r   r   r   r   round_ndigits  s,   

z$round_to_impl.<locals>.round_ndigits)compile_internal)r#   r   r$   r%   r\  r   r   r   round_to_impl  s   !r^  c                    r  )Nc                    s$   |j \}| | }|||d S r9   )r%   rA   fmul)r#   r   r$   r%   argtyfactorconstr   r   impl  s   zgen_deg_rad.<locals>.implr   )rc  rd  r   rb  r   gen_deg_rad  s   re  g     f@c                    sV   |t jv rt j|dd}|g}n
tj |t|d} fddt||D }||fS )z4
    Convert integer indices into tuple of intp
    r   )rg   rQ   )rQ   c                    s"   g | ]\}}  ||tjqS r   )rI  r	   intp)rt   tir   r#   r   r   rv     s    z&_normalize_indices.<locals>.<listcomp>)r	   integer_domainrP   r
   unpack_tuplelenzip)r#   r   indtyindsindicesr   ri  r   _normalize_indices  s   
rq  c                    r  )Nc                    s   |j \}}}|\}}}	|j}
t| |||\}}|
|kr#td|
|f |jt|kr5td|jt|f | || ||}tj| ||||dd} | ||
||	S )Nzexpect %s but got %sz#indexing %d-D array with %d-D indexT)
wraparound)	r%   rg   rq  	TypeErrorndimrl  
make_arrayr
   get_item_pointer)r#   r   r$   r%   arytyrn  valtyaryro  r   rg   rp  laryptrdispatch_fnr   r   imp  s   
z_atomic_dispatcher.<locals>.impr   )r}  r~  r   r|  r   _atomic_dispatcher  s   r  c                 C   \   |t jkr|j}|t|||fS |t jkr&|j}|t|||fS |d||dS )Nr  	monotonic)	r	   r   rC   rD   r   declare_atomic_add_float32r   declare_atomic_add_float64
atomic_rmwr#   r   rg   r{  r   rG   r   r   r   ptx_atomic_add_tuple,     

r  c                 C   r  )Nr	  r  )	r	   r   rC   rD   r   declare_atomic_sub_float32r   declare_atomic_sub_float64r  r  r   r   r   ptx_atomic_sub=  r  r  c                 C   L   |t jjv r|j}|j}ttd| }|||||fS td| d)Ndeclare_atomic_inc_intzUnimplemented atomic inc with  array	r   cudadeclunsigned_int_numba_typesr   rC   getattrr   rD   rs  r#   r   rg   r{  r   bwrG   r  r   r   r   ptx_atomic_incN     r  c                 C   r  )Ndeclare_atomic_dec_intzUnimplemented atomic dec with r  r  r  r   r   r   ptx_atomic_dec\  r  r  c                    s@   t  fdd}tjtjtjfD ]}t| tj|tj| qd S )Nc                    s2   |t jjv r| ||dS td  d| d)Nr  zUnimplemented atomic z with r  r   r  integer_numba_typesr  rs  r#   r   rg   r{  r   r   r   r   impl_ptx_atomick  s   z+ptx_atomic_bitwise.<locals>.impl_ptx_atomic)r  r	   rf  rP   Tupler  ArrayAny)stubr   r  r   r   r   r   ptx_atomic_bitwisej  s
   r  andorxorc                 C   s,   |t jjv r|d||dS td| d)Nxchgr  zUnimplemented atomic exch with r  r  r  r   r   r   ptx_atomic_exch{  s   r  c                 C      |j }|tjkr|t|||fS |tjkr#|t|||fS |tjtj	fv r4|j
d||ddS |tjtjfv rE|j
d||ddS td| Nr&  r  orderingumaxz&Unimplemented atomic max with %s array)rC   r	   r   rD   r   declare_atomic_max_float64r   declare_atomic_max_float32rB   int64r  uint32uint64rs  r  r   r   r   ptx_atomic_max     

r  c                 C   r  Nr'  r  r  uminz&Unimplemented atomic min with %s array)rC   r	   r   rD   r   declare_atomic_min_float64r   declare_atomic_min_float32rB   r  r  r  r  rs  r  r   r   r   ptx_atomic_min  r  r  c                 C   r  r  )rC   r	   r   rD   r   declare_atomic_nanmax_float64r   declare_atomic_nanmax_float32rB   r  r  r  r  rs  r  r   r   r   ptx_atomic_nanmax  r  r  c                 C   r  r  )rC   r	   r   rD   r   declare_atomic_nanmin_float64r   declare_atomic_nanmin_float32rB   r  r  r  r  rs  r  r   r   r   ptx_atomic_nanmin  r  r  c                 C   s   |j \}}}|\}}}	|j}
| || ||}| tjd}t| ||||f}|jtj	j
v r@|j}|jj}t||||||	S td|
 )Nr   z3Unimplemented atomic compare_and_swap with %s array)r%   rg   ru  rA   r	   rf  r
   rv  r   r  r  rC   r   r   atomic_cmpxchgrs  )r#   r   r$   r%   rw  oldtyrx  ry  oldr   rg   rz  r  r{  rG   r   r   r   r   ptx_atomic_cas_tuple  s   
r  c                 C   s@   t jt t  t dgdddd}|d }|||g d S )Nr   znanosleep.u32 $0;r   Tr   r   )r   r   r   r   r   rD   )r#   r   r$   r%   	nanosleepnsr   r   r   ptx_nanosleep  s
   r  Fc           "   	      s  t tj|d}|dko|ot|dk}|dkr|std j| }	t|tjtj	fp5t|	t
jp5|tjk}
|tjvrC|
sCtd|  |}t||}|tjkr\tj|||d}nD|j}t||||} |}d|d  > |_|r{d|_nt|tj|_t |t!d|}|"t#t!d|}|$||g}t%&tj'} |}|(|}|}g }t)t*|D ]\}}|+| ||9 }qdd	 t*|D } fd
d	|D }|r	tj,t-t!dg dddd}|.|$|g t!d} /tj0|}|1||g}n	 fdd	|D }t|}tj2||dd}  3|  |}! j4|!|"||!j5j6|| /tj0|d d |!7 S )Nr   r   zarray length <= 0zunsupported type: %sra   externalr   c                 S   s   g | ]}|qS r   r   rs   r   r   r   rv   8  s    z"_generic_array.<locals>.<listcomp>c                       g | ]	}  tj|qS r   rA   r	   rf  rs   r#   r   r   rv   9      r   zmov.u32 $0, %dynamic_smem_size;r   Tr   r   c                    r  r   r  rs   r  r   r   rv   K  r  C)rg   rt  layout)datarf   stridesitemsizememinfo)8r   operatorrW   rl  rS   data_model_managerrO   r	   RecordBooleanr   StructModelr  number_domainrs  get_data_typer   	ArrayTyper   r{   r
   alloca_oncerC   add_global_variableget_abi_sizeof
bit_lengthalignlinkageConstant	Undefinedinitializerr   insert_addrspace_convr   r   PointerTyperD   llcreate_target_datadata_layoutget_abi_size	enumeratereversedappendr   r   r   rA   rf  udivr  ru  populate_arrayr  type	_getvalue)"r#   r   rf   rg   rh   ri   rj   	elemcountdynamic_smem
data_modelother_supported_typelldtypelarytydataptrrG   gvmemr  convaddrspaceptr
targetdatar  
laststriderstridesrh  lastsizer  kstridesget_dynshared_sizedynsmem_size	kitemsizekshapert  rw  ry  r   r  r   rl     sx   









rl   c                 C   s   |   S r   )r   )r#   r   r   pyvalr   r   r   cuda_dispatcher_const[  s   r  )F)	functoolsr   r  rT  llvmliter   llvmlite.bindingbindingr  numba.core.imputilsr   r   numba.core.typing.npydeclr   r   numba.core.datamodelr   
numba.corer	   r
   cudadrvr   numbar   
numba.cudar   r   r   numba.cuda.typesr   r   r   registryr  lower_getattr
lower_attrlower_constantr   Moduler&   r*   r-   r0   r4   r7   r<   r>   r@   cg	this_gridrH   rK   gridrB   rV   rX   gridsizer]   rc  
array_liker  r^   r_   rc   sharedarrayIntegerLiteralr  rp   r  rP   ry   localr|   r}   syncthreadsr   syncthreads_counti4r   syncthreads_andr   syncthreads_orr   threadfence_blockr   threadfence_systemr   threadfencer   syncwarpr   r   shfl_sync_intrinsicr   f4f8r   vote_sync_intrinsicr6  r   match_any_syncr   match_all_syncr   r   r   lanemask_ltr   popcr   r   r   r   r  Floatr   r   r   Integerr   r   r  fp16haddhsubhmulhnegr  habsr  hfmar  r  r  heqhnehgehgthlehltr%  hmaxhminr   r   r*  cbrtr,  brevu4r1  u8r4  clzr7  ffsr9  r<  selpr@  r&  rD  rK  r'  rN  rP  rW  rR  r^  re  pi_deg2rad_rad2degradiansdegreesrq  r  atomicr  rf  r  r	  r  incr  decr  r  and_r   r  exchr  r  r  nanmaxr  nanminr  compare_and_swapr  r  r  r  rl   r  r   r   r   r   <module>   s   













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

		
		












	.










	













%






e