
    xKg'                     Z   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j6                  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jL                          yy)    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                  j                  t              }t        j                  d      }||   dz   | |<   y )N         ?)r   const
array_likeCONST1Dgrid)ACis      h/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudadrv/test_linker.pysimple_const_memr#      s4    

g&A		!AQ4#:AaD    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  }h ||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<   y )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1X
a
a
a
a
a
a
a
a
a
a
a
a
a
a
a
q
q
q
q
q) * 2glR'",AdiilOdiilOrBw|b(2--OdiilOrBw|b(2--OdiilOrBw|b(2--Or$   c                     t         j                  j                  d|      }t        j                  d      }|dk(  rt	        d      D ]  }|||<   	 t        j
                          ||   | |<   y )Nd   r   r   )r   sharedarrayr   r&   syncthreads)arydtysmr!   js        r"   simple_smemrL   H   s\    			3	$B		!AAvsABqE UCFr$   c                     t        j                  d      \  }}t         j                  j                  dt              }|dz   |dz   z  |||f<   t        j
                          |||f   | ||f<   y )N   )r      r   )r   r   rE   rF   r   rG   )rH   r!   rK   rJ   s       r"   coop_smem2drP   R   sd    99Q<DAq			8W	-BA!a% Bq!tH1a4C1Ir$   c                 8    t        j                  d      }|| |<   y Nr   )r   r   )rH   r!   s     r"   simple_maxthreadsrS   Z   s    		!ACFr$   i  c                     t         j                  j                  t        |      }t	        |j
                  d         D ]
  }| |   ||<    t	        |j
                  d         D ]
  }||   ||<    y Nr   )r   localrF   	LMEM_SIZEr&   shape)r   BrI   r    r!   s        r"   simple_lmemrZ   b   s`    

