
    J/Ph'                     `   d dl Zd dl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
mZ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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)unittest)skip_on_cudasimskip_if_cuda_includes_missing)CUDATestCasetest_data_dir)CudaAPIErrorLinkerLinkerError)
NvrtcError)require_context)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      d/var/www/html/test/jupyter/venv/lib/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_registersrB      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_smemrL   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   rE   rF   r   rG   )rH   r!   rK   rJ   s       r"   coop_smem2drP   R   sl    9Q<<DAq			8W	-	-BA!a% Bq!tH1a4C1IIIr$   c                 8    t          j        d          }|| |<   d S Nr   )r   r   )rH   r!   s     r"   simple_maxthreadsrS   Z   s    	!ACFFFr$   i  c                     t           j                            t          |          }t	          |j        d                   D ]}| |         ||<   t	          |j        d                   D ]}||         ||<   d S Nr   )r   localrF   	LMEM_SIZEr&   shape)r   BrI   r    r!   s        r"   simple_lmemrZ   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d Z	d Z
d	 Zd
 Zd Ze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 rR   )r   r   rh   )r'   yr!   s      r"   fooz%TestLinker._test_linking.<locals>.foo   s3    	!AaDDDC!IIDDDDDr$   {   r   iA  )r   r   r   i  )
r   declare_devicerh   strr   jitnprF   r   
assertTrue)rd   eagerrj   argsrm   r   rY   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rt   rv   rd   s    r"   test_linking_lazy_compilez$TestLinker.test_linking_lazy_compile   s    '''''r$   c                 2    |                      d           d S )NTrx   ry   rz   s    r"   test_linking_eager_compilez%TestLinker.test_linking_eager_compile   s    &&&&&r$   c                 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 )
Nrh   ri   z
jitlink.curj   c                     t          j        d          }|t          |           k     r ||                   | |<   d S d S rR   )r   r   len)rr'   r!   rh   s      r"   kernelz*TestLinker.test_linking_cu.<locals>.kernel   s?    	!A3q66zzs1Q4yy! zr$   r   r   )r       rN   )r   ro   rp   r   rq   rr   aranger   
zeros_liketestingassert_array_equal)rd   rj   r   r'   r   expectedrh   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 )Nrh   ri   zwarn.cuT)recordvoid(int32)r   c                      |            d S N r'   rh   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   ro   rp   r   warningscatch_warningsr   rq   assertEqualr   assertInmessage)rd   rj   wr   rh   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 )Nrh   ri   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   ro   rp   r   assertRaisesr   rq   	exceptionru   r   )rd   rj   r,   r   msgrh   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;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 .cuhvoid()z
header.cuhr   c                      d S r   r   r   r$   r"   r   z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernel       r$   assertRaisesRegexRuntimeErrorr   rq   rd   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$   )rp   r   r   rq   )rd   rj   r   s      r"   test_linking_cu_cuda_includez'TestLinker.test_linking_cu_cuda_include   sL    =#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 rU   r   )r'   s    r"   r-   z2TestLinker.test_try_to_link_nonexistent.<locals>.f   s    !r$   znonexistent.a not found)r   r
   r   rq   r   r   ru   )rd   r,   r-   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   rq   rB   
specializerr   emptyr&   assertGreaterget_regs_per_threadrd   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   rq   rB   r   rr   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   rq   rB   r   r   )rd   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 rR   )	r   r   r   rq   r#   get_const_mem_sizeassertGreaterEqualr   nbytes)rd   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   rq   rB   r   rr   r   r&   get_shared_mem_per_blockr   )rd   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   rr   r   rq   rL   r   r   )rd   r   r   r   s       r"   test_get_shared_mem_per_blockz(TestLinker.test_get_shared_mem_per_block  si    51:vbh//00 48C==--";;==#.....r$   c                    t          j        t                    }|                    t	          j        dt          j                  t          j                  }|                                }| 	                    |d           d S )NrD   r   i   )
r   rq   rL   r   rr   zerosr   r   r   r   )rd   r   compiled_specializedr   s       r"   #test_get_shared_mem_per_specializedz.TestLinker.test_get_shared_mem_per_specialized  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   rq   rP   get_max_threads_per_blockr   )rd   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   rq   rS   r   rr   r   r   r   r   r   )rd   r   r   nelemrH   r,   s         r"   test_max_threads_exceededz$TestLinker.test_max_threads_exceeded   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 rR   )r   r   r   rr   r   rq   rZ   get_local_mem_per_threadr   itemsizerW   r   )rd   r   r   local_mem_size	calc_sizes        r"   test_get_local_mem_per_threadz(TestLinker.test_get_local_mem_per_thread*  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   rq   rZ   r   rr   r   rW   r   r   r   r   r   r   )rd   r   r   r   r   s        r"   "test_get_local_mem_per_specializedz-TestLinker.test_get_local_mem_per_specialized1  s    8K(('22HYbh///HYbh///J    .FFHHHRZ((1I=		:::::r$   N)__name__
__module____qualname___NUMBA_NVIDIA_BINDING_0_ENVr   rf   rv   r{   r}   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r$   r"   r\   r\   j   s       #BC"H  _/ / /.( ( (' ' '3 3 3*J J J$? ? ?$     #  #"C C C? ? ?A A A
A A A
A A A
@ @ @- - -/ / // / /+ + +
3 3 3; ; ;; ; ; ; ;r$   r\   __main__)'numpyrr   r   numba.cuda.testingr   r   r   r   r   numba.cuda.cudadrv.driverr   r	   r
   numba.cuda.cudadrv.errorr   
numba.cudar   numba.tests.supportr   numbar   r   r   r   r   r   r   r   r   r#   rB   rL   rP   rS   rW   rZ   r\   r   mainr   r$   r"   <module>r      s        ' ' ' ' ' ' O O O O O O O O : : : : : : : :4 4 4 4 4 4 4 4 4 4 / / / / / / & & & & & & 8 8 8 8 8 8 D D D D D D D D D D D D D D D D D D ")Bbj
)
)
)  -. -. -.`      
 	   788N; N; N; N; N; N; N; 98N;b zHMOOOOO r$   