+
    :i'                     b   ^ RI t^ RIt^ RIHt ^ RIHtHt ^ RIHtHt ^ RI	H
t
HtHt ^ RIHt ^ RIHt ^ RIHt ^ R	IHtHtHtHtHtHtHt ]P6                  ! ^
]P,                  R
7      tR tR tR tR t R t!Rt"R t#]! R4       ! R R]4      4       t$]%R8X  d   ]PL                  ! 4        R# R# )    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                     \         P                  P                  \        4      p\         P                  ! ^4      pW,          R,           W&   R# )         ?N)r   const
array_likeCONST1Dgrid)ACis   &  b/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/tests/cudadrv/test_linker.pysimple_const_memr"      s0    

g&A		!A4#:AD    c                    R pR pR p	R p
R pR pR pR pR pR pR pR pR pR pR p^
p^
p^
p^
p^
p\        V4       F  pWr,          pW,          pW,          p	W,          p
W,          pW,          pW,          pW,          pW,          pVV,          pVV,          pVV,          pVV,          pVV,          pVV,          pVV,          pVV,          pVV,          pVV,          pVV,          pK  	  Wx,           V	,           V
,           V,           V \        P                  ! ^4      &   V \        P                  ! ^4      ;;,          W,           V,           V,           V,           ,          uu&   V \        P                  ! ^4      ;;,          VV,           V,           V,           V,           ,          uu&   V \        P                  ! ^4      ;;,          VV,           V,           V,           V,           ,          uu&   R# )r   N)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_registersrA      s   	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B	B1X









a
a
a
a
a
a
q
q
q
q
q) * glR'",AdiilOdiilOrw|b(2--OdiilOrBw|b(2--OdiilOrBw|b(2--Or#   c                     \         P                  P                  ^dV4      p\         P                  ! ^4      pV^ 8X  d   \	        ^d4       F  pWBV&   K	  	  \         P
                  ! 4        W#,          W&   R# )d   N)r   sharedarrayr   r%   syncthreads)arydtysmr    js   &&   r!   simple_smemrK   H   sV    			3	$B		!AAvsAqE UCFr#   c                     \         P                  ! ^4      w  r\         P                  P                  R\        4      pV^,           V^,           ,          W1V3&   \         P
                  ! 4        W1V3,          WV3&   R# )   N)
      )r   r   rD   rE   r   rF   )rG   r    rJ   rI   s   &   r!   coop_smem2drP   R   s\    99Q<DA			8W	-BA!a% B!tHa4C1Ir#   c                 :    \         P                  ! ^4      pWV&   R# r   N)r   r   )rG   r    s   & r!   simple_maxthreadsrS   Z   s    		!AFr#   i  c                    \         P                  P                  \        V4      p\	        VP
                  ^ ,          4       F  pW,          W4&   K  	  \	        VP
                  ^ ,          4       F  pW4,          W&   K  	  R# r   N)r   localrE   	LMEM_SIZEr%   shape)r   BrH   r   r    s   &&&  r!   simple_lmemrZ   b   sX    

C(A1771:t 1771:t r#   z$Linking unsupported in the simulatorc                      a  ] tR t^jt o RR/t]R 4       tR tR tR t	R t
R tR	 tR
 tR t]R 4       tR tR tR tR tR tR tR tR tR tR tR tR tR tRtV tR# )
