B
      ›\|  ã               @   sn   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 ddd	œZG d
d„ deƒZG dd„ deƒZdS )é    )Úbinding)Úcore)ÚBaseCPUCodegenÚCodeLibrary)Úutilsé   )Únvvmznvptx-nvidia-cudaznvptx64-nvidia-cuda)é    é@   c               @   s,   e Zd Zdd„ Zdd„ Zdd„ Zdd„ Zd	S )
ÚCUDACodeLibraryc             C   s   d S )N© )ÚselfZ	ll_moduler   r   ú1lib/python3.7/site-packages/numba/cuda/codegen.pyÚ_optimize_functions   s    z#CUDACodeLibrary._optimize_functionsc             C   sH   t  ¡ }d|_d|_d|_d|_d|_t  ¡ }| |¡ | 	| j
¡ d S )Nr   FT)ÚllZPassManagerBuilderÚ	opt_levelZdisable_unit_at_a_timeZdisable_unroll_loopsZloop_vectorizeZslp_vectorizeZModulePassManagerZpopulateÚrunÚ_final_module)r   ZpmbZpmr   r   r   Ú_optimize_final_module   s    
z&CUDACodeLibrary._optimize_final_modulec             C   s0   x*| j jD ]}d|jkr
|j dd¡|_q
W d S )NÚ.Ú_)r   Úglobal_variablesÚnameÚreplace)r   Zgvr   r   r   Ú_finalize_specific    s    
z"CUDACodeLibrary._finalize_specificc             C   s   d S )Nr   )r   r   r   r   Úget_asm_str&   s    zCUDACodeLibrary.get_asm_strN)Ú__name__Ú
__module__Ú__qualname__r   r   r   r   r   r   r   r   r      s   r   c               @   s<   e Zd ZdZeZdd„ Zdd„ Zdd„ Zdd	„ Z	d
d„ Z
dS )ÚJITCUDACodegenz£
    This codegen implementation for CUDA actually only generates optimized
    LLVM IR.  Generation of PTX code is done separately (see numba.cuda.compiler).
    c             C   s0   t |jƒg kstdƒ‚tj| _t | j¡| _d S )NzModule isn't empty)	Úlistr   ÚAssertionErrorr   Zdefault_data_layoutÚ_data_layoutr   Zcreate_target_dataZ_target_data)r   Úllvm_moduler   r   r   Ú_init4   s    zJITCUDACodegen._initc             C   s(   t  |¡}ttj |_| jr$| j|_|S )N)ÚlcZModuleÚCUDA_TRIPLEr   ZMACHINE_BITSZtripler"   Zdata_layout)r   r   Z	ir_moduler   r   r   Ú_create_empty_module9   s
    
z#JITCUDACodegen._create_empty_modulec             C   s   t ‚d S )N)ÚNotImplementedError)r   r   r   r   Ú_module_pass_manager@   s    z#JITCUDACodegen._module_pass_managerc             C   s   t ‚d S )N)r(   )r   r#   r   r   r   Ú_function_pass_managerC   s    z%JITCUDACodegen._function_pass_managerc             C   s   d S )Nr   )r   Úmoduler   r   r   Ú_add_moduleF   s    zJITCUDACodegen._add_moduleN)r   r   r   Ú__doc__r   Z_library_classr$   r'   r)   r*   r,   r   r   r   r   r   ,   s   r   N)Zllvmliter   r   Zllvmlite.llvmpyr   r%   Znumba.targets.codegenr   r   Znumbar   Zcudadrvr   r&   r   r   r   r   r   r   Ú<module>   s   