
    sh                         S SK JrJrJrJr  S SKJrJrJrJ	r	  S SK
JrJrJr  S SKJrJr  S SKJr  SrSr\" S5       " S	 S
\5      5       r " S S\5      r\S:X  a  \R.                  " 5         gg)    )byrefc_intc_void_psizeof)host_to_devicedevice_to_hostdriverlaunch_kernel)devicesdrvapir	   )unittestCUDATestCase)skip_on_cudasima  
    .version 1.4
    .target sm_10, map_f64_to_f32

    .entry _Z10helloworldPi (
    .param .u64 __cudaparm__Z10helloworldPi_A)
    {
    .reg .u32 %r<3>;
    .reg .u64 %rd<6>;
    .loc	14	4	0
$LDWbegin__Z10helloworldPi:
    .loc	14	6	0
    cvt.s32.u16 	%r1, %tid.x;
    ld.param.u64 	%rd1, [__cudaparm__Z10helloworldPi_A];
    cvt.u64.u16 	%rd2, %tid.x;
    mul.lo.u64 	%rd3, %rd2, 4;
    add.u64 	%rd4, %rd1, %rd3;
    st.global.s32 	[%rd4+0], %r1;
    .loc	14	7	0
    exit;
$LDWend__Z10helloworldPi:
    } // _Z10helloworldPi
a  
.version 3.0
.target sm_20
.address_size 64

    .file	1 "/tmp/tmpxft_000012c7_00000000-9_testcuda.cpp3.i"
    .file	2 "testcuda.cu"

.entry _Z10helloworldPi(
    .param .u64 _Z10helloworldPi_param_0
)
{
    .reg .s32 	%r<3>;
    .reg .s64 	%rl<5>;


    ld.param.u64 	%rl1, [_Z10helloworldPi_param_0];
    cvta.to.global.u64 	%rl2, %rl1;
    .loc 2 6 1
    mov.u32 	%r1, %tid.x;
    mul.wide.u32 	%rl3, %r1, 4;
    add.s64 	%rl4, %rl2, %rl3;
    st.global.u32 	[%rl4], %r1;
    .loc 2 7 2
    ret;
}
z,CUDA Driver API unsupported in the simulatorc                   d   ^  \ rS rSrU 4S jrU 4S jrS rS rS rS r	S r
S	 rS
 rS rSrU =r$ )TestCudaDriver?   c                 .  > [         TU ]  5         U R                  [        [        R
                  5      S:  5        [        R                  " 5       U l        U R                  R                  nUR                  u  p#US:  a  [        U l        g [        U l        g )Nr      )supersetUp
assertTruelenr   gpusget_contextcontextdevicecompute_capabilityptx2ptxptx1)selfr   ccmajor_	__class__s       }/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudadrv/test_cuda_driver.pyr   TestCudaDriver.setUpA   sg    GLL)A-.**,$$..
a<DHDH    c                 &   > [         TU ]  5         U ?g N)r   tearDownr   )r!   r$   s    r%   r*   TestCudaDriver.tearDownL   s    Lr'   c                 |   U R                   R                  U R                  5      nUR                  S5      n[        S-  " 5       nU R                   R                  [        U5      5      n[        XC[        U5      5        UR                  nSn[        R                  (       a3  [        [        U5      5      n[        R                  R                  U5      n[        UR                   SSSSSSSUU/5
        [#        X4[        U5      5        [%        U5       H  u  pxU R'                  Xx5        M     UR)                  5         g )N_Z10helloworldPid   r      )r   create_module_ptxr   get_functionr   memallocr   r   device_ctypes_pointer_driverUSE_NV_BINDINGr   intbindingCUstreamr
   handler   	enumerateassertEqualunload)	r!   modulefunctionarraymemoryptrstreamivs	            r%   test_cuda_driver_basic%TestCudaDriver.test_cuda_driver_basicP   s    //9&&'9:&&ve}5vfUm4**!!3s8$C__--f5Fhoo1a1ae	 	ufUm4e$DAQ" % 	r'   c                    U R                   R                  U R                  5      nUR                  S5      n[        S-  " 5       nU R                   R                  5       nUR                  5          U R                   R                  [        U5      5      n[        XS[        U5      US9  UR                  n[        R                  (       a  [        [        U5      5      n[        UR                   SSSSSSSUR                   U/5
        S S S 5        [#        UW[        U5      US9  [%        U5       H  u  pxU R'                  Xx5        M     g ! , (       d  f       NI= f)Nr-   r.   )rB   r/   r   )r   r0   r   r1   r   create_streamauto_synchronizer2   r   r   r3   r4   r5   r   r6   r
   r9   r   r:   r;   )	r!   r=   r>   r?   rB   r@   rA   rC   rD   s	            r%   "test_cuda_driver_stream_operations1TestCudaDriver.test_cuda_driver_stream_operationsm   s
   //9&&'9:++-$$&\\**6%=9F6&-G..C%%s3x((//q!q! --%! ' 	uffUmFCe$DAQ" %# '&s   /BE		
