U
    ,d$                     @   s   d dl Zd dlmZmZmZ d dlmZmZ d dl	m
Z
 e
jrHd\ZZnd\ZZee ZeefZG dd deZed	kre  dS )
    N)cudafloat32void)unittestCUDATestCase)config)      )2       c                   @   s   e Zd Zdd ZdS )TestCudaMatMulc           
   
   C   s*  t ttd d d d df td d d d df td d d d df dd }tjd tjtjttftjd}tjtjttftjd}t	|}t 
 }| T t ||}t ||}t ||}|ttfttf|f ||| ||| W 5 Q R X t||}	tjj||	dd d S )N   c                 S   s<  t jjttd}t jjttftd}t jj}t jj}t j	j}t j	j}t j
j}	t j
j}
|||	  }|||
  }td}ttD ]}|tk r|tk r| |||t  f |||f< |||t  |f |||f< t   |tk r|tk rttD ] }||||f |||f  7 }qt   q||tk r8|tk r8||||f< d S )N)shapedtyper   )r   ZsharedarraySM_SIZEr   tpbZ	threadIdxxyZblockIdxZblockDimrangebpgnZsyncthreads)ABCZsAZsBZtxtyZbxZbyZbwZbhr   r   accij r   G/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudapy/test_matmul.pycu_square_matrix_mul   s,    
z6TestCudaMatMul.test_func.<locals>.cu_square_matrix_mul*   )r   gh㈵>)Zrtol)r   Zjitr   r   nprandomseedr   r   Z
empty_likestreamZauto_synchronizeZ	to_devicer   r   Zcopy_to_hostdotZtestingZassert_allclose)
selfr!   r   r   r   r&   ZdAZdBZdCZCansr   r   r    	test_func   s    F


zTestCudaMatMul.test_funcN)__name__
__module____qualname__r)   r   r   r   r    r      s   r   __main__)Znumpyr#   Znumbar   r   r   Znumba.cuda.testingr   r   Z
numba.corer   ZENABLE_CUDASIMr   r   r   r   r   r*   mainr   r   r   r    <module>   s   
8