o
    8ήc                     @   s   d dl Z d dlZd dl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mZ edG dd	 d	ejZG d
d deZedkrKe  dS dS )    N)configcuda)TypingError)NVVM)unittestskip_on_cudasimCUDATestCase)irz,This is testing CUDA backend code generationc                   @   s   e Zd Zdd ZdS )TestConstStringCodegenc                 C   s\  ddl m} ddlm}m} |j}|d}d}|||}||| t	dt
|}| t|d ttd g }	t||	d	}
t|
 }|j|||d
}|| t	dt
|}| t|d t||	d}
t|
 }|||}|| t	dt
|}| t|d |t
|d}tt	d|}| t|d d S )Nr   )cuda_target)llvm_to_ptxADDRSPACE_CONSTANT zA Little Brown Foxz7@\"__conststring__.*internal.*constant.*\[19\s+x\s+i8\]      test_insert_const_string)	addrspace"test_insert_string_const_addrspaceasciiz\.const.*__conststring__)numba.cuda.descriptorr   numba.cuda.cudadrv.nvvmr   r   target_contextcreate_moduleinsert_const_stringrefindallstrassertEquallenr	   FunctionTypeIntType
as_pointerFunction	IRBuilderappend_basic_blockinsert_addrspace_convretinsert_string_const_addrspacedecodelist)selfr   r   r   	targetctxmod
textstringgv0resfntyfnbuildermatchesptx r5   P/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_const_string.pytest_const_string   sB   


z(TestConstStringCodegen.test_const_stringN)__name__
__module____qualname__r7   r5   r5   r5   r6   r
   
   s    r
   c                   @   s   e Zd Zdd Zdd Zdd Zdd Zd	d
 Zdd Ze	ddd Z
e	ddd Ze	ddd Ze	ddd Ze	ddd ZdS )TestConstStringc                 C   s$   t jst js| d d S d S d S )Nz+Character sequences unsupported on NVVM 3.4)r   ENABLE_CUDASIMr   	is_nvvm70skipTestr*   r5   r5   r6   skip_on_nvvm34>   s   zTestConstString.skip_on_nvvm34c                 C   s   t  jr| d d S d S )Nz/Character sequences are permitted with NVVM 7.0)r   r=   r>   r?   r5   r5   r6   skip_on_nvvm70C   s   zTestConstString.skip_on_nvvm70c                 C   l   |    tjdd }d}tj|d dd}|d|f | t|}d|d d< d	|d< tj|| d S )
Nc                 S   &   t d}|t| k rd| |< d S d S )Nr   XYZr   gridr   arrir5   r5   r6   
str_assignJ      
zDTestConstString.test_assign_const_unicode_string.<locals>.str_assignr   r   <U12dtyperD   r   r@   r   jitnpzeros
zeros_liketestingassert_equal)r*   rJ   	n_stringsrH   expectedr5   r5   r6    test_assign_const_unicode_stringG      