TestLinkerNUMBA_CUDA_USE_NVIDIA_BINDING0c                6    \         P                  ! RR7      p?R# )z9Simply go through the constructor and destructor
        )ccN)      )r	   new)selflinkers   & r!   test_linker_basicTestLinker.test_linker_basicn   s     v&r#   c                   \         P                  ! R R4      s\        \        R,          4      pV'       d   R.pM. p\         P
                  ! VRV./ R 4       p\        P                  ! ^{.\        P                  R7      p\        P                  ! R.\        P                  R7      pVR	,          ! WV4       V P                  V^ ,          R
8H  4       R# )barint32(int32)zjitlink.ptxzvoid(int32[:], int32[:])linkc                 x    \         P                  ! ^4      pW;;,          \        W,          4      ,          uu&   R# rR   )r   r   ri   )r&   yr    s   && r!   foo%TestLinker._test_linking.<locals>.foo   s!    		!ADCIDr#   r   iA  N)r   r   i  )
r   declare_deviceri   strr   jitnprE   r   
assertTrue)rd   eagerrk   argsrn   r   rY   s   &&     r!   _test_linkingTestLinker._test_linkingu   s    !!%8==01./DD	4	%tf	%	 
&	 HHcU"((+HHcU"((+D	!!-.r#   c                *    V P                  R R7       R# )Fru   Nrw   rd   s   &r!   test_linking_lazy_compile$TestLinker.test_linking_lazy_compile   s    'r#   c                *    V P                  R R7       R# )Trz   Nr{   r|   s   &r!   test_linking_eager_compile%TestLinker.test_linking_eager_compile   s    &r#   c                  a \         P                  ! R R4      o\        \        R,          4      p\         P                  ! V.R7      V3R l4       p\
        P                  ! ^
\
        P                  R7      p\
        P                  ! V4      pVR,          ! WC4       V^,          p\
        P                  P                  WE4       R# )ri   rj   z
jitlink.curk   c                 x   < \         P                  ! ^4      pV\        V 4      8  d   S! W,          4      W&   R# R# rR   )r   r   len)rr&   r    ri   s   && r!   kernel*TestLinker.test_linking_cu.<locals>.kernel   s-    		!A3q6z14y r#   r   N)r       )r   rp   rq   r   rr   rs   aranger   
zeros_liketestingassert_array_equal)rd   rk   r   r&   r   expectedri   s   &     @r!   test_linking_cuTestLinker.test_linking_cu   s    !!%8=</0	v		! 
	! IIb)MM!ua q5


%%a2r#   c                  a \         P                  ! R R4      o\        \        R,          4      p\        P
                  ! RR7      ;_uu_ 4       p\        4        \         P                  ! RV.R7      V3R l4       pRRR4       V P                  \        X4      ^R	4       V P                  R
\        V^ ,          P                  4      4       V P                  R\        V^ ,          P                  4      4       R#   + '       g   i     L; i)ri   rj   zwarn.cuT)recordvoid(int32)r   c                    < S! V 4       R # N r&   ri   s   &r!   r   6TestLinker.test_linking_cu_log_warning.<locals>.kernel   
    Ar#   NzExpected warnings from NVRTCzNVRTC log messageszdeclared but never referenced)r   rp   rq   r   warningscatch_warningsr   rr   assertEqualr   assertInmessage)rd   rk   wr   ri   s   &   @r!   test_linking_cu_log_warning&TestLinker.test_linking_cu_log_warning   s    !!%8=9,-$$D11Q$&XXm4&1 2	 2 	Q$BC*C!,=>5s1Q4<<7HI 21s   -C99D		c                  a \         P                  ! R R4      o\        \        R,          4      pV P	                  \
        4      ;_uu_ 4       p\         P                  ! RV.R7      V3R l4       pRRR4       XP                  P                  ^ ,          pV P                  RV4       V P                  RV4       V P                  R	V4       R#   + '       g   i     Le; i)
ri   rj   zerror.cur   r   c                    < S! V 4       R # r   r   r   s   &r!   r   0TestLinker.test_linking_cu_error.<locals>.kernel   r   r#   NzNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")
r   rp   rq   r   assertRaisesr   rr   	exceptionrv   r   )rd   rk   r+   r   msgri   s   &    @r!   test_linking_cu_error TestLinker.test_linking_cu_error   s    !!%8=:-.z**aXXm4&1 2 +
 kkq!1378#>8#> +*s   #CC	c                    R pV P                  \        V4      ;_uu_ 4        \        P                  ! RR.R7      R 4       pRRR4       R#   + '       g   i     R# ; i)z/Don't know how to link file with extension .cuhvoid()z
