o
    ùe~k  ã                   @   sˆ   d dl Z d dlmZmZmZ d dlZd dl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dd„ Zdd„ Zdd„ ZdS )é    N)ÚcudaÚcudartÚnvrtcc                 C   sŒ   t | tjƒr| tjjkrtd | ¡ƒ‚d S t | tjƒr*| tjjkr(td | ¡ƒ‚d S t | t	j
ƒr?| t	j
jkr=td | ¡ƒ‚d S td | ¡ƒ‚)NzCuda Error: {}zCudart Error: {}zNvrtc Error: {}zUnknown error type: {})Ú
isinstancer   ZCUresultZCUDA_SUCCESSÚRuntimeErrorÚformatr   ÚcudaError_tZcudaSuccessr   ZnvrtcResultZNVRTC_SUCCESS)Úerr© r
   úLD:\Projects\ConvertPro\env\Lib\site-packages\cuda/tests/test_kernelParams.pyÚ
ASSERT_DRV   s   ÿÿÿr   c                 C   s¤  t  t jj|¡\}}t|ƒ t  t jj|¡\}}t|ƒ t ¡ \}}}t|ƒ |dk}|r/dnd}td|› d|› |› dƒ}	t 	t
 | ¡ddg g ¡\}}
t|ƒ d	|	g}t |
t|ƒ|¡\}t |
¡\}}t|ƒ d
| }t |
|¡\}t|ƒ | ¡ }t|ƒdkr†t|ƒ t|ƒ |r§t |
¡\}}t|ƒ d
| }t |
|¡\}t|ƒ nt |
¡\}}t|ƒ d
| }t |
|¡\}t|ƒ t  tj |¡¡\}}t|ƒ |S )Né   ÚsmZcomputez--gpu-architecture=Ú_Úasciis   allKernelStrings.cur   s   --fmad=falseó    )r   ÚcuDeviceGetAttributeÚCUdevice_attributeZ,CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJORr   Z,CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINORr   ZnvrtcVersionÚbytesZnvrtcCreateProgramÚstrÚencodeZnvrtcCompileProgramÚlenZnvrtcGetProgramLogSizeZnvrtcGetProgramLogÚdecodeÚprintZnvrtcGetCUBINSizeZnvrtcGetCUBINZnvrtcGetPTXSizeZnvrtcGetPTXZcuModuleLoadDataÚnpÚcharÚarray)ZallKernelStringsÚdevr	   ÚmajorÚminorr   Znvrtc_minorZ	use_cubinÚprefixZarch_argÚprogÚoptsZerr_logZlogSizeÚlogÚresultZdataSizeÚdataÚmoduler
   r
   r   Úcommon_nvrtc   sH   
r'   c            
      C   sˆ  t  d¡\} t| ƒ t  d¡\} }t| ƒ t  d|¡\} }t| ƒ d}t||ƒ}t  |d¡\} }t| ƒ t  d¡\} }t| ƒ t  |ddddddd|dd¡\} t| ƒ t  |ddddddd|d d¡\} t| ƒ t	 
¡ }t  |d¡\} }}	t| ƒ |	t	 t	j
¡ks‚J ‚t  ||t	 t	j
¡|¡\} t| ƒ t  |¡\} t| ƒ |jdks¤J ‚t  |¡\} t| ƒ t  |¡\} t| ƒ t  |¡\} t| ƒ d S )Nr   z˜    static __device__ bool isDone;
    extern "C" __global__
    void empty_kernel()
    {
        isDone = true;
        if (isDone) return;
    }
    s   empty_kernelr   )r
   r
   s   isDoneT)r   ÚcuInitr   ÚcuDeviceGetÚcuCtxCreater'   ÚcuModuleGetFunctionÚcuStreamCreateÚcuLaunchKernelÚctypesÚc_boolZcuModuleGetGlobalÚsizeofÚcuMemcpyDtoHAsyncÚcuStreamSynchronizeÚvalueÚcuStreamDestroyÚcuModuleUnloadÚcuCtxDestroy)
r	   ÚcuDeviceÚcontextÚkernelStringr&   ÚkernelÚstreamZisDone_hostZisDonePtr_deviceZisDonePtr_device_sizer
   r
   r   Útest_kernelParams_emptyF   sR   

