o
    Y۷i                     @   sx   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ndZeefZG dd deZed	kr:e  dS dS )
    N)cudafloat64void)unittestCUDATestCase)config      c                   @   s   e Zd Zdd ZdS )TestCudaLaplacec              	      s  t jtttddddd  t ttd d d d f td d d d f td d d d f  fdd}tjr@d\}}d}nd	\}}d
}tj||ftjd}tj||ftjd}|}d}d}	t|D ]}
d||
df< d||
df< qdd}t	t	f}||d  ||d  f}t|}t 
 }t ||}t ||}t ||}|	|kr||k r| |jtjk ||||f ||| |j||d |  t| }	|}|}|}|d7 }|	|kr||k sd S d S d S d S )NT)deviceinlinec                 S   s   | |kr| S |S )N )abr   r   Z/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_laplace.pyget_max   s   z3TestCudaLaplace.test_laplace_small.<locals>.get_maxc                    s  t jjttd}t jj}t jj}t jj}t jj}| j	d }| j	d }	t 
d\}
}d|||f< |dkry||d k ry|
dkry|
|	d k ryd| ||
d f | ||
d f  | |d |
f  | |d |
f   |||
f< |||
f | ||
f  |||f< t   td }|dkr||k r |||f ||| |f |||f< |d }t   |dkstd }|dkr||k r|dkrˈ |||f |||| f |||f< |d }t   |dks|dkr|dkr|d |||f< d S d S d S )Ndtyper         g      ?)r   r   )r   sharedarraySM_SIZEr   	threadIdxxyblockIdxshapegridsyncthreadstpb)AAnewerrorerr_smtytxbxbynmijtr   r   r   jocabi_relax_core   sD   

(  &&z=TestCudaLaplace.test_laplace_small.<locals>.jocabi_relax_core)r   r      )   r1   i  r   gư>g      ?r   r   )stream)r   jitr   r   r   ENABLE_CUDASIMnpzerosranger    r2   	to_device
assertTruer   copy_to_hostsynchronizeabsmax)selfr/   NNNMiter_maxr!   r"   r)   tolr#   r,   iterblockdimgriddim
error_gridr2   dAdAnewderror_gridtmpr   r.   r   test_laplace_small   sH   
@(
 z"TestCudaLaplace.test_laplace_smallN)__name__
__module____qualname__rK   r   r   r   r   r
      s    r
   __main__)numpyr5   numbar   r   r   numba.cuda.testingr   r   
numba.corer   r4   r    r   r
   rL   mainr   r   r   r   <module>   s    g