o
    Y۷i$                     @   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r$d\ZZnd\ZZee ZeefZG dd deZed	krBe  dS 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fttf|f ||| ||| W d    n1 sw   Y  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 ]K}|tk rd|tk rd| |||t  f |||f< |||t  |f |||f< t   |tk r|tk rttD ]}||||f |||f  7 }qtt   q>|tk r|tk r||||f< d S d S d S )N)shapedtyper   )r   sharedarraySM_SIZEr   tpb	threadIdxxyblockIdxblockDimrangebpgnsyncthreads)ABCsAsBtxtybxbybwbhr   r   accij r+   Y/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_matmul.pycu_square_matrix_mul   s.   
z6TestCudaMatMul.test_func.<locals>.cu_square_matrix_mul*   )r   gh㈵>)rtol)r   jitr   r   nprandomseedr   r   
empty_likestreamauto_synchronize	to_devicer   r   copy_to_hostdottestingassert_allclose)
selfr-   r   r   r   r5   dAdBdCCansr+   r+   r,   	test_func   s    F


zTestCudaMatMul.test_funcN)__name__
__module____qualname__rA   r+   r+   r+   r,   r      s    r   __main__)numpyr1   numbar   r   r   numba.cuda.testingr   r   
numba.corer   ENABLE_CUDASIMr   r   r   r   r   rB   mainr+   r+   r+   r,   <module>   s    
8