U
    9%e*                     @   sf  d dl mZ d dlmZmZmZmZmZmZ d dl	m
Z
mZmZmZmZmZ d dlmZ d dlmZmZmZ d dlmZ d dlmZmZmZ d dlmZ d d	lmZ d
d Z G dd deZ!G dd deZ"dd Z#edddG dd deZ$edddG dd deZ%G dd deZ&ed(ddZ'ed)ddZ(d*d d!Z)d"d# Z*d$d% Z+G d&d' d'e,Z-dS )+    )ConcreteTemplate)typestypingfuncdescconfigcompilersigutils)sanitize_compile_result_entriesCompilerBaseDefaultPassBuilderFlagsOptionCompileResult)global_compiler_lock)LoweringPassPassManagerregister_pass)NumbaInvalidConfigWarning)IRLegalizationNativeLoweringAnnotateTypes)warn)get_current_devicec                 C   s"   | d krd S t | tst| S d S N)
isinstancedictAssertionError)x r   R/var/www/html/Darija-Ai-API/env/lib/python3.8/site-packages/numba/cuda/compiler.py_nvvm_options_type   s    r    c                   @   s(   e Zd ZeedddZeedddZdS )	CUDAFlagsNzNVVM options)typedefaultdoczCompute Capability)__name__
__module____qualname__r   r    nvvm_optionstuplecompute_capabilityr   r   r   r   r!      s   r!   c                   @   s   e Zd Zedd ZdS )CUDACompileResultc                 C   s   t | S r   )idselfr   r   r   entry_point7   s    zCUDACompileResult.entry_pointN)r%   r&   r'   propertyr/   r   r   r   r   r+   6   s   r+   c                  K   s   t | } tf | S r   )r	   r+   )entriesr   r   r   cuda_compile_result<   s    r2   TF)Zmutates_CFGZanalysis_onlyc                   @   s    e Zd ZdZdd Zdd ZdS )CUDABackendZcuda_backendc                 C   s   t |  d S r   r   __init__r-   r   r   r   r5   F   s    zCUDABackend.__init__c              
   C   sJ   |d }t j|jf|j }t|j|j|jj|j	|j
|j||jd|_dS )zH
        Back-end: Packages lowering output in a compile result
        cr)typing_contexttarget_contextZtyping_errortype_annotationlibrarycall_helper	signaturefndescT)r   r<   return_typeargsr2   	typingctx	targetctxstatusZfail_reasonr9   r:   r;   r=   r6   )r.   stateZloweredr<   r   r   r   run_passI   s    
