U
    ?hï  ã                   @   s¤   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
„ d
eƒƒZG dd„ deƒZedkr e ¡  dS )é    )Ú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                       sh   e Zd Z‡ fdd„Z‡ fdd„Zdd„ Zdd„ Zd	d
„ Zdd„ Zdd„ Z	dd„ Z
dd„ Zdd„ Z‡  ZS )ÚTestCudaDriverc                    sT   t ƒ  ¡  |  ttjƒdk¡ t ¡ | _| jj}|j	\}}|dkrJt
| _nt| _d S )Nr   é   )ÚsuperÚsetUpÚ
assertTrueÚlenr
   ZgpusÚget_contextÚcontextÚdeviceZcompute_capabilityÚptx2ÚptxÚptx1)Úselfr   ZccmajorÚ_©Ú	__class__© ú[/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/tests/cudadrv/test_cuda_driver.pyr   A   s    


zTestCudaDriver.setUpc                    s   t ƒ  ¡  | `d S ©N)r   ÚtearDownr   )r   r   r   r    r"   L   s    
zTestCudaDriver.tearDownc           	      C   sÂ   | j  | j¡}| d¡}td ƒ }| j  t|ƒ¡}t||t|ƒƒ |j}d}t	j
rjtt|ƒƒ}t	j |¡}t|jddddddd||gƒ
 t||t|ƒƒ t|ƒD ]\}}|  ||¡ q | ¡  d S )NÚ_Z10helloworldPiéd   r   é   )r   Úcreate_module_ptxr   Úget_functionr   Úmemallocr   r   Údevice_ctypes_pointerÚ_driverÚUSE_NV_BINDINGr   ÚintZbindingZCUstreamr	   Úhandler   Ú	enumerateÚassertEqualZunload)	r   ÚmoduleÚfunctionÚarrayÚmemoryÚptrÚstreamÚiÚvr   r   r    Útest_cuda_driver_basicP   s2    

    ûz%TestCudaDriver.test_cuda_driver_basicc           	      C   sÒ   | j  | j¡}| d¡}td ƒ }| j  ¡ }| ¡ b | j  t|ƒ¡}t	||t|ƒ|d |j
}tjrrtt|ƒƒ}t|jddddddd|j|gƒ
 W 5 Q R X t||t|ƒ|d t|ƒD ]\}}|  ||¡ q¸d S )Nr#   r$   )r5   r%   r   )r   r&   r   r'   r   Úcreate_streamZauto_synchronizer(   r   r   r)   r*   r+   r   r,   r	   r-   r   r.   r/   )	r   r0   r1   r2   r5   r3   r4   r6   r7   r   r   r    Ú"test_cuda_driver_stream_operationsm   s0    



    ûz1TestCudaDriver.test_cuda_driver_stream_operationsc                 C   sD   | j  ¡ }|  dt|ƒ¡ |  dt|ƒ¡ |  |¡ |  |j¡ d S )NzDefault CUDA streamr   )	r   Zget_default_streamÚassertInÚreprr/   r,   r   ÚassertFalseÚexternal©r   Zdsr   r   r    Útest_cuda_driver_default_stream‰   s
    

z.TestCudaDriver.test_cuda_driver_default_streamc                 C   sD   | j  ¡ }|  dt|ƒ¡ |  dt|ƒ¡ |  |¡ |  |j¡ d S )NzLegacy default CUDA streamr%   )	r   Zget_legacy_default_streamr;   r<   r/   r,   r   r=   r>   r?   r   r   r    Ú&test_cuda_driver_legacy_default_stream”   s
    

z5TestCudaDriver.test_cuda_driver_legacy_default_streamc                 C   sD   | j  ¡ }|  dt|ƒ¡ |  dt|ƒ¡ |  |¡ |  |j¡ d S )NzPer-thread default CUDA streamr   )	r   Zget_per_thread_default_streamr;   r<   r/   r,   r   r=   r>   r?   r   r   r    Ú*test_cuda_driver_per_thread_default_streamœ   s
    

z9TestCudaDriver.test_cuda_driver_per_thread_default_streamc                 C   sd   | j  ¡ }|  dt|ƒ¡ |  dt|ƒ¡ |  dt|ƒ¡ |  dt|ƒ¡ |  |¡ |  |j	¡ d S )NzCUDA streamZDefaultZExternalr   )
r   r9   r;   r<   ÚassertNotInZassertNotEqualr,   r   r=   r>   )r   Úsr   r   r    Útest_cuda_driver_stream¤   s    

z&TestCudaDriver.test_cuda_driver_streamc                 C   sŽ   t jrt d¡}t|ƒ}nt ¡ }t t|ƒd¡ |j}| j	 
|¡}|  dt|ƒ¡ |  dt|ƒ¡ |  |t|ƒ¡ |  |¡ |  |j¡ d S )Nr   zExternal CUDA streamZefault)r*   r+   r   ZcuStreamCreater,   r   Z	cu_streamr   Úvaluer   Zcreate_external_streamr;   r<   rC   r/   r   r>   )r   r-   r4   rD   r   r   r    Ú test_cuda_driver_external_stream®   s    


z/TestCudaDriver.test_cuda_driver_external_streamc                 C   st   | j  | j¡}| d¡}| j  |dd¡}|  |dk¡ dd„ }| j  ||dd¡\}}|  |dk¡ |  |dk¡ d S )Nr#   é€   r   c                 S   s   | S r!   r   )Úbsr   r   r    Úb2dÊ   s    z6TestCudaDriver.test_cuda_driver_occupancy.<locals>.b2d)r   r&   r   r'   Z$get_active_blocks_per_multiprocessorr   Zget_max_potential_block_size)r   r0   r1   rF   rJ   ÚgridÚblockr   r   r    Útest_cuda_driver_occupancyÂ   s    
 ÿ
 ÿz)TestCudaDriver.test_cuda_driver_occupancy)Ú__name__Ú
__module__Ú__qualname__r   r"   r8   r:   r@   rA   rB   rE   rG   rM   Ú__classcell__r   r   r   r    r   ?   s   
r   c                   @   s   e Zd Zdd„ ZdS )Ú
TestDevicec                 C   s\   d}|d }|d }|d }d|› d|› d|› d|› d|› d}t  ¡ j}|  |j|¡ d S )Nz[0-9a-f]{%d}é   é   é   z^GPU-ú-ú$)r
   r   r   ÚassertRegexÚuuid)r   ÚhZh4Zh8Zh12Zuuid_formatÚdevr   r   r    Útest_device_get_uuidÔ   s    $
zTestDevice.test_device_get_uuidN)rN   rO   rP   r\   r   r   r   r    rR   Ó   s   rR   Ú__main__N)Úctypesr   r   r   r   Znumba.cuda.cudadrv.driverr   r   r   r	   Znumba.cuda.cudadrvr
   r   r*   Znumba.cuda.testingr   r   r   r   r   r   rR   rN   Úmainr   r   r   r    Ú<module>   s    