
    ܙd.                     h   d dl Zd dlZd dlmZmZmZ d dlmZmZm	Z	 d dlm
Z
mZ d dlmZmZmZmZ d dlmZ d dlmZmZ d dlmZmZmZmZmZmZmZ  ej        d	ej        
          Zd Z d Z!d Z"d Z#d Z$dZ%d Z& ed           G d de
                      Z'e(dk    r ej)                     dS dS )    N)skip_if_mvc_enabledskip_unless_cc_53unittest)skip_on_cudasimskip_unless_cuda_pythonskip_if_cuda_includes_missing)CUDATestCasetest_data_dir)CudaAPIErrorLinkerLinkerError
NvrtcError)require_context)TestCaseignore_internal_warnings)cudavoidfloat64int64int32typeoffloat32
   dtypec                     t           j                            t                    }t          j        d          }||         dz   | |<   d S )N         ?)r   const
array_likeCONST1Dgrid)ACis      Dlib/python3.11/site-packages/numba/cuda/tests/cudadrv/test_linker.pysimple_const_memr'      s9    
g&&A	!AQ4#:AaDDD    c                    d}d}d}	d}
d}d}d}d}d}d}d}d}d}d}d}d}d}d}d}d}t          |          D ]f}||z  }||z  }|	|z  }	|
|z  }
||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }g||z   |	z   |
z   |z   | t          j        d          <   | t          j        d          xx         ||z   |z   |z   |z   z  cc<   | t          j        d          xx         ||z   |z   |z   |z   z  cc<   | t          j        d          xx         ||z   |z   |z   |z   z  cc<   d S )Nr   r   r   )ranger   r"   )xabcdefa1a2a3a4a5b1b2b3b4b5c1c2c3c4c5d1d2d3d4d5r%   s                               r&   func_with_lots_of_registersrF      s   	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B1XX  
a
a
a
a
a
a
a
a
a
a
a
a
a
a
a
q
q
q
q
q2glR'",AdillOdillOOOrBw|b(2--OOOdillOOOrBw|b(2--OOOdillOOOrBw|b(2--OOOOOr(   c                     t           j                            d|          }t          j        d          }|dk    rt	          d          D ]}|||<   t          j                     ||         | |<   d S )Nd   r   r   )r   sharedarrayr"   r*   syncthreads)arydtysmr%   js        r&   simple_smemrP   H   so    			3	$	$B	!AAvvs 	 	ABqEEUCFFFr(   c                     t          j        d          \  }}t           j                            dt                    }|dz   |dz   z  |||f<   t          j                     |||f         | ||f<   d S )N   )r      r   )r   r"   rI   rJ   r   rK   )rL   r%   rO   rN   s       r&   coop_smem2drT   R   sl    9Q<<DAq			8W	-	-BA!a% Bq!tH1a4C1IIIr(   c                 8    t          j        d          }|| |<   d S Nr   )r   r"   )rL   r%   s     r&   simple_maxthreadsrW   Z   s    	!ACFFFr(   i  c                     t           j                            t          |          }t	          |j        d                   D ]}| |         ||<   t	          |j        d                   D ]}||         ||<   d S Nr   )r   localrJ   	LMEM_SIZEr*   shape)r#   BrM   r$   r%   s        r&   simple_lmemr^   b   s{    
C((A171:  t!171:  t! r(   z$Linking unsupported in the simulatorc                      e Zd ZddiZed             Zd Zd Zd Z e	d          d             Z
 e	d          d	             Z e	d          d
             Z ed           ej        e          d                         Ze ej        e          d                         Zd Zd Ze e	d          d                         Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z d Z!d Z"d Z#dS ) 
TestLinkerNUMBA_CUDA_USE_NVIDIA_BINDING0c                 2    t          j        d          }~dS )z9Simply go through the constructor and destructor
        )      )ccN)r   new)selflinkers     r&   test_linker_basiczTestLinker.test_linker_basicn   s     v&&&FFr(   c                    t          j        dd          at          t          dz            }|rdg}ng }t          j        |d|gid             }t          j        dgt          j                  }t          j        d	gt          j                  } |d
         ||           | 	                    |d         dk               d S )Nbarint32(int32)zjitlink.ptxzvoid(int32[:], int32[:])linkc                 t    t          j        d          }| |xx         t          ||                   z  cc<   d S rV   )r   r"   rl   )r+   yr%   s      r&   fooz%TestLinker._test_linking.<locals>.foo   s3    	!AaDDDC!IIDDDDDr(   {   r   iA  )r   r   r   i  )