zCUDABackend.run_passN)r%   r&   r'   _namer5   rD   r   r   r   r   r3   A   s   r3   c                   @   s$   e Zd ZdZdZdd Zdd ZdS )CreateLibraryz
    Create a CUDACodeLibrary for the NativeLowering pass to populate. The
    NativeLowering pass will create a code library if none exists, but we need
    to set it up with nvvm_options from the flags if they are present.
    create_libraryc                 C   s   t |  d S r   r4   r-   r   r   r   r5   g   s    zCreateLibrary.__init__c                 C   s8   |j  }|jj}|jj}|j||d|_|j  dS )N)r(   T)	rA   codegenZfunc_idZfunc_qualnameflagsr(   rG   r:   Zenable_object_caching)r.   rC   rH   namer(   r   r   r   rD   j   s    

zCreateLibrary.run_passN)r%   r&   r'   __doc__rE   r5   rD   r   r   r   r   rF   ]   s   rF   c                   @   s   e Zd Zdd Zdd ZdS )CUDACompilerc                 C   sh   t }td}|| j}|j|j || j}|j|j | | j}|j|j |  |gS )Ncuda)	r   r   Zdefine_untyped_pipelinerC   ZpassesextendZdefine_typed_pipelinedefine_cuda_lowering_pipelinefinalize)r.   ZdpbpmZuntyped_passesZtyped_passesZlowering_passesr   r   r   define_pipelinesv   s    zCUDACompiler.define_pipelinesc                 C   sP   t d}|td |td |td |td |td |  |S )NZcuda_loweringz$ensure IR is legal prior to loweringzannotate typeszcreate libraryznative loweringzcuda backend)r   Zadd_passr   r   rF   r   r3   rP   )r.   rC   rQ   r   r   r   rO      s    z*CUDACompiler.define_cuda_lowering_pipelineN)r%   r&   r'   rR   rO   r   r   r   r   rL   u   s   rL   Nc	                 C   s   |d krt dddlm}	 |	j}
|	j}t }d|_d|_d|_|sH|rNd|_	|rXd|_
|rdd|_nd|_|rtd|_|r~d|_|r||_||_ddlm} |d	  tj|
|| |||i td
}W 5 Q R X |j}|  |S )Nz#Compute Capability must be supplied   cuda_targetTpythonnumpyr   )target_overriderM   )r@   rA   funcr?   r>   rI   localsZpipeline_class)
ValueError
descriptorrU   r7   r8   r!   Z
no_compileZno_cpython_wrapperZno_cfunc_wrapperZ	debuginfoZdbg_directives_onlyZerror_modelZforceinlinefastmathr(   r*   Znumba.core.target_extensionrX   r   Zcompile_extrarL   r:   rP   )pyfuncr>   r?   debuglineinfoinliner]   r(   ccrU   r@   rA   rI   rX   cresr:   r   r   r   compile_cuda   sJ    
	rd   c              
   C   s   |r|rd}t t| ||r"dndd}	t|\}
}|p@tj}t| ||
||||	|d}|jj}|r||s||t	j
kr|td|r|j}n6|j}| j}|j}|j}||j|j|||	||\}}|j|d}||fS )a  Compile a Python function to PTX for a given set of argument types.

    :param pyfunc: The Python function to compile.
    :param sig: The signature representing the function's input and output
                types.
    :param debug: Whether to include debug info in the generated PTX.
    :type debug: bool
    :param lineinfo: Whether to include a line mapping from the generated PTX
                     to the source code. Usually this is used with optimized
                     code (since debug mode would automatically include this),
                     so we want debug info in the LLVM but only the line
                     mapping in the final PTX.
    :type lineinfo: bool
    :param device: Whether to compile a device function. Defaults to ``False``,
                   to compile global kernel functions.
    :type device: bool
    :param fastmath: Whether to enable fast math flags (ftz=1, prec_sqrt=0,
                     prec_div=, and fma=1)
    :type fastmath: bool
    :param cc: Compute capability to compile for, as a tuple
               ``(MAJOR, MINOR)``. Defaults to ``(5, 0)``.
    :type cc: tuple
    :param opt: Enable optimizations. Defaults to ``True``.
    :type opt: bool
    :return: (ptx, resty): The PTX code and inferred return type
    :rtype: tuple
    z{debug=True with opt=True (the default) is not supported by CUDA. This may result in a crash - set debug=False or opt=False.   r   )r]   opt)r_   r`   r]   r(   rb   z'CUDA kernel must have void return type.)rb   )r   r   r   Znormalize_signaturer   ZCUDA_DEFAULT_PTX_CCrd   r<   r>   r   void	TypeErrorr:   r8   __code__co_filenameco_firstlinenoZprepare_cuda_kernelr=   Zget_asm_str)r^   sigr_   r`   devicer]   rb   rf   msgr(   r?   r>   rc   ZrestylibZtgtcodefilenameZlinenumZkernelZptxr   r   r   compile_ptx   s>    


    rr   c              
   C   s    t  j}t| ||||||ddS )zCompile a Python function to PTX for a given set of argument types for
    the current device's compute capabilility. This calls :func:`compile_ptx`
    with an appropriate ``cc`` value for the current device.T)r_   r`   rm   r]   rb   rf   )r   r*   rr   )r^   rl   r_   r`   rm   r]   rf   rb   r   r   r   compile_ptx_for_current_device  s    
   rs   c                 C   s   t | ||jS r   ) declare_device_function_templatekeyrJ   restypeargtypesr   r   r   declare_device_function  s    ry   c                    st   ddl m} |j}|j}tj|f| t|  G  fdddt}tj	| ||d}|
 | |
 | |S )NrS   rT   c                       s   e Zd Z ZgZdS )zBdeclare_device_function_template.<locals>.device_function_templateN)r%   r&   r'   ru   Zcasesr   Zextfnrl   r   r   device_function_template*  s   r{   rv   )r\   rU   r7   r8   r   r<   ExternFunctionr   r   ZExternalFunctionDescriptorZinsert_user_function)rJ   rw   rx   rU   r@   rA   r{   r=   r   rz   r   rt   #  s    
  rt   c                   @   s   e Zd Zdd ZdS )r|   c                 C   s   || _ || _d S r   )rJ   rl   )r.   rJ   rl   r   r   r   r5   7  s    zExternFunction.__init__N)r%   r&   r'   r5   r   r   r   r   r|   6  s   r|   )FFFFNN)FFFFNT)FFFFT).Znumba.core.typing.templatesr   Z
numba.corer   r   r   r   r   r   Znumba.core.compilerr	   r
   r   r   r   r   Znumba.core.compiler_lockr   Znumba.core.compiler_machineryr   r   r   Znumba.core.errorsr   Znumba.core.typed_passesr   r   r   warningsr   Znumba.cuda.apir   r    r!   r+   r2   r3   rF   rL   rd   rr   rs   ry   rt   objectr|   r   r   r   r   <module>   sL     	

!       :      C      

