
    `i                         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                     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                   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                 D   t                                                       |                     t          t          j                  dk               t	          j                    | _        | 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__s       }/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/numba/cuda/tests/cudadrv/test_cuda_driver.pyr   zTestCudaDriver.setUpA   sz    GL))A-...*,,$.
a<<DHHHDHHH    c                 L    t                                                       | `d S N)r   tearDownr   )r    r#   s    r$   r(   zTestCudaDriver.tearDownL   s"    LLLr%   c                    | j                             | j                  }|                    d          }t	          dz              }| j                             t          |                    }t          ||t          |                     |j        }d}t          j
        r;t          t          |                    }t          j                            |          }t          |j        ddddddd||g
  
         t#          ||t          |                     t%          |          D ]\  }}|                     ||           |                                 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vs	            r$   test_cuda_driver_basicz%TestCudaDriver.test_cuda_driver_basicP   s;   //99&&'9::&&ve}}55vufUmm444*! 	63s88$$C_--f55Fho1a1ae	 	 	 	uffUmm444e$$ 	# 	#DAqQ""""r%   c                    | j                             | j                  }|                    d          }t	          dz              }| j                                         }|                                5  | j                             t          |                    }t          ||t          |          |           |j
        }t          j        rt          t          |                    }t          |j        ddddddd|j        |g
  
         d d d            n# 1 swxY w Y   t#          ||t          |          |           t%          |          D ]\  }}|                     ||           d S )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&&'9::++--$$&& 	! 	!\**6%==99F65&--GGGG.C% )s3xx(((/q!q! -%! ! !	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	! 	uffUmmFCCCCe$$ 	# 	#DAqQ""""	# 	#s   3BDD Dc                 "   | j                                         }|                     dt          |                     |                     dt          |                     |                     |           |                     |j                   d S )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"XX666CGG$$$ 	%%%%%r%   c                 "   | j                                         }|                     dt          |                     |                     dt          |                     |                     |           |                     |j                   d S )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   sz    \33552DHH===CGG$$$%%%%%r%   c                 "   | j                                         }|                     dt          |                     |                     dt          |                     |                     |           |                     |j                   d S )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   sz    \77996RAAACGG$$$%%%%%r%   c                    | j                                         }|                     dt          |                     |                     dt          |                     |                     dt          |                     |                     dt          |                     |                     |           |                     |j	                   d S )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    L&&((mT!WW---DGG,,,T!WW---As1vv&&&$$$$$r%   c                 B   t           j        r$t          j        d          }t	          |          }n<t          j                    }t          j        t          |          d           |j        }| j	        
                    |          }|                     dt          |                     |                     dt          |                     |                     |t	          |                     |                     |           |                     |j                   d S )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++CC%''F!%--333,CL//44,d1gg666477+++c!ff%%%
#####r%   c                    | j                             | j                  }|                    d          }| j                             |dd          }|                     |dk               d }| j                             ||dd          \  }}|                     |dk               |                     |dk               d S )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&&'9::AA(BEsL L	"""	 	 	 l??#@CSJ Jeq!!!	"""""r%   )__name__
__module____qualname__r   r(   rB   rF   rO   rR   rU   r\   rc   rn   __classcell__)r#   s   @r$   r   r   ?   s        	 	 	 	 	      :# # #8	& 	& 	&& & && & &% % %$ $ $(# # # # # # #r%   r   c                       e Zd Zd ZdS )
TestDevicec                     d}|dz  }|dz  }|dz  }d| d| d| d| d| d}t          j                    j        }|                     |j        |           d S )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   s     UU"f8b8828888R88#888!##*;/////r%   N)ro   rp   rq   r   rg   r%   r$   rt   rt      s#        0 0 0 0 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      sW   1 1 1 1 1 1 1 1 1 1 1 16 6 6 6 6 6 6 6 6 6 6 6 A A A A A A A A A A 5 5 5 5 5 5 5 5 . . . . . .0: ?@@P# P# P# P# P#\ P# P# A@P#f0 0 0 0 0 0 0 0. zHMOOOOO r%   