z0TestConstString.test_assign_const_unicode_stringc                 C   rB   )
Nc                 S   rC   )Nr      XYZrE   rG   r5   r5   r6   bytes_assign^   rK   zCTestConstString.test_assign_const_byte_string.<locals>.bytes_assignr   r   S12rM   r[   rO       rP   )r*   r\   rW   rH   rX   r5   r5   r6   test_assign_const_byte_string[   rZ   z-TestConstString.test_assign_const_byte_stringc                 C   p   |    tjdd }tdtjfdtdfg}td|}|d | tjdd	g|d
}tj	|| d S )Nc                 S   4   d| d d< d| d d< d| d d< d| d d< d S )Nr   r   xABCy   rD   r5   ar5   r5   r6   fr      z=TestConstString.test_assign_const_string_in_record.<locals>.frb   rd   rL   re   r   r   )r   rc   )re   rD   rM   )
r@   r   rQ   rR   rN   int32rS   asarrayrU   assert_array_equalr*   rh   dtrg   	referencer5   r5   r6   "test_assign_const_string_in_recordo      
z2TestConstString.test_assign_const_string_in_recordc                 C   r`   )Nc                 S   ra   )Nr   r   rb      ABCrd   re   r[   r5   rf   r5   r5   r6   rh      ri   z<TestConstString.test_assign_const_bytes_in_record.<locals>.frb   rd   r]   re   rj   )r   rs   )re   r[   rM   )
r@   r   rQ   rR   rN   float32rS   rl   rU   rm   rn   r5   r5   r6   !test_assign_const_bytes_in_record   rr   z1TestConstString.test_assign_const_bytes_in_recordz'Const strings not disallowed on cudasimc                 C   b   |    tjdd }tjddd}| td |d | W d    d S 1 s*w   Y  d S )Nc                 S      d| d< d S )NrD   r   r5   rH   r5   r5   r6   rJ         zOTestConstString.test_const_unicode_string_disallowed_nvvm34.<locals>.str_assignr   rL   rM   .*is a char sequence.*rj   rA   r   rQ   rR   rS   assertRaisesRegexr   )r*   rJ   rH   r5   r5   r6   +test_const_unicode_string_disallowed_nvvm34      
"z;TestConstString.test_const_unicode_string_disallowed_nvvm34c                 C   rv   )Nc                 S   rw   )Nr[   r   r5   rx   r5   r5   r6   r\      ry   zNTestConstString.test_const_byte_string_disallowed_nvvm34.<locals>.bytes_assignr   r]   rM   rz   rj   r{   )r*   r\   rH   r5   r5   r6   (test_const_byte_string_disallowed_nvvm34   r~   z8TestConstString.test_const_byte_string_disallowed_nvvm34c                 C      |    tjdd }tdtjfdtdfg}tjd|d}| td |d	 | W d    d S 1 s9w   Y  d S )
Nc                 S      d| d d< d S )NrD   r   rd   r5   rx   r5   r5   r6   rJ         zVTestConstString.test_record_const_unicode_string_disallowed_nvvm34.<locals>.str_assignrb   rd   rL   r   rM   rz   rj   	rA   r   rQ   rR   rN   rk   rS   r|   r   r*   rJ   ro   rH   r5   r5   r6   2test_record_const_unicode_string_disallowed_nvvm34      
"zBTestConstString.test_record_const_unicode_string_disallowed_nvvm34c                 C   r   )
Nc                 S   r   )Nr[   r   rd   r5   rx   r5   r5   r6   rJ      r   zMTestConstString.test_record_const_bytes_disallowed_nvvm34.<locals>.str_assignrb   rd   r]   r   rM   rz   rj   r   r   r5   r5   r6   )test_record_const_bytes_disallowed_nvvm34   r   z9TestConstString.test_record_const_bytes_disallowed_nvvm34c                 C   s   |    tjdd }tdtdfdtjfg}tdtjfd|fg}tjd|d	}| td
 |d | W d    d S 1 sEw   Y  d S )Nc                 S   s   d| d d d< d S )NrD   r   rd   zr5   rx   r5   r5   r6   rJ      s   z]TestConstString.test_nested_record_const_unicode_string_disallowed_nvvm34.<locals>.str_assignr   rL   rg   rb   rd   r   rM   rz   rj   r   )r*   rJ   subdtro   rH   r5   r5   r6   9test_nested_record_const_unicode_string_disallowed_nvvm34   s   
"zITestConstString.test_nested_record_const_unicode_string_disallowed_nvvm34N)r8   r9   r:   r@   rA   rY   r_   rq   ru   r   r}   r   r   r   r   r5   r5   r5   r6   r;   <   s"    



r;   __main__)r   numpyrR   numbar   r   numba.core.errorsr   r   r   numba.cuda.testingr   r   r   llvmliter	   TestCaser
   r;   r8   mainr5   r5   r5   r6   <module>   s    1 