C(A1771:t! 1771: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y)
TestLinkerNUMBA_CUDA_USE_NVIDIA_BINDING0c                 2    t        j                  d      }~y)z9Simply go through the constructor and destructor
        )      )ccN)r	   new)selflinkers     r"   test_linker_basiczTestLinker.test_linker_basicn   s     v&r$   c                 z   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
   ||       | j                  |d   dk(         y )Nbarint32(int32)zjitlink.ptxzvoid(int32[:], int32[:])linkc                 `    t        j                  d      }| |xx   t        ||         z  cc<   y rR   )r   r   rh   )r'   yr!   s      r"   fooz%TestLinker._test_linking.<locals>.foo   s%    		!AaDC!IDr$   {   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==01./DD	4	%tf	%	 
&	 HHcU"((+HHcU"((+D	!Q!-.r$   c                 (    | j                  d       y )NFrt   rv   rd   s    r"   test_linking_lazy_compilez$TestLinker.test_linking_lazy_compile   s    'r$   c                 (    | j                  d       y )NTrx   ry   rz   s    r"   test_linking_eager_compilez%TestLinker.test_linking_eager_compile   s    &r$   c                 j   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                  j                  ||       y )
Nrh   ri   z
jitlink.curj   c                 j    t        j                  d      }|t        |       k  r ||         | |<   y y rR   )r   r   len)rr'   r!   rh   s      r"   kernelz*TestLinker.test_linking_cu.<locals>.kernel   s1    		!A3q6z1Q4y! r$   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=</0	v		! 
	! IIb)MM!ua q5


%%a2r$   c                    t        j                  dd      t        t        dz        }t	        j
                  d      5 }t                t        j                  d|g      fd       }d d d        | j                  t              d	d
       | j                  dt        |d   j                               | j                  dt        |d   j                               y # 1 sw Y   vxY w)Nrh   ri   zwarn.cuT)recordvoid(int32)r   c                      |        y 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=9,-$$D1Q$&XXm4&1 2	 2 	Q$BC*C!,=>5s1Q4<<7HI 21s    +C  C)c                    t        j                  dd      t        t        dz        }| j	                  t
              5 }t        j                  d|g      fd       }d d d        j                  j                  d   }| j                  d|       | j                  d	|       | j                  d
|       y # 1 sw Y   YxY w)Nrh   ri   zerror.cur   r   c                      |        y 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=:-.z*aXXm4&1 2 +
 kkq!1378#>8#> +*s   !B88Cc                     d}| j                  t        |      5  t        j                  ddg      d        }d d d        y # 1 sw Y   y xY w)Nz/Don't know how to link file with extension .cuhvoid()z
header.cuhr   c                       y 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   sC    H##L,?XXhl^4 5 @??   AA
c                     d}| j                  t        |      5  t        j                  ddg      d        }d d d        y # 1 sw Y   y xY w)Nz-Don't know how to link file with no extensionr   datar   c                       y 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   sC    F##L,?XXhfX. / @??r   c                 d    t        t        dz        }t        j                  d|g      d        }y )Nzcuda_include.cur   r   c                       y r   r   r   r$   r"   r   z7TestLinker.test_linking_cu_cuda_include.<locals>.kernel   s    r$   )rp   r   r   rq   )rd   rj   r   s      r"   test_linking_cu_cuda_includez'TestLinker.test_linking_cu_cuda_include   s3    =#445 
($	(	 
)	r$   c                     | j                  t              5 }t        j                  ddg      d        }d d d        | j	                  dj
                  j                         y # 1 sw Y   0xY w)Nvoid(int32[::1])znonexistent.ar   c                     d| d<   y 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XX(/@A B , 	/1A1AB	 ,+s   A$$A-c                     t        j                  t              } |j                  t	        j
                  d      gt        d       }| j                  |j                         d       y)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   sO    
 8878&8&&rxx|?eAh?87792>r$   c                      t        j                  d      t              } |j                  t	        j
                  d      gt        d       }| j                  |j                         d       y )Nr   max_registersr   r   	r   rq   rB   r   rr   r   r&   assertLessEqualr   r   s     r"   test_set_registers_57z TestLinker.test_set_registers_57   T    -488"-.IJ&8&&rxx|?eAh?X99;R@r$   c                      t        j                  d      t              } |j                  t	        j
                  d      gt        d       }| j                  |j                         d       y )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
              }| j                  |j                         d       y )Nr   r   r   )r   r   r   r   rq   rB   r   r   )rd   sigr   s      r"   test_set_registers_eagerz#TestLinker.test_set_registers_eager   sQ    73Q3<ueUEJ2488Cr23NOX99;R@r$   c                     t        t        d d d         } t        j                  |      t              }|j                         }| j                  |t        j                         y 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  sK    73Q3<  488C=!12!446?r$   c                     t        j                  t              } |j                  t	        j
                  d      gt        d       }|j                         }| j                  |d       y )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  sR    8878&8&&rxx|?eAh?";;=!,r$   c                     t        t        d d d   t        t        j                              } t	        j
                  |      t              }|j                         }| j                  |d       y )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  sQ    51:vbhh/0 488C=-";;=#.r$   c                     t        j                  t              }|j                  t	        j
                  dt        j                        t        j                        }|j                         }| j                  |d       y )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  sW    88K('22HHS)2:: 7.GGI#.r$   c                      t        j                  d      t              }|j                         }| j	                  |d       y )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  s6    348823K@88:;*r$   c                 6    t        j                  d      t              }|j                         }|dz   }t	        j
                  |t        j                        }	  |d|f   |       y # t        $ r&}| j                  d|j                         Y d }~y 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    /488./0AB88:ahhuBHH-	3HQXs# 	3MM*AEE22	3s   A) )	B2BBc                 ^   t        t        d d d   t        d d d   t        t        j                              } t	        j
                  |      t              }|j                         }t        j                  t        j                        j                  t        z  }| j                  ||       y 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*  sx    51:uSqSz6"((+;< 488C=-!::<HHRXX&//);		:r$   c                    t        j                  t              }|j                  t	        j
                  t        t        j                        t	        j
                  t        t        j                        t        j                        }|j                         }t	        j                  t        j                        j                  t        z  }| j                  ||       y )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    88K('22HHYbhh/HHYbhh/JJ  .FFHHHRZZ(1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*J$?$ # #C?A
A
A
@-//+
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 :4 4 / & 8 D D D "))Bbjj
)-.`
 	 78N; N; 9N;b zHMMO r$   