üür<   c           "      C   s>  t  d¡\}t|ƒ t  d¡\}}t|ƒ t  d|¡\}}t|ƒ | rnt d¡t d¡t d¡t 	d¡t 
d¡t d¡t d¡t d¡t d¡t d	¡t d	¡t d
¡t d
¡t d¡t tdƒ¡t tdƒ¡t d¡f}nddddddddddddddtdƒtdƒdf}tjtjtjtj	tj
tjtjtjtjtjtjtjtjtjtjtjtjf}d}d}d|v r!| r¹|| jn|| }|| tjkrÒ| dtt|ƒƒd d¡}nG|| tjkrå| dtt|ƒƒd¡}n4|| tjkrú| dt|ƒdd … d¡}n|| tjkr| dtt|ƒƒd¡}n| dtt|ƒƒd¡}|d7 }d|v s²t||ƒ}	t  |	d¡\}}
t|ƒ t  d¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj	¡¡\}}t|ƒ t   t !tj
¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ t   t !tj¡¡\}}t|ƒ ||||||||||||||||f}d}|| }|| }t  "|
ddddddd|||fd¡\}t|ƒ t#dd„ |d d… D ƒƒ} t$t%| ƒƒD ]}!t  &| |! ||! t !||! ¡|¡\}t|ƒ qrt  '|¡\}t|ƒ t$t%| ƒƒD ]B}!| r¦||! jn||! }||! tjkrÓ| rÁ|| |! jks¿J ‚qœ|t| |! jd ƒd ksÑJ ‚qœ|| |! jksÝJ ‚qœt  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  (|¡\}t|ƒ t  )|¡\}t|ƒ t  *|	¡\}t|ƒ t  +|¡\}t|ƒ d S )Nr   Tó   Zu   Ä€iÿÿÿéÿ   r   é   é   é   é   çw¾Ÿ/Ý^@l   ï>[= éZ   éH   éU   éR   éB   éA   éV   é!   a      extern "C" __global__
    void basic(bool b,
               char c, wchar_t wc,
               signed char byte, unsigned char ubyte,
               short s, unsigned short us,
               int i, unsigned int ui,
               long l, unsigned long ul,
               long long ll, unsigned long long ull,
               size_t size,
               float f, double d,
               void *p,
               bool *pb,
               char *pc, wchar_t *pwc,
               signed char *pbyte, unsigned char *pubyte,
               short *ps, unsigned short *pus,
               int *pi, unsigned int *pui,
               long *pl, unsigned long *pul,
               long long *pll, unsigned long long *pull,
               size_t *psize,
               float *pf, double *pd)
    {
        assert(b == {});
        assert(c == {});
        assert(wc == {});
        assert(byte == {});
        assert(ubyte == {});
        assert(s == {});
        assert(us == {});
        assert(i == {});
        assert(ui == {});
        assert(l == {});
        assert(ul == {});
        assert(ll == {});
        assert(ull == {});
        assert(size == {});
        assert(f == {});
        assert(d == {});
        assert(p == (void*){});
        *pb = b;
        *pc = c;
        *pwc = wc;
        *pbyte = byte;
        *pubyte = ubyte;
        *ps = s;
        *pus = us;
        *pi = i;
        *pui = ui;
        *pl = l;
        *pul = ul;
        *pll = ll;
        *pull = ull;
        *psize = size;
        *pf = f;
        *pd = d;
    }
    z{}Úfs   basic)NNNNNNNNNNNNNNNNc                 S   s   g | ]}|ƒ ‘qS r
   r
   )Ú.0Z	valueTyper
   r
   r   Ú
<listcomp>8  s    z&kernelParams_basic.<locals>.<listcomp>éÿÿÿÿiè  ),r   r(   r   r)   r*   r.   r/   Úc_charÚc_wcharÚc_byteÚc_ubyteÚc_shortÚc_ushortÚc_intÚc_uintÚc_longÚc_ulongÚ
c_longlongÚc_ulonglongÚc_size_tÚc_floatÚfloatÚc_void_pÚc_doubler3   Úreplacer   ÚordÚintr'   r+   r,   Z
cuMemAllocr0   r-   ÚtupleÚranger   r1   r2   Z	cuMemFreer4   r5   r6   )"Úuse_ctypes_as_valuesr	   r7   r8   ZassertValues_hostZassertTypes_hostZbasicKernelStringÚidxÚvalr&   r:   r;   ZpbZpcZpwcZpbyteZpubyteZpsZpusÚpiZpuiÚplZpulZpllÚpullÚpsizeÚpfÚpdZassertValues_deviceZassertTypes_deviceZbasicKernelValuesZbasicKernelTypesZhost_paramsÚir
   r
   r   ÚkernelParams_basic€   sF  ÷÷
÷9
ô
ø	
ü$$rp   c                   C   ó   t dd d S )NF©rf   ©rp   r
   r
   r
   r   Útest_kernelParams_basicq  ó   rt   c                   C   rq   )NTrr   rs   r
   r
   r
   r   Útest_kernelParams_basic_ctypesu  ru   rv   c                  C   sÎ  t  d¡\} t| ƒ t  d¡\} }t| ƒ t  d|¡\} }t| ƒ t  t jj|¡\} }t| ƒ t 	t
 t
