ó
\K]c           @   s   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 i d d 6d	 d
 6Z d e f d     YZ d e f d     YZ d S(   i˙˙˙˙(   t   binding(   t   core(   t   BaseCPUCodegent   CodeLibrary(   t   utilsi   (   t   nvvms   nvptx-nvidia-cudai    s   nvptx64-nvidia-cudai@   t   CUDACodeLibraryc           B   s,   e  Z d    Z d   Z d   Z d   Z RS(   c         C   s   d  S(   N(    (   t   selft	   ll_module(    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _optimize_functions   s    c         C   sf   t  j   } d | _ t | _ t | _ t | _ t | _ t  j	   } | j
 |  | j |  j  d  S(   Ni   (   t   llt   PassManagerBuildert	   opt_levelt   Falset   disable_unit_at_a_timet   Truet   disable_unroll_loopst   loop_vectorizet   slp_vectorizet   ModulePassManagert   populatet   runt   _final_module(   R   t   pmbt   pm(    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _optimize_final_module   s    					c         C   sE   x> |  j  j D]0 } d | j k r | j j d d  | _ q q Wd  S(   Nt   .t   _(   R   t   global_variablest   namet   replace(   R   t   gv(    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _finalize_specific    s    c         C   s   d  S(   N(   t   None(   R   (    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   get_asm_str&   s    (   t   __name__t
   __module__R	   R   R    R"   (    (    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyR      s   			t   JITCUDACodegenc           B   sA   e  Z d  Z e Z d   Z d   Z d   Z d   Z d   Z	 RS(   sŁ   
    This codegen implementation for CUDA actually only generates optimized
    LLVM IR.  Generation of PTX code is done separately (see numba.cuda.compiler).
    c         C   sF   t  | j  g  k s! t d   t j |  _ t j |  j  |  _ d  S(   Ns   Module isn't empty(	   t   listR   t   AssertionErrorR   t   default_data_layoutt   _data_layoutR
   t   create_target_datat   _target_data(   R   t   llvm_module(    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _init4   s    !c         C   s;   t  j |  } t t j | _ |  j r7 |  j | _ n  | S(   N(   t   lct   Modulet   CUDA_TRIPLER   t   MACHINE_BITSt   tripleR)   t   data_layout(   R   R   t	   ir_module(    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _create_empty_module9   s
    	c         C   s
   t   d  S(   N(   t   NotImplementedError(   R   (    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _module_pass_manager@   s    c         C   s
   t   d  S(   N(   R6   (   R   R,   (    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _function_pass_managerC   s    c         C   s   d  S(   N(    (   R   t   module(    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   _add_moduleF   s    (
   R#   R$   t   __doc__R   t   _library_classR-   R5   R7   R8   R:   (    (    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyR%   ,   s   				N(   t   llvmliteR    R
   t   llvmlite.llvmpyR   R.   t   numba.targets.codegenR   R   t   numbaR   t   cudadrvR   R0   R   R%   (    (    (    s1   lib/python2.7/site-packages/numba/cuda/codegen.pyt   <module>   s   

