
    J/Ph8                        d dl mZmZmZmZ d dlmZ d dlmZ d dl	m
Z
mZmZ d dlZd dlmZ ddlmZmZ  ej        d	ej        fd
ej        dfg          Z G d de          Z G d de          Zedk    r e
j                     dS dS )    )cudaint32float64void)TypingError)types)unittestCUDATestCaseskip_on_cudasimN)numpy_support   )test_struct_model_type
TestStructij)      c                   >    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	S )
TestSharedMemoryIssuec                     t          j        d          d             t           j        fd            } |d                      d S )NT)devicec                  R    t           j                            dt                    } d S Nr   dtyper   sharedarrayr   )	inner_arrs    _/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_sm.pyinnerzGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.inner   s     ))!5)99III    c                  h    t           j                            dt                    }               d S r   r   )	outer_arrr!   s    r    outerzGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.outer   s,    ))!5)99IEGGGGGr"   r   r   )r   jit)selfr%   r!   s     @r    "test_issue_953_sm_linkage_conflictz8TestSharedMemoryIssue.test_issue_953_sm_linkage_conflict   sf    					: 	: 
		: 
	 	 	 	 
	 	dr"   c                     t           j        fd            }t          j        dt          j                  } |d         |           |                     |d         |           d S )Nc                 h    t           j                            t                    }|j        | d<   d S Nr   r   )r   r   r   r   size)aarrshapes     r    sz9TestSharedMemoryIssue._check_shared_array_size.<locals>.s   s,    +##E#77C8AaDDDr"   r   r   r&   r   )r   r'   npzerosr   assertEqual)r(   r0   expectedr1   results    `   r    _check_shared_array_sizez.TestSharedMemoryIssue._check_shared_array_size   sq    		 	 	 	 
	 !28,,,$H-----r"   c                 2    |                      dd           d S Nr   r7   r(   s    r    %test_issue_1051_shared_size_broken_1dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_1d&   s    %%a+++++r"   c                 2    |                      dd           d S )N)r   r      r:   r;   s    r    %test_issue_1051_shared_size_broken_2dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_2d)   s    %%fa00000r"   c                 2    |                      dd           d S )N)r   r         r:   r;   s    r    %test_issue_1051_shared_size_broken_3dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_3d,   s    %%i44444r"   c                     t           j        fd            }t          j        dt          j                  } |d         |           |                     |d         |           d S )Nc                 ^    t           j                                      }|j        | d<   d S r,   )r   r   r   r-   )r.   r/   r0   tys     r    r1   z>TestSharedMemoryIssue._check_shared_array_size_fp16.<locals>.s1   s,    +##E#44C8AaDDDr"   r   r   r&   r   )r   r'   r2   r3   float16r4   )r(   r0   r5   rF   r1   r6   s    ` `  r    _check_shared_array_size_fp16z3TestSharedMemoryIssue._check_shared_array_size_fp160   sw    		 	 	 	 	 
	 !2:...$H-----r"   c                     |                      ddt          j                   |                      ddt          j                   d S r9   )rH   r   rG   r2   r;   s    r    test_issue_fp16_supportz-TestSharedMemoryIssue.test_issue_fp16_support:   s<    **1a???**1a<<<<<r"   c                     dd}dd}t           j        fd            }t          j        |t          j                  }t          j        |          } |||f         |           t          j                     dS )z
        Test issue of warp misalign address due to nvvm not knowing the
        alignment(? but it should have taken the natural alignment of the type)
        r   0   rA   r   c                     t           j                            ft                    }t           j                            dt                    }t           j        j        }d}t                    D ]}||||f         z  }|d         |z   | d<   d S )N   r   )r   r   r   r   	threadIdxxrange)d_block_costs
