
    ܙd                     R   d dl mZ d dlmZ d dlmZ d dlmZ d dlmZ	 d dl
Z
dZdZ ed	           e
j        ej        d
k    d           e
j         ej                     d           G d de
j                                                      Zedk    r e
j                     dS dS )    )nvvm)skip_on_cudasim)utils)ir)bindingNzNcall void @llvm.memset.p0i8.i64(i8* align 4 %arg.x.41, i8 0, i64 %0, i1 false)zFcall void @llvm.memset.p0i8.i64(i8* %arg.x.41, i8 0, i64 %0, i1 false)z"libNVVM not supported in simulator    zCUDA not support for 32-bitz
No libNVVMc                        e Zd Zd Zd Zd ZdS )TestNvvmWithoutCudac                     t          j        t                    }|                     d|           |                                D ]0}d|v r*|                     |d                    dd                     1dS )z
        Test llvm.memset changes in llvm7.
        In LLVM7 the alignment parameter can be implicitly provided as
        an attribute to pointer in the first argument.
        zcall void @llvm.memsetzi32 4, i1 false\) z\s+N)r   llvm100_to_34_iroriginalassertIn
splitlinesassertRegexpMatchesreplace)selffixedlns      Alib/python3.11/site-packages/numba/cuda/tests/nocuda/test_nvvm.pytest_nvvm_memset_fixup_for_34z1TestNvvmWithoutCuda.test_nvvm_memset_fixup_for_34   s     %h//.666""$$ 	 	B'2--(((00f==  	 	    c                     |                      t                    5 }t          j        t                     ddd           n# 1 swxY w Y   |                     t          |j                  d           dS )zy
        We require alignment to be specified as a parameter attribute to the
        dest argument of a memset.
        Nz+No alignment attribute found on memset dest)assertRaises
ValueErrorr   r   missing_alignr   str	exception)r   es     r   +test_nvvm_memset_fixup_for_34_missing_alignz?TestNvvmWithoutCuda.test_nvvm_memset_fixup_for_34_missing_align'   s    
 z** 	1a!-000	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	c!+&&C	E 	E 	E 	E 	Es   AAAc                    t          j        t          j        t          j        d          d          t	          t          d                              }t          j                    }d|_        t          j	        |           t          j
        ||j        d          }d|_        ||_        t          j                    j        |_        t!          j        t%          |                    }t          j        t%          |                    }d                    d t          d          D                       }d| d	                    d
          }|                     ||           d S )N      znvptx64-nvidia-cuda
myconstantTz, c                 ,    g | ]}t          |          S  )r   ).0is     r   
<listcomp>zBTestNvvmWithoutCuda.test_nvvm_accepts_encoding.<locals>.<listcomp>L   s    999c!ff999r   zmyconstant[256] = {}zutf-8)r   Constant	ArrayTypeIntType	bytearrayrangeModuletripler   add_ir_versionGlobalVariabletypeglobal_constantinitializerNVVMdata_layoutllvmparse_assemblyr   llvm_to_ptxjoinencoder   )r   cmgvparsedptxelementsr$   s           r   test_nvvm_accepts_encodingz.TestNvvmWithoutCuda.test_nvvm_accepts_encoding2   s"    KRZ]]C88!%**--/ /IKK(Aq!&,77!	/ $SVV,,s6{{++ 9999eCjj999::8H888??HH
j#&&&&&r   N)__name__
__module____qualname__r   r    rD   r&   r   r   r
   r
      sD          "	E 	E 	E' ' ' ' 'r   r
   __main__)numba.cuda.cudadrvr   numba.cuda.testingr   
numba.corer   llvmliter   r   r9   unittestr   r   skipIfMACHINE_BITSis_availableTestCaser
   rE   mainr&   r   r   <module>rS      s-   # # # # # # . . . . . .             $ $ $ $ $ $ <9 566#r)+HII&T&(((,779' 9' 9' 9' 9'(+ 9' 9' 87 JI 769'x zHMOOOOO r   