o
    ۾i                     @   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rRe  dS 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                    sV   t    | ttjdk t | _| jj}|j	\}}|dkr&t
| _d S t| _d S )Nr      )supersetUp
assertTruelenr
   gpusget_contextcontextdevicecompute_capabilityptx2ptxptx1)selfr   ccmajor_	__class__ ]/home/ubuntu/.local/lib/python3.10/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
r5tt|}t	j|}t|jddddddd||g
 t||t| t|D ]
\}}| || qP|  d S )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vr"   r"   r#   test_cuda_driver_basicP   s*   

z%TestCudaDriver.test_cuda_driver_basicc           	      C   s   | j | j}|d}td  }| j  }| 6 | j t|}t	||t||d |j
}tjr9tt|}t|jddddddd|j|g
 W d    n1 sSw   Y  t||t||d t|D ]
\}}| || qfd S )Nr&   r'   )r;   r(   r   )r   r)   r   r*   r   create_streamauto_synchronizer+   r   r   r,   r-   r.   r   r/   r	   r2   r   r3   r4   )	r   r6   r7   r8   r;   r9   r:   r<   r=   r"   r"   r#   "test_cuda_driver_stream_operationsm   s,   



z1TestCudaDriver.test_cuda_driver_stream_operationsc                 C   D   | j  }| dt| | dt| | | | |j d S )NzDefault CUDA streamr   )	r   get_default_streamassertInreprr4   r/   r   assertFalseexternalr   dsr"   r"   r#   test_cuda_driver_default_stream   s
   

z.TestCudaDriver.test_cuda_driver_default_streamc                 C   rB   )NzLegacy default CUDA streamr(   )	r   get_legacy_default_streamrD   rE   r4   r/   r   rF   rG   rH   r"   r"   r#   &test_cuda_driver_legacy_default_stream   
   

z5TestCudaDriver.test_cuda_driver_legacy_default_streamc                 C   rB   )NzPer-thread default CUDA streamr   )	r   get_per_thread_default_streamrD   rE   r4   r/   r   rF   rG   rH   r"   r"   r#   *test_cuda_driver_per_thread_default_stream   rM   z9TestCudaDriver.test_cuda_driver_per_thread_default_streamc                 C   sd   | j  }| dt| | dt| | dt| | dt| | | | |j	 d S )NzCUDA streamDefaultExternalr   )
r   r?   rD   rE   assertNotInassertNotEqualr/   r   rF   rG   )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 streamefault)r-   r.   r   cuStreamCreater/   r   	cu_streamr   valuer   create_external_streamrD   rE   rR   r4   r   rG   )r   r2   r:   rT   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*   $get_active_blocks_per_multiprocessorr   get_max_potential_block_size)r   r6   r7   rY   r^   gridblockr"   r"   r#   test_cuda_driver_occupancy   s   

z)TestCudaDriver.test_cuda_driver_occupancy)__name__
__module____qualname__r   r%   r>   rA   rJ   rL   rO   rU   r[   rc   __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   hh4h8h12uuid_formatdevr"   r"   r#   test_device_get_uuid   s   $
zTestDevice.test_device_get_uuidN)rd   re   rf   rv   r"   r"   r"   r#   rh      s    rh   __main__N)ctypesr   r   r   r   numba.cuda.cudadrv.driverr   r   r   r	   numba.cuda.cudadrvr
   r   r-   numba.cuda.testingr   r   r   r   r   r   rh   rd   mainr"   r"   r"   r#   <module>   s     