s_featuress_initialcostrO   
predictionr   examples_per_blocknum_weightss         r    
costs_funcz9TestSharedMemoryIssue.test_issue_2393.<locals>.costs_funcH   s    **,>+L+24 4J K--a99M(IJ;'' 7 7jA66

,Q/*<M!r"   r   N)r   r'   r2   r3   r   	to_devicesynchronize)r(   
num_blocksthreads_per_blockrX   block_costsrR   rV   rW   s         @@r    test_issue_2393z%TestSharedMemoryIssue.test_issue_2393>   s    
 
		= 	= 	= 	= 	= 
	= hz<<<{331
:001-@@@r"   N)__name__
__module____qualname__r)   r7   r<   r?   rC   rH   rJ   r^    r"   r    r   r      s        
 
 
. . ., , ,1 1 15 5 5. . .= = =    r"   r   c                       e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	 Zd
 Zd Zd Zd Z ed          d             Z ed          d             ZdS )TestSharedMemoryc                 j   t          |          }dt          |z            }t          j        |j                  t
          j        fd            }t          j        |          } ||f         ||           |                                }t          j
                            ||           d S )N   c                 x   t           j                            	          }t           j        j        }t           j        j        }t           j        j        }||z  |z   }|t          |           k     r| |         ||<   t          j                     |dk    r#t          	          D ]}||         |||z  |z   <   d S d S r,   
r   r   r   rO   rP   blockIdxblockDimlensyncthreadsrQ   )
rP   ysmtxbxbdr   r   dtnthreadss
           r    use_sm_chunk_copyz8TestSharedMemory._test_shared.<locals>.use_sm_chunk_copyj   s    ""82"66B!BBB R"A3q66zz12 Qwwx + +A%'UAb2gkNN w+ +r"   )rk   intnps
from_dtyper   r   r'   device_array_likecopy_to_hostr2   testingassert_array_equal)	r(   r/   nelemnblocksrt   d_resulthost_resultrr   rs   s	          @@r    _test_sharedzTestSharedMemory._test_shared_   s     Ceh&''^CI&&		+ 	+ 	+ 	+ 	+ 
	+& )#..,'8+,S(;;;++--

%%c;77777r"   c                 F   t          j        dt                    }t          t	          |                    D ]S}|||         _        t          j        dt           j                  }|                    dd          |z  ||         _	        T| 
                    |           d S )N   r   r>   r   r   )r2   recarrayrecordwith2darrayrQ   rk   r   arangefloat32reshaper   r   )r(   r/   rP   r   s       r    test_shared_recarrayz%TestSharedMemory.test_shared_recarray   s    k#%6777s3xx 	+ 	+ACFH	%rz222AyyA*CFHH#r"   c                     t           j                            ddt           j                  }|                     |           d S )Nr   )   )r-   r   )r2   randomrandintbool_r   )r(   r/   s     r    test_shared_boolz!TestSharedMemory.test_shared_bool   s:    irx@@#r"   c                     |j         |j        j        z  } |ddd|f         |           t          j                            ||           d S )Nr   r   )r-   r   itemsizer2   rz   r{   )r(   funcr/   r5   nshareds        r    _test_dynshared_slicez&TestSharedMemory._test_dynshared_slice   sQ    
 (SY//Q1gs###

%%h44444r"   c                     t           j        d             }t          j        dt          j                  }t          j        ddgt          j                  }|                     |||           d S )Nc                     t           j                            dt                    }|dd         }|dd         }d|d<   d|d<   |d         | d<   |d         | d<   d S Nr   r   r   r   r   rP   dynsmemsm1sm2s       r    slice_writez@TestSharedMemory.test_dynshared_slice_write.<locals>.slice_write   sf    k'''77G!A#,C!A#,CCFCF1:AaD1:AaDDDr"   r   r   r   r   r'   r2   r3   r   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_writez+TestSharedMemory.test_dynshared_slice_write   sm    		 	 
	 hq)))8QF"(333"";X>>>>>r"   c                     t           j        d             }t          j        dt          j                  }t          j        ddgt          j                  }|                     |||           d S )Nc                     t           j                            dt                    }|dd         }|dd         }d|d<   d|d<   |d         | d<   |d         | d<   d S r   r   r   s       r    
slice_readz>TestSharedMemory.test_dynshared_slice_read.<locals>.slice_read   sf    k'''77G!A#,C!A#,CGAJGAJq6AaDq6AaDDDr"   r   r   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_readz*TestSharedMemory.test_dynshared_slice_read   sm    		 	 
	 hq)))8QF"(333"":sH=====r"   c                     t           j        d             }t          j        dt          j                  }t          j        g dt          j                  }|                     |||           d S )Nc                     t           j                            dt                    }|dd         }|dd         }d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   d S )Nr   r   r   r   r   r   r   s       r    slice_diff_sizeszJTestSharedMemory.test_dynshared_slice_diff_sizes.<locals>.slice_diff_sizes   s|    k'''77G!A#,C!A#,CGAJGAJGAJq6AaDq6AaDq6AaDDDr"   r   r   )r   r   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_diff_sizesz0TestSharedMemory.test_dynshared_slice_diff_sizes   sp     

	 
	 

	 hq)))8IIIRX666""#3S(CCCCCr"   c                     t           j        d             }t          j        dt          j                  }t          j        g dt          j                  }|                     |||           d S )Nc                    t           j                            dt                    }|dd         }|dd         }d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   d S )Nr   r   r   r   rA   r   r   r   s       r    slice_overlapzDTestSharedMemory.test_dynshared_slice_overlap.<locals>.slice_overlap   s    k'''77G!A#,C!A#,CGAJGAJGAJGAJq6AaDq6AaDq6AaDq6AaDq6AaDDDr"      r   )r   r   r   r   rA   r   )r(   r   r/   r5   s       r    test_dynshared_slice_overlapz-TestSharedMemory.test_dynshared_slice_overlap   sm    		 	 
	 hq)))8OOO28<<<""=#x@@@@@r"   c                     t           j        d             }t          j        dt          j                  }t          j        g dt          j                  }|                     |||           d S )Nc                    t           j                            dt                    }|dd         }|dd         }d|d<   d|d<   d|d<   d|d<   d|d<   d|d	<   d|d<   d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d	         | d	<   |d         | d<   d S )
Nr   r   r   r   rA   r>   c   r   r   r   r   s       r    
slice_gapsz>TestSharedMemory.test_dynshared_slice_gaps.<locals>.slice_gaps   s    k'''77G!A#,C!A#,C GAJGAJGAJGAJGAJGAJGAJCFCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDDDr"   rN   r   )r   r   r   r   r   rA   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_gapsz*TestSharedMemory.test_dynshared_slice_gaps   sr     
	 	 
	6 hq)))8444BHEEE"":sH=====r"   c                     t           j        d             }t          j        dt          j                  }t          j        g dt          j                  }|                     |||           d S )Nc                     t           j                            dt                    }|dd d         }|ddd         }d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   d S )Nr   r   r   r   r   rA   r   r   s       r    slice_write_backwardszTTestSharedMemory.test_dynshared_slice_write_backwards.<locals>.slice_write_backwards  s    k'''77G!%R%.C!Ab&/CCFCFCFCF1:AaD1:AaD1:AaD1:AaDDDr"   rA   r   )r   r   rA   r   r   )r(   r   r/   r5   s       r    $test_dynshared_slice_write_backwardsz5TestSharedMemory.test_dynshared_slice_write_backwards  sp     
	 	 
	 hq)))8LLL999""#8#xHHHHHr"   c                     t           j        d             }t          j        dt          j                  }t          j        g dt          j                  }|                     |||           d S )Nc                 F   t           j                            dt                    }|d d d         }d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   d S )	Nr   r   r   r   r   r   rA   r   r   rP   r   r   s      r    slice_nonunit_stridezRTestSharedMemory.test_dynshared_slice_nonunit_stride.<locals>.slice_nonunit_stride!  s    k'''77G##A#,C GAJGAJGAJGAJGAJGAJCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDDDr"   r>   r   )r   r   r   r   r   r   r   )r(   r   r/   r5   s       r    #test_dynshared_slice_nonunit_stridez4TestSharedMemory.test_dynshared_slice_nonunit_stride  ss     
	 	 
	. hq)))8111BBB""#7hGGGGGr"   c                     t           j        d             }t          j        dt          j                  }t          j        g dt          j                  }|                     |||           d S )Nc                 F   t           j                            dt                    }|dd d         }d|d<   d|d<   d|d<   d|d<   d|d	<   d|d
<   d|d<   d|d<   d|d<   |d         | d<   |d         | d<   |d         | d<   |d         | d<   |d	         | d	<   |d
         | d
<   d S )Nr   r   r   r   r   r   r   rA   r   r   r   s      r    slice_nonunit_reverse_stridezbTestSharedMemory.test_dynshared_slice_nonunit_reverse_stride.<locals>.slice_nonunit_reverse_stride@  s    k'''77G"&b&/C GAJGAJGAJGAJGAJGAJCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDDDr"   r>   r   )r   r   r   r   r   r   r   )r(   r   r/   r5   s       r    +test_dynshared_slice_nonunit_reverse_stridez<TestSharedMemory.test_dynshared_slice_nonunit_reverse_stride=  ss     
	 	 
	. hq)))8111BBB""#?hOOOOOr"   c                   
 t          j        d          }t          |          }d}t          ||z            }t	          j        |j                  
||j        j        z  }t          |dz            }t          j	        
fd            }t          j
        |          } |||d|f         |||           |                                }	t           j                            ||	           d S )Nr   rf   r   c                    t           j                            d          }|d|         }|||dz           }t           j        j        }t           j        j        }t           j        j        }||z  |z   }	|	t          |           k     r ||k     r| |	         ||<   n| |	         |||z
  <   t          j                     |dk    r7t          |          D ])}
||
         |||z  |
z   <   ||
         |||z  |
z   |z   <   (d S d S )Nr   r   r   rh   )rP   rm   	chunksizer   r   r   ro   rp   rq   r   r   rr   s              r    sm_slice_copyz7TestSharedMemory.test_issue_5073.<locals>.sm_slice_copyk  s   k'''44G!I+&C)IM12C!BBB R"A3q66zz	>>dCGG*+A$CY' Qwwy)) 8 8A%(VAb2gkN14QAb2gkI-.. w8 8r"   r   )r2   r   rk   ru   rv   rw   r   r   r   r'   rx   ry   rz   r{   )r(   r/   r|   rs   r}   r   r   r   r~   r   rr   s             @r    test_issue_5073z TestSharedMemory.test_issue_5073\  s     iooCeh&''^CI&&SY//1%%			8 	8 	8 	8 
	82 )#..4gxG34S(INNN++--

%%c;77777r"   zCan't check typing in simulatorc                    d}d }|                      t          |          5   t          j        t	                                |           d d d            n# 1 swxY w Y   d}d }|                      t          |          5   t          j        t	                                |           d d d            d S # 1 swxY w Y   d S )Nz+.*Cannot infer the type of variable 'arr'.*c                  l    t           j                            dt          j        d                    } d S )N
   Or   )r   r   r   r2   r   r/   s    r    unsupported_typezBTestSharedMemory.test_invalid_array_type.<locals>.unsupported_type  s(    +##Bbhsmm#<<CCCr"   z*.*Invalid NumPy dtype specified: 'int33'.*c                  H    t           j                            dd          } d S )Nr   int33r   )r   r   r   r   s    r    invalid_string_typezETestSharedMemory.test_invalid_array_type.<locals>.invalid_string_type  s     +##Bg#66CCCr"   )assertRaisesRegexr   r   r'   r   )r(   rgxr   r   s       r    test_invalid_array_typez(TestSharedMemory.test_invalid_array_type  sO   ;	= 	= 	=##K55 	/ 	/DHTVV-...	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ 	/ ;	7 	7 	7##K55 	2 	2DHTVV0111	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2s#   *AAA*B99B= B=z+Struct model array unsupported in simulatorc           	         dt          j        t          t          d d d         t          d d d                             fd            }t	          j        fd          }t	          j        fd          } |df         ||           t          |          D ]!\  }}|                     ||z
  dz
             "t          |          D ]$\  }}|                     ||z
  dz
  dz             %d S )N@   r   c                    t           j                            t                    }t          j        d          }|z
  dz
  }|t          |           k     rz|t          |          k     rit          t          |          t          |dz                      }|||<   t          j                     ||         j	        | |<   ||         j
        ||<   d S d S d S )Nr   r   r   )r   r   r   r   gridrk   r   r   rl   rP   rm   )outxoutyr/   r   riobjrs   s         r    write_then_reverse_read_staticzVTestSharedMemory.test_struct_model_type_static.<locals>.write_then_reverse_read_static  s     +##H4J#KKC	!AA!B3t99}}SYY q5Q<<88A """b')Qb')Q }r"   r   r   r   )r   r'   r   r   r2   r3   	enumerater4   )r(   r   arrxarryr   rP   rm   rs   s          @r    test_struct_model_type_staticz.TestSharedMemory.test_struct_model_type_static  s#   	$uSSqSz51:..	/	/	$ 	$ 	$ 	$ 
0	/	$" x7333x73333&q({3D$???dOO 	2 	2DAqQ1q 01111dOO 	8 	8DAqQA!1Q 67777	8 	8r"   N)r_   r`   ra   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rb   r"   r    rd   rd   ^   s8       "8 "8 "8H    5 5 5? ? ?"> > >"D D D(A A A,!> !> !>FI I I,H H H>P P P>,8 ,8 ,8\ _6772 2 872 _BCC8 8 DC8 8 8r"   rd   __main__)numbar   r   r   r   numba.core.errorsr   
numba.corer   numba.cuda.testingr	   r
   r   numpyr2   numba.npr   rv   extensions_usecasesr   r   r   r   r   r   rd   r_   mainrb   r"   r    <module>r      sZ   , , , , , , , , , , , , ) ) ) ) ) )       F F F F F F F F F F     ) ) ) ) ) ) C C C C C C C CBHsBHo"BJ79 : : L L L L LL L L L^Z8 Z8 Z8 Z8 Z8| Z8 Z8 Z8z
 zHMOOOOO r"   