header.cuhr   c                      R # r   r   r   r#   r!   r   >TestLinker.test_linking_unknown_filetype_error.<locals>.kernel       r#   NassertRaisesRegexRuntimeErrorr   rr   rd   expected_errr   s   &  r!   #test_linking_unknown_filetype_error.TestLinker.test_linking_unknown_filetype_error   sI    H##L,??XXhl^4 5 @???    AA	c                    R pV P                  \        V4      ;_uu_ 4        \        P                  ! RR.R7      R 4       pRRR4       R#   + '       g   i     R# ; i)z-Don't know how to link file with no extensionr   datar   c                      R # r   r   r   r#   r!   r   DTestLinker.test_linking_file_with_no_extension_error.<locals>.kernel   r   r#   Nr   r   s   &  r!   )test_linking_file_with_no_extension_error4TestLinker.test_linking_file_with_no_extension_error   sI    F##L,??XXhfX. / @???r   c                p    \        \        R ,          4      p\        P                  ! RV.R7      R 4       pR# )zcuda_include.cur   r   c                      R # r   r   r   r#   r!   r   7TestLinker.test_linking_cu_cuda_include.<locals>.kernel   s    r#   N)rq   r   r   rr   )rd   rk   r   s   &  r!   test_linking_cu_cuda_include'TestLinker.test_linking_cu_cuda_include   s3    =#445 
($	(	 
)	r#   c                   V P                  \        4      ;_uu_ 4       p\        P                  ! R R.R7      R 4       pRRR4       V P	                  RXP
                  P                  4       R#   + '       g   i     L8; i)void(int32[::1])znonexistent.ar   c                     ^ V ^ &   R# rU   r   )r&   s   &r!   r,   2TestLinker.test_try_to_link_nonexistent.<locals>.f   s    !r#   Nznonexistent.a not found)r   r
   r   rr   r   r   rv   )rd   r+   r,   s   &  r!   test_try_to_link_nonexistent'TestLinker.test_try_to_link_nonexistent   s`    {++qXX(/@A B , 	/1A1AB	 ,+s    A..A>	c                    \         P                  ! \        4      pVP                  ! \        P
                  ! ^ 4      .\        ^4      O5!  pV P                  VP                  4       ^94       R# )zEnsure 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.N)	r   rr   rA   
specializers   emptyr%   assertGreaterget_regs_per_threadrd   compileds   & r!   test_set_registers_no_max$TestLinker.test_set_registers_no_max   sM    
 8878&&rxx|?eAh?87792>r#   c                    \         P                  ! ^9R7      ! \        4      pVP                  ! \        P
                  ! ^ 4      .\        ^4      O5!  pV P                  VP                  4       ^94       R# )9   max_registersN	r   rr   rA   r   rs   r   r%   assertLessEqualr   r   s   & r!   test_set_registers_57 TestLinker.test_set_registers_57   P    88"-.IJ&&rxx|?eAh?X99;R@r#   c                    \         P                  ! ^&R7      ! \        4      pVP                  ! \        P
                  ! ^ 4      .\        ^4      O5!  pV P                  VP                  4       ^&4       R# )&   r   Nr   r   s   & r!   test_set_registers_38 TestLinker.test_set_registers_38   r   r#   c           	         \        \        R ,          \        \        \        \        \        \        4      p\        P                  ! V^&R7      ! \
        4      pV P                  VP                  4       ^&4       R# )NNr   r   N)r   r   r   r   rr   rA   r   r   )rd   sigr   s   &  r!   test_set_registers_eager#TestLinker.test_set_registers_eager   sK    73<ueUEJ88Cr23NOX99;R@r#   c                    \        \        R ,          4      p\        P                  ! V4      ! \        4      pVP                  4       pV P                  V\        P                  4       R# r   N)	r   r   r   rr   r"   get_const_mem_sizeassertGreaterEqualr   nbytes)rd   r   r   const_mem_sizes   &   r!   test_get_const_mem_size"TestLinker.test_get_const_mem_size  sE    73< 88C=!12!446?r#   c                    \         P                  ! \        4      pVP                  ! \        P
                  ! ^ 4      .\        ^4      O5!  pVP                  4       pV P                  V^ 4       R# )r   N)	r   rr   rA   r   rs   r   r%   get_shared_mem_per_blockr   )rd   r   shared_mem_sizes   &  r!   test_get_no_shared_memory$TestLinker.test_get_no_shared_memory  sP    8878&&rxx|?eAh?";;=!,r#   c                    \        \        R ,          \        \        P                  4      4      p\        P
                  ! V4      ! \        4      pVP                  4       pV P                  VR4       R# )r   i  N)	r   r   r   rs   r   rr   rK   r   r   )rd   r   r   r   s   &   r!   test_get_shared_mem_per_block(TestLinker.test_get_shared_mem_per_block  sK    5:vbhh/088C=-";;=#.r#   c                   \         P                  ! \        4      pVP                  \        P
                  ! ^d\        P                  R7      \        P                  4      pVP                  4       pV P                  VR4       R# )rC   r   i   N)
