o Td\0@sddlmZddlmZmZmZmZmZmZddl m Z m Z m Z m Z mZmZddlmZddlmZmZmZmZddlmZmZddlmZmZmZddlmZdd l m!Z!d d Z"Gd d d e Z#GdddeZ$ddZ%edddGdddeZ&edddGdddeZ'edddGdddeZ(Gddde Z)e   d*ddZ*e  d+d d!Z+  d,d"d#Z,d$d%Z-d&d'Z.Gd(d)d)e/Z0dS)-)ConcreteTemplate)typestypingfuncdescconfigcompilersigutils)sanitize_compile_result_entries CompilerBaseDefaultPassBuilderFlagsOption CompileResult)global_compiler_lock) LoweringPass AnalysisPass PassManager register_pass)NumbaInvalidConfigWarning TypingError)IRLegalizationNativeLowering AnnotateTypes)warn)get_current_devicecCs|durdSt|ts J|SN) isinstancedict)xrf/home/ncw/WWW/www-new/content/articles/pi-bbp/venv/lib/python3.10/site-packages/numba/cuda/compiler.py_nvvm_options_typesr!c@s(eZdZeedddZeedddZdS) CUDAFlagsNz NVVM options)typedefaultdoczCompute Capability)__name__ __module__ __qualname__r r! nvvm_optionstuplecompute_capabilityrrrr r"s r"c@seZdZeddZdS)CUDACompileResultcCst|Sr)idselfrrr entry_point7szCUDACompileResult.entry_pointN)r&r'r(propertyr0rrrr r,6sr,cKst|}tdi|S)Nr)r r,)entriesrrr cuda_compile_result<sr3TF) mutates_CFG analysis_onlyc@ eZdZdZddZddZdS) CUDABackend cuda_backendcCt|dSrr__init__r.rrr r;FzCUDABackend.__init__c CsL|d}tj|jg|jR}t|j|j|jj|j |j |j ||j d|_ dS)zH Back-end: Packages lowering output in a compile result cr)typing_contexttarget_context typing_errortype_annotationlibrary call_helper signaturefndescT)rrD return_typeargsr3 typingctx targetctxstatus fail_reasonrArBrCrEr=)r/stateloweredrDrrr run_passIs zCUDABackend.run_passNr&r'r(_namer;rNrrrr r7A r7c@s$eZdZdZdZddZddZdS) 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_librarycCr9rr:r.rrr r;gr<zCreateLibrary.__init__cCs8|j}|jj}|jj}|j||d|_|jdS)N)r)T) rIcodegenfunc_id func_qualnameflagsr)rSrBenable_object_caching)r/rLrTnamer)rrr rNjs  zCreateLibrary.run_passN)r&r'r(__doc__rPr;rNrrrr rR]s  rRc@r6)CUDALegalizationcuda_legalizationcCr9r)rr;r.rrr r;zr<zCUDALegalization.__init__csXddlm}|jr dS|j}fdd|D]\}t|tjr)|jqdS)Nr)NVVMFcsVt|tjtjfrd}t|t|tjr'|jD] }|djqdSdS)Nz is a char sequence type. This type is not supported with CUDA toolkit versions < 11.2. To use this type, you need to update your CUDA toolkit - try 'conda install cudatoolkit=11' if you are using conda to manage your environment.) rrUnicodeCharSeqCharSeqrRecordfieldsitemsr#)dtypemsgsubdtype check_dtypekrr rhs  z.CUDALegalization.run_pass..check_dtype) numba.cuda.cudadrv.nvvmr] is_nvvm70typemaprcrrArrayrd)r/rLr]typmapvrrgr rN}s   zCUDALegalization.run_passNrOrrrr r[urQr[c@seZdZddZddZdS) CUDACompilercCstt}td}||j}|j|j||j}|j|j|td| |j}|j|j| |gS)NcudazCUDA legalization) r rdefine_untyped_pipelinerLpassesextenddefine_typed_pipelineadd_passr[define_cuda_lowering_pipelinefinalize)r/dpbpmuntyped_passes typed_passeslowering_passesrrr define_pipeliness    zCUDACompiler.define_pipelinescCsPtd}|td|td|td|td|td||S)N cuda_loweringz$ensure IR is legal prior to loweringzannotate typeszcreate libraryznative loweringz cuda backend)rrvrrrRrr7rx)r/rLrzrrr rws    z*CUDACompiler.define_cuda_lowering_pipelineN)r&r'r(r~rwrrrr rps rpNc  Cs|durtdddlm} | j} | j} t} d| _d| _d| _|s$|r'd| _ |r,d| _ |r2d| _ nd| _ |r:d| _ |r?d| _ |rD|| _|| _ddlm} | d tj| | |||| itd }Wdn1siwY|j}||S) Nz#Compute Capability must be suppliedr^ cuda_targetTpythonnumpyr)target_overriderq)rHrIfuncrGrFrWlocalspipeline_class) ValueError descriptorrr>r?r" no_compileno_cpython_wrapperno_cfunc_wrapper debuginfodbg_directives_only error_model forceinlinefastmathr)r+numba.core.target_extensionrr compile_extrarprBrx)pyfuncrFrGdebuglineinfoinlinerr)ccrrHrIrWrcresrBrrr compile_cudasL    rc Cs|r |r d}tt|||rdndd} t|\} } |p tj}t|| | |||| |d} | jj} | r>|s>| t j kr>t d|rD| j }n| j }|j}|j}|j}|| j | j||| ||\}}|j|d}|| fS)aCompile 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)ropt)rrrr)rz'CUDA kernel must have void return type.)r)rrrnormalize_signaturerCUDA_DEFAULT_PTX_CCrrDrFrvoid TypeErrorrBr?__code__ co_filenameco_firstlinenoprepare_cuda_kernelrE get_asm_str)rsigrrdevicerrrrer)rGrFrrestylibtgtcodefilenamelinenumkernelptxrrr compile_ptxs6     rc Cs tj}t|||||||ddS)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)rrrrrr)rr+r)rrrrrrrrrrr compile_ptx_for_current_device9s rcCst|||jSr) declare_device_function_templatekeyrYrestypeargtypesrrr declare_device_functionCr<rcsvddlm}|j}|j}tj|g|Rt|Gfdddt}tj |||d}| || ||S)Nr^rcseZdZZgZdS)zBdeclare_device_function_template..device_function_templateN)r&r'r(rcasesrextfnrrr device_function_templateNs rr) rrr>r?rrDExternFunctionrrExternalFunctionDescriptorinsert_user_function)rYrrrrHrIrrErrr rGs    rc@seZdZddZdS)rcCs||_||_dSr)rYr)r/rYrrrr r;[s zExternFunction.__init__N)r&r'r(r;rrrr rZs r)FFFFNN)FFFFNT)FFFFT)1numba.core.typing.templatesr numba.corerrrrrrnumba.core.compilerr r r r r rnumba.core.compiler_lockrnumba.core.compiler_machineryrrrrnumba.core.errorsrrnumba.core.typed_passesrrrwarningsrnumba.cuda.apirr!r"r,r3r7rRr[rprrrrrobjectrrrrr sF         "" : C