j¡¡\} }t| ƒ t t tj¡tj¡\} }t| ƒ t t tj¡tj¡\} }t| ƒ |rztjj|t d¡tj|dt ¡ tj|df}n/t |d¡\} }t| ƒ t |d¡\} }	t| ƒ tjj|t d¡tj|dt ¡ tj|	df}d t
jd t
jd t
jf}
d|d _d|d _d|d _d}t||ƒ}t  |d	¡\} }t| ƒ t  d¡\} }t| ƒ t  |ddddddd|||
fd¡\} t| ƒ t
 ¡ }t t
 |¡|t
 t
 ¡ ¡tjj|¡\} t| ƒ t   |¡\} t| ƒ t |j!¡}|rf|d |ks/J ‚t"|d ƒt"|d ƒks>J ‚|d j|d
 jksKJ ‚|d j|d
 jksXJ ‚|d j|d
 jkseJ ‚nCtj|d}tj|d}|d |ks{J ‚t"|d ƒt"|ƒksˆJ ‚|d j|jks“J ‚|d j|jksžJ ‚|d j|jks©J ‚t #|¡\} t| ƒ t $|¡\} t| ƒ t $|¡\} t| ƒ t  %|¡\} t| ƒ t  &|¡\} t| ƒ t  '|¡\} t| ƒ d S )Nr   éø   )Z_ptrr   rA   r?   r@   a\      extern "C" __global__
    void structsCuda(cudaError_t err, cudaError_t *perr,
                     cudaSurfaceObject_t surface, cudaSurfaceObject_t *pSurface,
                     dim3 dim, dim3* pdim)
    {
        *perr = err;
        *pSurface = surface;
        pdim->x = dim.x;
        pdim->y = dim.y;
        pdim->z = dim.z;
    }
    s   structsCudarB   )(r   r(   r   r)   r*   r   r   Ú&CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSINGr   Z
cudaMallocr.   r0   rV   ÚcudaHostAllocZcudaSurfaceObject_tÚcudaHostAllocMappedZdim3r   ZcudaErrorUnknownÚcudaHostGetDevicePointerr_   ÚxÚyÚzr'   r+   r,   r-   ZcudaMemcpyAsyncÚ	addressofZcudaMemcpyKindZcudaMemcpyDeviceToHostr2   r3   rc   ZcudaFreeÚcudaFreeHostr4   r5   r6   )r	   r7   r8   ÚuvaSupportedZperrZpSurface_hostZ
pDim3_hostÚkernelValuesZpSurface_deviceZpDim3_deviceÚkernelTypesr9   r&   r:   r;   Zhost_errZcuda_errZsurface_hostZ	dim3_hostr
   r
   r   Útest_kernelParams_types_cuday  sœ   þþþ



ü(r„   c                  C   s¾  t  d¡\} t| ƒ t  d¡\} }t| ƒ t  d|¡\} }t| ƒ t  t jj|¡\} }t| ƒ d}t||ƒ}t  	|d¡\} }t| ƒ t  
d¡\} }t| ƒ G dd„ dtjƒ}t t |¡tj¡\} }	t| ƒ |ro|dƒ|	f}
nt |	d¡\} }t| ƒ |dƒ|f}
d tjf}t  |ddddddd||
|fd¡\} t| ƒ t  |¡\} t| ƒ | |	¡}|
d j|jksµJ ‚t |	¡\} t| ƒ t  |¡\} t| ƒ t  |¡\} t| ƒ t  |¡\} t| ƒ d S )Nr   zÂ    struct testStruct {
        int value;
    };

    extern "C" __global__
    void structCustom(struct testStruct src, struct testStruct *dst)
    {
        dst->value = src.value;
    }
    s   structCustomc                   @   ó   e Zd ZdejfgZdS )z3test_kernelParams_struct_custom.<locals>.testStructr3   N©Ú__name__Ú
__module__Ú__qualname__r.   rV   Ú_fields_r
   r
   r
   r   Ú
testStruct  ó    r‹   rB   r   )r   r(   r   r)   r*   r   r   rx   r'   r+   r,   r.   Ú	Structurer   ry   r0   rz   r{   r_   r-   r2   Úfrom_addressr3   r€   r4   r5   r6   )r	   r7   r8   r   r9   r&   r:   r;   r‹   ÚpStruct_hostr‚   ÚpStruct_devicerƒ   Zstruct_sharedr
   r
   r   Útest_kernelParams_struct_customâ  sT   