r   rr   rK   r   rs   zerosr   r   r   r   )rd   r   compiled_specializedr   s   &   r!   #test_get_shared_mem_per_specialized.TestLinker.test_get_shared_mem_per_specialized  sW    88K('22HHS)2:: 7.GGI#.r#   c                    \         P                  ! R 4      ! \        4      pVP                  4       pV P	                  V^ 4       R# )zvoid(float32[:,::1])N)r   rr   rP   get_max_threads_per_blockr   )rd   r   max_threadss   &  r!   test_get_max_threads_per_block)TestLinker.test_get_max_threads_per_block  s4    8823K@88:;*r#   c                R   \         P                  ! R 4      ! \        4      pVP                  4       pV^,           p\        P
                  ! V\        P                  R7      p V^V3,          ! V4       R#   \         d(   pT P                  RTP                  4        Rp?R# Rp?ii ; i)r   r   cuLaunchKernelN)
r   rr   rS   r   rs   r   r   r   r   r   )rd   r   r   nelemrG   r+   s   &     r!   test_max_threads_exceeded$TestLinker.test_max_threads_exceeded   s~    88./0AB88:ahhuBHH-	3QXs# 	3MM*AEE22	3s   !A4 4B&?B!!B&c                n   \        \        R ,          \        R ,          \        \        P                  4      4      p\        P
                  ! V4      ! \        4      pVP                  4       p\        P                  ! \        P                  4      P                  \        ,          pV P                  W44       R# r   )r   r   r   rs   r   rr   rZ   get_local_mem_per_threadr   itemsizerW   r   )rd   r   r   local_mem_size	calc_sizes   &    r!   test_get_local_mem_per_thread(TestLinker.test_get_local_mem_per_thread*  sl    5:uSz6"((+;<88C=-!::<HHRXX&//);	:r#   c                   \         P                  ! \        4      pVP                  \        P
                  ! \        \        P                  R 7      \        P
                  ! \        \        P                  R 7      \        P                  4      pVP                  4       p\        P                  ! \        P                  4      P                  \        ,          pV P                  W44       R# )r   N)r   rr   rZ   r   rs   r   rW   r   r   r  r   r  r   )rd   r   r   r  r  s   &    r!   "test_get_local_mem_per_specialized-TestLinker.test_get_local_mem_per_specialized1  s    88K('22HHYbhh/HHYbhh/JJ  .FFHHHRZZ(11I=	:r#   r   N) __name__
__module____qualname____firstlineno___NUMBA_NVIDIA_BINDING_0_ENVr   rf   rw   r}   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r
  __static_attributes____classdictcell__)__classdict__s   @r!   r\   r\   j   s     #BC"H /.('3*J$?$ # #C?A
A
A
@-//+
3;; ;r#   r\   __main__)'numpyrs   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"   rA   rK   rP   rS   rW   rZ   r\   r  mainr   r#   r!   <module>r     s      ' O :4 4 / & 8 D D D ))Bbjj
)-.`
 	 78N; N; 9N;b zMMO r#   