r   declare_devicerl   strr
   jitnprJ   r   
assertTrue)rh   eagerrn   argsrq   r#   r]   s          r&   _test_linkingzTestLinker._test_linkingu   s    !%88==011 	./DDD	4	%tf	%	%	 	 
&	%	 HcU"(+++HcU"(+++D	!Q!-.....r(   c                 2    |                      d           d S )NFrx   rz   rh   s    r&   test_linking_lazy_compilez$TestLinker.test_linking_lazy_compile   s    '''''r(   c                 2    |                      d           d S )NTr|   r}   r~   s    r&   test_linking_eager_compilez%TestLinker.test_linking_eager_compile   s    &&&&&r(   zNVIDIA Binding needed for NVRTCc                 x   t          j        dd          t          t          dz            }t          j        |g          fd            }t          j        dt
          j                  }t          j        |          } |d         ||           |d	z  }t
          j	        
                    ||           d S )
Nrl   rm   
jitlink.curn   c                     t          j        d          }|t          |           k     r ||                   | |<   d S d S rV   )r   r"   len)rr+   r%   rl   s      r&   kernelz*TestLinker.test_linking_cu.<locals>.kernel   s?    	!A3q66zzs1Q4yy! zr(   r   r   )r       rR   )r   rs   rt   r
   ru   rv   aranger   
zeros_liketestingassert_array_equal)rh   rn   r   r+   r   expectedrl   s         @r&   test_linking_cuzTestLinker.test_linking_cu   s    !%88=</00	v				! 	! 	! 	! 
		! Ib)))M!ua q5

%%a22222r(   c                    t          j        dd          t          t          dz            }t	          j        d          5 }t                       t          j        d|g          fd            }d d d            n# 1 swxY w Y   |                     t          |          d	d
           | 
                    dt          |d         j                             | 
                    dt          |d         j                             d S )Nrl   rm   zwarn.cuT)recordvoid(int32)r   c                      |            d S N r+   rl   s    r&   r   z6TestLinker.test_linking_cu_log_warning.<locals>.kernel       Ar(   r   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   rs   rt   r
   warningscatch_warningsr   ru   assertEqualr   assertInmessage)rh   rn   wr   rl   s       @r&   test_linking_cu_log_warningz&TestLinker.test_linking_cu_log_warning   s;   !%88=9,--$D111 	Q$&&&Xm4&111    21		 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	Q$BCCC*C!,=,=>>>5s1Q4<7H7HIIIIIs   1B  BBc                    t          j        dd          t          t          dz            }|                     t
                    5 }t          j        d|g          fd            }d d d            n# 1 swxY w Y   |j        j        d         }| 	                    d|           | 	                    d	|           | 	                    d
|           d S )Nrl   rm   zerror.cur   r   c                      |            d S r   r   r   s    r&   r   z0TestLinker.test_linking_cu_error.<locals>.kernel   r   r(   r   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")
r   rs   rt   r
   assertRaisesr   ru   	exceptionry   r   )rh   rn   r0   r   msgrl   s        @r&   test_linking_cu_errorz TestLinker.test_linking_cu_error   s   !%88=:-..z** 	aXm4&111    21	 	 	 	 	 	 	 	 	 	 	 	 	 	 	
 kq!137778#>>>8#>>>>>s   #A77A;>A;z0NVRTC not available when ctypes binding is used.)envvarsc                     t          t          dz            }d}|                     t          |          5  t	          j        d|g          d             }d d d            d S # 1 swxY w Y   d S )Nr   zBLinking CUDA source files is not supported with the ctypes bindingvoid()r   c                      d S r   r   r   r(   r&   r1   z8TestLinker.test_linking_cu_ctypes_unsupported.<locals>.f       r(   )rt   r
   assertRaisesRegexNotImplementedErrorr   ru   )rh   rn   r   r1   s       r&   "test_linking_cu_ctypes_unsupportedz-TestLinker.test_linking_cu_ctypes_unsupported   s     =</00 ##$7== 	 	XhdV,,,  -,	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	s   !A##A'*A'c                     d}|                      t          |          5  t          j        d          d             }d d d            d S # 1 swxY w Y   d S )NzLUse of float16 requires the use of the NVIDIA CUDA bindings and setting the zvoid(float16[::1])c                 P    t           j                            | d                    d S rY   )r   fp16hexp10r+   s    r&   hexp10_vectorszCTestLinker.test_linking_float16_unsupported.<locals>.hexp10_vectors   s"    	  1&&&&&r(   )r   r   r   ru   )rh   r   r   s      r&    test_linking_float16_unsupportedz+TestLinker.test_linking_float16_unsupported   s    +##$7== 	' 	'X*++' ' ,+'	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	' 	's   A		AAc                     d}|                      t          |          5  t          j        ddg          d             }d d d            d S # 1 swxY w Y   d S )Nz/Don't know how to link file with extension .cuhr   z
header.cuhr   c                      d S r   r   r   r(   r&   r   z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernel   r   r(   r   RuntimeErrorr   ru   rh   expected_errr   s      r&   #test_linking_unknown_filetype_errorz.TestLinker.test_linking_unknown_filetype_error   s    H##L,?? 	 	Xhl^444  54	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	   !AAAc                     d}|                      t          |          5  t          j        ddg          d             }d d d            d S # 1 swxY w Y   d S )Nz-Don't know how to link file with no extensionr   datar   c                      d S r   r   r   r(   r&   r   zDTestLinker.test_linking_file_with_no_extension_error.<locals>.kernel   r   r(   r   r   s      r&   )test_linking_file_with_no_extension_errorz4TestLinker.test_linking_file_with_no_extension_error   s    F##L,?? 	 	XhfX...  /.	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	r   c                 t    t          t          dz            }t          j        d|g          d             }d S )Nzcuda_include.cur   r   c                      d S r   r   r   r(   r&   r   z7TestLinker.test_linking_cu_cuda_include.<locals>.kernel  s    Dr(   )rt   r
   r   ru   )rh   rn   r   s      r&   test_linking_cu_cuda_includez'TestLinker.test_linking_cu_cuda_include   sN     =#4455 
($	(	(	(	 	 
)	(	 	 	r(   c                     |                      t                    5 }t          j        ddg          d             }d d d            n# 1 swxY w Y   |                     d|j        j                   d S )Nvoid(int32[::1])znonexistent.ar   c                     d| d<   d S rY   r   r   s    r&   r1   z2TestLinker.test_try_to_link_nonexistent.<locals>.f	  s    !r(   znonexistent.a not found)r   r   r   ru   r   r   ry   )rh   r0   r1   s      r&   test_try_to_link_nonexistentz'TestLinker.test_try_to_link_nonexistent  s    {++ 	qX(/@AAA  BA	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	/1ABBBBBs   !AAAc                     t          j        t                    } |j        t	          j        d          gt          d          R  }|                     |                                d           dS )a  Ensure that the jitted kernel used in the test_set_registers_* tests
        uses more than 57 registers - this ensures that test_set_registers_*
        are really checking that they reduced the number of registers used from
        something greater than the maximum.r      9   N)	r   ru   rF   
specializerv   emptyr*   assertGreaterget_regs_per_threadrh   compileds     r&   test_set_registers_no_maxz$TestLinker.test_set_registers_no_max  sd    
 8788&8&rx||?eAhh???877992>>>>>r(   c                      t          j        d          t                    } |j        t	          j        d          gt          d          R  }|                     |                                d           d S )Nr   max_registersr   r   	r   ru   rF   r   rv   r   r*   assertLessEqualr   r   s     r&   test_set_registers_57z TestLinker.test_set_registers_57  o    -48"---.IJJ&8&rx||?eAhh???X99;;R@@@@@r(   c                      t          j        d          t                    } |j        t	          j        d          gt          d          R  }|                     |                                d           d S )N&   r   r   r   r   r   s     r&   test_set_registers_38z TestLinker.test_set_registers_38  r   r(   c           	          t          t          d d d         t          t          t          t          t          t                    } t          j        |d          t
                    }|                     |                                d           d S )Nr   r   r   )r   r   r   r   ru   rF   r   r   )rh   sigr   s      r&   test_set_registers_eagerz#TestLinker.test_set_registers_eager!  sl    733Q3<ueUEJJ248Cr2223NOOX99;;R@@@@@r(   c                     t          t          d d d                   } t          j        |          t                    }|                                }|                     |t          j                   d S rV   )	r   r   r   ru   r'   get_const_mem_sizeassertGreaterEqualr!   nbytes)rh   r   r   const_mem_sizes       r&   test_get_const_mem_sizez"TestLinker.test_get_const_mem_size&  s`    733Q3<   48C==!122!4466?????r(   c                     t          j        t                    } |j        t	          j        d          gt          d          R  }|                                }|                     |d           d S )Nr   r   r   )	r   ru   rF   r   rv   r   r*   get_shared_mem_per_blockr   )rh   r   shared_mem_sizes      r&   test_get_no_shared_memoryz$TestLinker.test_get_no_shared_memory,  sg    8788&8&rx||?eAhh???";;==!,,,,,r(   c                    t          t          d d d         t          t          j                            } t	          j        |          t                    }|                                }|                     |d           d S )Nr   i  )	r   r   r   rv   r   ru   rP   r   r   )rh   r   r   r   s       r&   test_get_shared_mem_per_blockz(TestLinker.test_get_shared_mem_per_block2  si    51:vbh//00 48C==--";;==#.....r(   c                    t          j        t                    }|                    t	          j        dt          j                  t          j                  }|                                }| 	                    |d           d S )NrH   r   i   )
r   ru   rP   r   rv   zerosr   r   r   r   )rh   r   compiled_specializedr   s       r&   #test_get_shared_mem_per_specializedz.TestLinker.test_get_shared_mem_per_specialized8  sl    8K(('22HS)))2: 7  7.GGII#.....r(   c                      t          j        d          t                    }|                                }|                     |d           d S )Nzvoid(float32[:,::1])r   )r   ru   rT   get_max_threads_per_blockr   )rh   r   max_threadss      r&   test_get_max_threads_per_blockz)TestLinker.test_get_max_threads_per_block?  sI    348233K@@88::;*****r(   c                 J    t          j        d          t                    }|                                }|dz   }t	          j        |t          j                  }	  |d|f         |           d S # t          $ r&}|                     d|j	                   Y d }~d S d }~ww xY w)Nr   r   r   cuLaunchKernel)
r   ru   rW   r   rv   r   r   r   r   r   )rh   r   r   nelemrL   r0   s         r&   test_max_threads_exceededz$TestLinker.test_max_threads_exceededD  s    /48.//0ABB88::ahuBH---	3HQXs##### 	3 	3 	3MM*AE222222222	3s   A2 2
B"<BB"c                 |   t          t          d d d         t          d d d         t          t          j                            } t	          j        |          t                    }|                                }t          j        t          j                  j	        t          z  }|                     ||           d S rV   )r   r   r   rv   r   ru   r^   get_local_mem_per_threadr   itemsizer[   r   )rh   r   r   local_mem_size	calc_sizes        r&   test_get_local_mem_per_threadz(TestLinker.test_get_local_mem_per_threadN  s    51:uSSqSz6"(+;+;<< 48C==--!::<<HRX&&/);		:::::r(   c                    t          j        t                    }|                    t	          j        t          t          j                  t	          j        t          t          j                  t          j                  }|	                                }t	          j
        t          j                  j        t          z  }|                     ||           d S )Nr   )r   ru   r^   r   rv   r   r[   r   r   r   r   r   r   )rh   r   r   r   r   s        r&   "test_get_local_mem_per_specializedz-TestLinker.test_get_local_mem_per_specializedU  s    8K(('22HYbh///HYbh///J    .FFHHHRZ((1I=		:::::r(   N)$__name__
__module____qualname___NUMBA_NVIDIA_BINDING_0_ENVr   rj   rz   r   r   r   r   r   r   r   r   run_test_in_subprocessr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r(   r&   r`   r`   j   s       #BC"H  _/ / /.( ( (' ' ' >??3 3 @?3* >??J J @?J$ >??? ? @??0 KLL$X$-HIII  JI ML $X$-HIII' ' JI '     #>??  @? #"C C C? ? ?A A A
A A A
A A A
@ @ @- - -/ / // / /+ + +
3 3 3; ; ;; ; ; ; ;r(   r`   __main__)*numpyrv   r   numba.cuda.testingr   r   r   r   r   r   r	   r
   numba.cuda.cudadrv.driverr   r   r   r   
numba.cudar   numba.tests.supportr   r   numbar   r   r   r   r   r   r   r   r!   r'   rF   rP   rT   rW   r[   r^   r`   r   mainr   r(   r&   <module>r     s        O O O O O O O O O O? ? ? ? ? ? ? ? ? ? : : : : : : : :@ @ @ @ @ @ @ @ @ @ @ @ & & & & & & B B B B B B B B D D D D D D D D D D D D D D D D D D ")Bbj
)
)
)  -. -. -.`      
 	   788r; r; r; r; r; r; r; 98r;j zHMOOOOO r(   