
    wg                         d dl mZmZmZmZ d dlmZ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 dZdZ ed       G d	 d
e             Z G d de      Zedk(  r ej.                          yy)    )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                   X     e Zd Z fdZ fdZd Zd Zd Zd Zd Z	d Z
d	 Zd
 Z xZS )TestCudaDriverc                 ,   t         |           | j                  t        t        j
                        dkD         t	        j                         | _        | j                  j                  }|j                  \  }}|dk\  rt        | _        y t        | _        y )Nr      )supersetUp
assertTruelenr   gpusget_contextcontextdevicecompute_capabilityptx2ptxptx1)selfr   ccmajor_	__class__s       n/home/mcse/projects/flask/flask-venv/lib/python3.12/site-packages/numba/cuda/tests/cudadrv/test_cuda_driver.pyr   zTestCudaDriver.setUpA   si    GLL)A-.**,$$..
a<DHDH    c                 &    t         |           | `y N)r   tearDownr   )r    r#   s    r$   r(   zTestCudaDriver.tearDownL   s    Lr%   c                 t   | j                   j                  | j                        }|j                  d      }t	        dz         }| j                   j                  t        |            }t        ||t        |             |j                  }d}t        j                  r3t        t        |            }t        j                  j                  |      }t        |j                   ddddddd||g
       t#        ||t        |             t%        |      D ]  \  }}| j'                  ||        |j)                          y )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_basicz%TestCudaDriver.test_cuda_driver_basicP   s   //9&&'9:&&ve}5vufUm4**!!3s8$C__--f5Fhoo1a1ae	 	uffUm4e$ 	#DAqQ"	# 	r%   c                    | j                   j                  | j                        }|j                  d      }t	        dz         }| j                   j                         }|j                         5  | j                   j                  t        |            }t        ||t        |      |       |j                  }t        j                  rt        t        |            }t        |j                   ddddddd|j                   |g
       d d d        t#        |t        |      |       t%        |      D ]  \  }}| j'                  ||        y # 1 sw Y   GxY w)Nr*   r+   )r?   r,   r   )r   r-   r   r.   r   create_streamauto_synchronizer/   r   r   r0   r1   r2   r   r3   r
   r6   r   r7   r8   )	r    r:   r;   r<   r?   r=   r>   r@   rA   s	            r$   "test_cuda_driver_stream_operationsz1TestCudaDriver.test_cuda_driver_stream_operationsm   s   //9&&'9:++-$$& 	!\\**6%=9F65&-G..C%%s3x((//q!q! --%!	! 	uffUmFCe$ 	#DAqQ"	##	! 	!s   .BE

Ec                     | j                   j                         }| j                  dt        |             | j	                  dt        |             | j                  |       | j                  |j                         y )NzDefault CUDA streamr   )	r   get_default_streamassertInreprr8   r3   r   assertFalseexternalr    dss     r$   test_cuda_driver_default_streamz.TestCudaDriver.test_cuda_driver_default_stream   s[    \\,,.+T"X6CG$ 	%r%   c                     | j                   j                         }| j                  dt        |             | j	                  dt        |             | j                  |       | j                  |j                         y )NzLegacy default CUDA streamr,   )	r   get_legacy_default_streamrI   rJ   r8   r3   r   rK   rL   rM   s     r$   &test_cuda_driver_legacy_default_streamz5TestCudaDriver.test_cuda_driver_legacy_default_stream   sY    \\3352DH=CG$%r%   c                     | j                   j                         }| j                  dt        |             | j	                  dt        |             | j                  |       | j                  |j                         y )NzPer-thread default CUDA streamr   )	r   get_per_thread_default_streamrI   rJ   r8   r3   r   rK   rL   rM   s     r$   *test_cuda_driver_per_thread_default_streamz9TestCudaDriver.test_cuda_driver_per_thread_default_stream   sY    \\7796RACG$%r%   c                 h   | j                   j                         }| j                  dt        |             | j	                  dt        |             | j	                  dt        |             | j                  dt        |             | j                  |       | j                  |j                         y )NzCUDA streamDefaultExternalr   )
r   rD   rI   rJ   assertNotInassertNotEqualr3   r   rK   rL   )r    ss     r$   test_cuda_driver_streamz&TestCudaDriver.test_cuda_driver_stream   s    LL&&(mT!W-DG,T!W-As1v&$r%   c                    t         j                  r!t        j                  d      }t	        |      }n?t        j                         }t        j                  t        |      d       |j                  }| j                  j                  |      }| j                  dt        |             | j                  dt        |             | j                  |t	        |             | j                  |       | j                  |j                          y )Nr   zExternal CUDA streamefault)r1   r2   r	   cuStreamCreater3   r   	cu_streamr   valuer   create_external_streamrI   rJ   rY   r8   r   rL   )r    r6   r>   r[   s       r$    test_cuda_driver_external_streamz/TestCudaDriver.test_cuda_driver_external_stream   s     !!**1-Ff+C%%'F!!%-3,,CLL//4,d1g647+c!f%

#r%   c                 j   | j                   j                  | j                        }|j                  d      }| j                   j	                  |dd      }| j                  |dkD         d }| j                   j                  ||dd      \  }}| j                  |dkD         | j                  |dkD         y )Nr*      r   c                     | S r'    )bss    r$   b2dz6TestCudaDriver.test_cuda_driver_occupancy.<locals>.b2d   s    Ir%   )r   r-   r   r.   $get_active_blocks_per_multiprocessorr   get_max_potential_block_size)r    r:   r;   ra   ri   gridblocks          r$   test_cuda_driver_occupancyz)TestCudaDriver.test_cuda_driver_occupancy   s    //9&&'9:AA(BEsL	"	 ll??#@CSJeq!	"r%   )__name__
__module____qualname__r   r(   rB   rF   rO   rR   rU   r\   rc   rn   __classcell__)r#   s   @r$   r   r   ?   s5    	:#8	&&&%$(#r%   r   c                       e Zd Zd Zy)
TestDevicec                     d}|dz  }|dz  }|dz  }d| d| d| d| d| d}t        j                         j                  }| j                  |j                  |       y )Nz[0-9a-f]{%d}         z^GPU--$)r   r   r   assertRegexuuid)r    hh4h8h12uuid_formatdevs          r$   test_device_get_uuidzTestDevice.test_device_get_uuid   sr     UU"fbT2$at1RD#a8!!#**;/r%   N)ro   rp   rq   r   rg   r%   r$   rt   rt      s    0r%   rt   __main__N)ctypesr   r   r   r   numba.cuda.cudadrv.driverr   r   r	   r
   numba.cuda.cudadrvr   r   r1   numba.cuda.testingr   r   r   r   r   r   rt   ro   mainrg   r%   r$   <module>r      s{    1 16 6 A A 5 .0: ?@P#\ P# AP#f0 0. zHMMO r%   