ó
\K]c           @` sñ   d  d l  m Z m Z m Z d  d l m Z d  d l m Z d  d l m	 Z	 m
 Z
 d  d l m Z d  d l m Z m Z e	 d ƒ e j e j d k d	 ƒ e j e j ƒ  d
 ƒ d e
 e j f d „  ƒ  Yƒ ƒ ƒ Z e d k rí e j ƒ  n  d S(   i    (   t   absolute_importt   print_functiont   division(   t   compile_kernel(   t   nvvm(   t   skip_on_cudasimt   SerialMixin(   t   unittest_support(   t   typest   utilss"   libNVVM not supported in simulatori    s   CUDA not support for 32-bits
   No libNVVMt   TestNvvmWithoutCudac           B` s   e  Z d  „  Z d „  Z RS(   c         C` so   d „  } t  | d t j d d d … f d d ƒ} | j j j } t j | ƒ } |  j d | j	 d ƒ ƒ d S(	   sp   
        A simple test to exercise nvvm.llvm_to_ptx()
        to trigger issues with mismatch NVVM API.
        c         S` s   d |  d <d  S(   Ni{   i    (    (   t   x(    (    s@   lib/python2.7/site-packages/numba/cuda/tests/nocuda/test_nvvm.pyt   foo   s    t   argsNi   t   linkR   t   ascii(    (
   R   R   t   int32t   _funct   ptxt   llvmirR   t   llvm_to_ptxt   assertInt   decode(   t   selfR   t   cukernR   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/nocuda/test_nvvm.pyt   test_nvvm_llvm_to_ptx   s
    	+c         C` sú   d „  } t  | d t j d d d … f d d ƒ} | j j j } |  j d | ƒ t j | ƒ } |  j d | ƒ x? | j	 ƒ  D]1 } d | k r |  j
 | d j d d	 ƒ ƒ q q Wx? | j	 ƒ  D]1 } d | k rÁ |  j
 | d
 j d d	 ƒ ƒ qÁ qÁ Wd S(   s²   
        Test llvm.memset changes in llvm7.
        In LLVM7 the alignment parameter can be implicitly provided as
        an attribute to pointer in the first argument.
        c         S` s(   x! t  |  j ƒ D] } d |  | <q Wd  S(   Ni    (   t   ranget   size(   R   t   i(    (    s@   lib/python2.7/site-packages/numba/cuda/tests/nocuda/test_nvvm.pyR   "   s    R   Ni   R   s   call void @llvm.memsets   i64 %\d+, i1 false\)t    s   \s+s   i32 4, i1 false\)(    (   R   R   R   R   R   R   R   R   t   llvm39_to_34_irt
   splitlinest   assertRegexpMatchest   replace(   R   R   R   t   originalt   fixedt   ln(    (    s@   lib/python2.7/site-packages/numba/cuda/tests/nocuda/test_nvvm.pyt   test_nvvm_memset_fixup   s     	+(   t   __name__t
   __module__R   R%   (    (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/nocuda/test_nvvm.pyR
   
   s   	t   __main__N(   t
   __future__R    R   R   t   numba.cuda.compilerR   t   numba.cuda.cudadrvR   t   numba.cuda.testingR   R   t   numbaR   t   unittestR   R	   t   skipIft   MACHINE_BITSt   is_availablet   TestCaseR
   R&   t   main(    (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/nocuda/test_nvvm.pyt   <module>   s   	"1