Ec                     U R                   R                  5       nU R                  S[        U5      5        U R	                  S[        U5      5        U R                  U5        U R                  UR                  5        g )NzDefault CUDA streamr   )	r   get_default_streamassertInreprr;   r6   r   assertFalseexternalr!   dss     r%   test_cuda_driver_default_stream.TestCudaDriver.test_cuda_driver_default_stream   s[    \\,,.+T"X6CG$ 	%r'   c                     U R                   R                  5       nU R                  S[        U5      5        U R	                  S[        U5      5        U R                  U5        U R                  UR                  5        g )NzLegacy default CUDA streamr/   )	r   get_legacy_default_streamrN   rO   r;   r6   r   rP   rQ   rR   s     r%   &test_cuda_driver_legacy_default_stream5TestCudaDriver.test_cuda_driver_legacy_default_stream   sY    \\3352DH=CG$%r'   c                     U R                   R                  5       nU R                  S[        U5      5        U R	                  S[        U5      5        U R                  U5        U R                  UR                  5        g )NzPer-thread default CUDA streamr   )	r   get_per_thread_default_streamrN   rO   r;   r6   r   rP   rQ   rR   s     r%   *test_cuda_driver_per_thread_default_stream9TestCudaDriver.test_cuda_driver_per_thread_default_stream   sY    \\7796RACG$%r'   c                 h   U R                   R                  5       nU R                  S[        U5      5        U R	                  S[        U5      5        U R	                  S[        U5      5        U R                  S[        U5      5        U R                  U5        U R                  UR                  5        g )NzCUDA streamDefaultExternalr   )
r   rH   rN   rO   assertNotInassertNotEqualr6   r   rP   rQ   )r!   ss     r%   test_cuda_driver_stream&TestCudaDriver.test_cuda_driver_stream   s    LL&&(mT!W-DG,T!W-As1v&$r'   c                 $   [         R                  (       a"  [        R                  " S5      n[	        U5      nOA[
        R                  " 5       n[        R                  " [        U5      S5        UR                  nU R                  R                  U5      nU R                  S[        U5      5        U R                  S[        U5      5        U R                  U[	        U5      5        U R                  U5        U R                  UR                   5        g )Nr   zExternal CUDA streamefault)r4   r5   r	   cuStreamCreater6   r   	cu_streamr   valuer   create_external_streamrN   rO   ra   r;   r   rQ   )r!   r9   rA   rc   s       r%    test_cuda_driver_external_stream/TestCudaDriver.test_cuda_driver_external_stream   s     !!**1-Ff+C%%'F!!%-3,,CLL//4,d1g647+c!f%

#r'   c                 f   U R                   R                  U R                  5      nUR                  S5      nU R                   R	                  USS5      nU R                  US:  5        S nU R                   R                  X$SS5      u  pVU R                  US:  5        U R                  US:  5        g )Nr-      r   c                     U $ r)    )bss    r%   b2d6TestCudaDriver.test_cuda_driver_occupancy.<locals>.b2d   s    Ir'   )r   r0   r   r1   $get_active_blocks_per_multiprocessorr   get_max_potential_block_size)r!   r=   r>   rj   rs   gridblocks          r%   test_cuda_driver_occupancy)TestCudaDriver.test_cuda_driver_occupancy   s    //9&&'9:AA(BEsL	"	 ll??@CSJq!	"r'   )r   r   )__name__
__module____qualname____firstlineno__r   r*   rE   rJ   rT   rX   r\   rd   rl   ry   __static_attributes____classcell__)r$   s   @r%   r   r   ?   s:    	:#8	&&&%$(# #r'   r   c                       \ rS rSrS rSrg)
TestDevice   c                     SnUS-  nUS-  nUS-  nSU SU SU SU SU S3n[         R                  " 5       R                  nU R                  UR                  U5        g )Nz[0-9a-f]{%d}         z^GPU--$)r   r   r   assertRegexuuid)r!   hh4h8h12uuid_formatdevs          r%   test_device_get_uuidTestDevice.test_device_get_uuid   sr     UU"fbT2$at1RD#a8!!#**;/r'   rq   N)r{   r|   r}   r~   r   r   rq   r'   r%   r   r      s    0r'   r   __main__N)ctypesr   r   r   r   numba.cuda.cudadrv.driverr   r   r	   r
   numba.cuda.cudadrvr   r   r4   numba.cuda.testingr   r   r   r    r   r   r   r{   mainrq   r'   r%   <module>r      sy    1 16 6 A A 5 .0: ?@P#\ P# AP#f0 0. zMMO r'   