
    xKgb                     z    d dl Zd dlmZmZmZ d dlmZmZ d Z	 G d de      Z
edk(  r ej                          yy)    N)cudafloat32void)unittestCUDATestCasec                    t        j                  t        j                  | | z        j                  | |       t         j                        }t        j                  t        j                  |       dz   |j
                        }||fS )Ndtyper   )nparrayarangereshaper   r
   )nABs      g/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_nondet.pygenerate_inputr      sZ    
1q5!))!Q/rzzBA
1!1Aa4K    c                       e Zd Zd Zy)TestCudaNonDetc           	      ^   t        j                  t        t        ddddf   t        ddddf   t        dd             d        }d}t	        |      \  }}t        j                  |j                  |j                        }d}d}t        j                  |      }t        j                  |      }	t        j                  |d      }
 |||f   |
||	       t        j                  |t        j                  |            }t
        j                  j                  |
j                         |       y)	ziTest issue with loop not running due to bad sign-extension at the for
        loop precondition.
        Nc                    t        j                  d      \  }}t         j                  j                  t         j                  j                  z  }t         j                  j
                  t         j                  j
                  z  }| j                  d   }| j                  d   }t        |||      D ]&  }	t        |||      D ]  }
||
|	f   ||	   z  | |
|	f<    ( y )N   r      )r   gridgridDimxblockDimyshaperange)cabstartXstartYgridXgridYheightwidthr   r   s              r   diagproductz0TestCudaNonDet.test_for_pre.<locals>.diagproduct   s    !YYq\NFFLLNNT]]__4ELLNNT]]__4EWWQZFGGAJE6552vv7A1g!nAadG 8 3r      r	   )    r,   )r   r   F)copy)r   jitr   r   r   r   emptyr    r
   	to_devicedotdiagtestingassert_array_almost_equalcopy_to_host)selfr+   Nr   r   FblockdimgriddimdAdBdFEs               r   test_for_prezTestCudaNonDet.test_for_pre   s    
 
$wq!t}gadmWQZ@	A		- 
B		- a 1HHQWWAGG,^^A^^A^^AE*&GX%&r2r2FF1bggaj!


,,R__->Br   N)__name__
__module____qualname__r@    r   r   r   r      s     Cr   r   __main__)numpyr   numbar   r   r   numba.cuda.testingr   r   r   r   rA   mainrD   r   r   <module>rJ      s?     % % 5!C\ !CH zHMMO r   