ü
r‘   c                 C   sö  t  d¡\}t|ƒ t  d¡\}}t|ƒ t  d|¡\}}t|ƒ t  t jj|¡\}}t|ƒ d}t||ƒ}t  	|d¡\}}t|ƒ t  
d¡\}}t|ƒ G dd„ dtjƒ}	t t tj¡tj¡\}}
t|ƒ t t tj¡tj¡\}}t|ƒ t t |	¡tj¡\}}t|ƒ |r¦t d¡t |
¡t tdƒ¡t |¡|	dƒt |¡f}n?t |
d¡\}}t|ƒ t |d¡\}}t|ƒ t |d¡\}}t|ƒ t d¡t |¡t tdƒ¡t |¡|	dƒt |¡f}tjt|ƒ ƒ }tt|ƒƒD ]}t || ¡||< qót  |ddddddd|| rt |¡n|d¡\}t|ƒ t  |¡\}t|ƒ |d jtj |
¡jks4J ‚|d	 jtj |¡jksCJ ‚|d
 j|	 |¡jksQJ ‚t |¡\}t|ƒ t  |¡\}t|ƒ t  |¡\}t|ƒ t   |¡\}t|ƒ d S )Nr   á'      struct testStruct {
        int value;
    };
    extern "C" __global__
    void testkernel(int i, int *pi,
                    float f, float *pf,
                    struct testStruct s, struct testStruct *ps)
    {
        *pi = i;
        *pf = f;
        ps->value = s.value;
    }
    ó
   testkernelc                   @   r…   )z>kernelParams_buffer_protocol_ctypes_common.<locals>.testStructr3   Nr†   r
   r
   r
   r   r‹   H  rŒ   r‹   r   rC   rB   r?   rA   )!r   r(   r   r)   r*   r   r   rx   r'   r+   r,   r.   r   r   ry   r0   rV   rz   r]   r_   r^   r{   r   re   r   r-   r2   r3   rŽ   r€   r4   r5   r6   )Úpass_by_addressr	   r7   r8   r   r9   r&   r:   r;   r‹   Ú	pInt_hostÚpFloat_hostr   r‚   ÚpInt_deviceÚpFloat_devicer   ÚpackagedParamsrg   r
   r
   r   Ú*kernelParams_buffer_protocol_ctypes_common&  sv   
þþürš   c                   C   s   t dd t dd d S )NT)r”   F)rš   r
   r
   r
   r   Ú(test_kernelParams_buffer_protocol_ctypes|  s   
r›   c                  C   s<  t  d¡\} t| ƒ t  d¡\} }t| ƒ t  d|¡\} }t| ƒ t  t jj|¡\} }t| ƒ d}t||ƒ}t  	|d¡\} }t| ƒ t  
d¡\} }t| ƒ t dtjfg¡}t t tj¡jtj¡\} }	t| ƒ t t tj¡jtj¡\} }
t| ƒ t |jtj¡\} }t| ƒ |r»tjdtjdtj|	gtjdtjtdƒtjdtj|
gtjdt dg|¡tj|gtjdf}nTt |	d¡\} }t| ƒ t |
d¡\} }t| ƒ t |d¡\} }t| ƒ tjdtjdtj|gtjdtjtdƒtjdtj|gtjdt dg|¡tj|gtjdf}tjd	d
„ |D ƒtjd}t  |ddddddd||d¡\} t| ƒ t  |¡\} t| ƒ G dd„ dƒ}|d t ||	dƒ¡ksPJ ‚|d t ||
dƒ¡ks_J ‚|d d tj||dƒ|dd kstJ ‚t |¡\} t| ƒ t  |¡\} t| ƒ t  |¡\} t| ƒ t  |¡\} t| ƒ d S )Nr   r’   r“   r3   r   )ÚdtyperC   rB   c                 S   s   g | ]}|j j‘qS r
   )r.   r%   )rM   Úargr
   r
   r   rN   »  s    z;test_kernelParams_buffer_protocol_numpy.<locals>.<listcomp>c                   @   s   e Zd Zdd„ ZdS )zFtest_kernelParams_buffer_protocol_numpy.<locals>.numpy_address_wrapperc                 S   s   |df|ddœ| _ d S )NF)r   )r%   ÚtypestrÚshape)Z__array_interface__)ÚselfÚaddressrž   r
   r
   r   Ú__init__È  s   þzOtest_kernelParams_buffer_protocol_numpy.<locals>.numpy_address_wrapper.__init__N)r‡   rˆ   r‰   r¢   r
   r
   r
   r   Únumpy_address_wrapperÇ  s    r£   z<i4r?   z<f4rA   )r   r(   r   r)   r*   r   r   rx   r'   r+   r,   r   rœ   Zint32r   ry   Úitemsizerz   Zfloat32r   Zuint32Zuint64r^   r{   r-   r2   r€   r4   r5   r6   )r	   r7   r8   r   r9   r&   r:   r;   r‹   r•   r–   r   r‚   r—   r˜   r   r™   r£   r
   r
   r   Ú'test_kernelParams_buffer_protocol_numpy€  st   
"þ"þü*r¥   )Zpytestr   r   r   Únumpyr   r.   r   r'   r<   rp   rt   rv   r„   r‘   rš   r›   r¥   r
   r
   r
   r   Ú<module>   s    ,: riDV