U ,dw@sddlZddlZddlZddlZddlZddlmZmZm Z m Z m Z m Z ddl mZmZddlmZddlmZddlmZddlmZmZddlmZdd lmZdd lmZmZdd l m!Z!dd l"m#Z#dd l$m%Z%ddl&m'Z'ddl(m)Z)m*Z*ddl+m Z,ddl-m.Z.ddl-m/Z/ddl0m1Z1Gdddej2Z3Gddde4Z5GdddZ6GdddeZ7GdddeZ8Gdddeej2Z9dS) N)config serializesigutilstypestypingutils)Cache CacheImpl)global_compiler_lock) Dispatcher)NumbaPerformanceWarning)Purposetypeofget_current_device)wrap_arg) compile_cuda CUDACompiler)driver) get_context) get_cudalib) cuda_target)missing_launch_config_msgnormalize_kernel_dimensions)rcuda) _dispatcher)warnc seZdZdZed*fdd ZeddZed d Zd d Z ed dZ eddZ e fddZ ddZddZeddZeddZddZddZdd Zd+d!d"Zd,d$d%Zd-d&d'Zd(d)ZZS)._Kernelz CUDA Kernel specialized for a given set of argument types. When called, this object launches the kernel on the device. NFTc  sL| r tdtd|_d|_||_||_||_||_|p@g|_ |j|j|| rVdndd} t |jt j |j|j|j||| d} | j }|jj}|j}|j}|| j| j|| ||| \}}|sg}d|k|_|jr|tdd d |D]}||q|j|_| j|_| j|_||_| j|_||_ | j|_| j |_ g|_!g|_"g|_#dS) Nz,Cannot compile a device function as a kernelFr)debuglineinfofastmathopt)r r!inliner" nvvm_optionsZcudaCGGetIntrinsicHandleZ cudadevrtT)Zstatic)$ RuntimeErrorsuper__init__Z objectmode entry_pointpy_funcargtypesr r! extensionsrrvoidtarget_context__code__ co_filenameco_firstlinenoZprepare_cuda_kernellibraryfndesc get_asm_str cooperativeappendrZadd_linking_filename entry_name signaturetype_annotation_type_annotation _codelibrary call_helper environment_referenced_environmentsZliftedZ reload_init)selfr*r+linkr r!r$r"r,Z max_registersr#devicer%cresZtgt_ctxcodefilenamelinenumlibkernelfilepath __class__9/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/dispatcher.pyr(%sf     z_Kernel.__init__cCs|jSN)r<r@rLrLrMr2xsz_Kernel.librarycCs|jSrN)r;rOrLrLrMr:|sz_Kernel.type_annotationcCs|jSrN)r?rOrLrLrM_find_referenced_environmentssz%_Kernel._find_referenced_environmentscCs |jSrN)r.codegenrOrLrLrMrQsz_Kernel.codegencCs t|jjSrN)tupler9argsrOrLrLrMargument_typessz_Kernel.argument_typesc sX||} t|| d| _|| _|| _|| _d| _|| _|| _ || _ || _ || _ | S)& Rebuild an instance. N) __new__r'r(r)r5r8r9r;r<r r!r=r,) clsr5r7r9 codelibraryr r!r=r,instancerJrLrM_rebuilds z_Kernel._rebuildc Cs(t|j|j|j|j|j|j|j|jdS)a Reduce the instance for serialization. Compiled definitions are serialized in PTX form. Type annotation are discarded. Thread, block and shared memory configuration are serialized. Stream information is discarded. )r5r7r9rXr r!r=r,) dictr5r8r9r<r r!r=r,rOrLrLrM_reduce_statess z_Kernel._reduce_statescCs|jdS)z7 Force binding to current CUDA context N)r< get_cufuncrOrLrLrMbindsz _Kernel.bindcCstS)z, Get current active context rrOrLrLrMrBsz_Kernel.devicecCs|jjjS)zN The number of registers used by each thread for this kernel. )r<r]attrsregsrOrLrLrMregs_per_threadsz_Kernel.regs_per_threadcCs |jS)z6 Returns the LLVM IR for this kernel. )r< get_llvm_strrOrLrLrM inspect_llvmsz_Kernel.inspect_llvmcCs|jj|dS)z7 Returns the PTX code for this kernel. cc)r<r4)r@rerLrLrM inspect_asmsz_Kernel.inspect_asmcCs |jS)zp Returns the SASS code for this kernel. Requires nvdisasm to be available on the PATH. )r<Zget_sassrOrLrLrM inspect_sasssz_Kernel.inspect_sasscCsb|jdkrtd|dkr tj}td|j|jf|dtd|dt|j|dtd|ddS) Produce a dump of the Python source of this function annotated with the corresponding Numba IR and type information. The dump is written to *file*, or *sys.stdout* if *file* is *None*. Nz Type annotation is not availablez%s %sfilezP--------------------------------------------------------------------------------zP================================================================================)r; ValueErrorsysstdoutprintr8rT)r@rjrLrLrM inspect_typess  z_Kernel.inspect_typesrcCsHt}|j}t|tr*tdd|}||||}|jj }||S)a Calculates the maximum number of blocks that can be launched for this kernel in a cooperative grid in the current context, for the given block and dynamic shared memory sizes. :param blockdim: Block dimensions, either as a scalar for a 1D block, or a tuple for 2D or 3D blocks. :param dynsmemsize: Dynamic shared memory size in bytes. :return: The maximum number of blocks in the grid. cSs||SrNrL)xyrLrLrMz5_Kernel.max_cooperative_grid_blocks..) rr<r] isinstancerR functoolsreduceZ$get_active_blocks_per_multiprocessorrBZMULTIPROCESSOR_COUNT)r@blockdimZ dynsmemsizectxcufuncZ active_per_smZsm_countrLrLrMmax_cooperative_grid_blockss   z#_Kernel.max_cooperative_grid_blockscs|j|jrTjd}j|\}}|ttjks>t t} |j d|dg} g} t |j |D]\} } | | | || | qhtjrtjd}nd}|r|jp|}tjjf||||| fd|ji|jrtt| ||| jdkrfddfddd D}fd dd D}| j}|j|\}}}|dkrNd }n$|\}}}tj|}d |||f}d |||f}|rd||dff|dd}n|f}||| D] }|qdS)NZ __errcode__r)streamr5cs<jdj|f\}}t}tt||||jS)Nz%s__%s__) moduleget_global_symbolr7ctypesc_intrdevice_to_host addressofvalue)r7Zmemszval)ryrLrM load_symbol#s z#_Kernel.launch..load_symbolcsg|]}d|qS)tidrL.0irrLrM +sz"_Kernel.launch..Zzyxcsg|]}d|qS)ctaidrLrrrLrMr,sz"In function %r, file %s, line %s, z%stid=%s ctaid=%sz%s: %s)r<r]r r7r|r}r~sizeofrAssertionErrormemsetziprT _prepare_argsrUSE_NV_BINDINGZbindingZCUstreamhandleZ launch_kernelr5rrrr=Z get_exceptionospathabspath)r@rSgriddimrwr{ sharedmemexcnameZexcmemZexcszZexcvalretr kernelargstvZ zero_streamZ stream_handlerrrDZexcclsZexc_argslocZlocinfosymrIlinenoprefixwbrL)ryrrMlaunchsb         z_Kernel.launchcCsht|jD]}|j||||d\}}q t|tjrt|||}tj }t d} t d} ||j } ||j j } t|} tjrt| } t | }|| || || || ||t|jD]}|||j|qt|jD]}|||j|qnPt|tjrBttd||}||n"|tjkrttt|tj}||n|tjkrt|}||n|tj krt!|}||n|tj"krt#t|}||n|tj$kr|t!|j%|t!|j&nL|tj'krL|t|j%|t|j&nt|tj(tj)frz|t*|tj+nt|tj,rt|||}|j-} tjrt t| } || nt|tj.rt/|t/|kst0t1||D]\}}|2|||||qnVt|tj3rZz|2|j |j4|||Wn t5k rVt5||YnXn t5||dS)zF Convert arguments to ctypes and append to kernelargs )r{rrzc_%sN)6reversedr,Z prepare_argsrtrArrayrZ to_devicer~ c_ssize_tc_void_psizeZdtypeitemsizerZdevice_pointerrintr6rangendimshapestridesZIntegergetattrZfloat16c_uint16npviewZuint16Zfloat64c_doubleZfloat32c_floatbooleanc_uint8Z complex64realimagZ complex128Z NPDatetimeZ NPTimedeltac_int64Zint64ZRecordZdevice_ctypes_pointerZ BaseTuplelenrrrZ EnumMemberrNotImplementedError)r@tyrr{rr extensionZdevaryZc_intpZmeminfoparentZnitemsrptrdataZaxcvalZdevrecrrrLrLrMrEs                     z_Kernel._prepare_args) NFFFFNNTF)N)r)rr)__name__ __module__ __qualname____doc__r r(propertyr2r:rPrQrT classmethodrZr\r^rBrarcrfrgrorzrr __classcell__rLrLrJrMrsFR          Hrc@s$eZdZddZddZddZdS)ForAllcCs6|dkrtd|||_||_||_||_||_dS)Nrz0Can't create ForAll with negative task count: %s)rk dispatcherntasksthread_per_blockr{r)r@rrtpbr{rrLrLrMr(szForAll.__init__cGs^|jdkrdS|jjr|j}n |jj|}||}|j|d|}||||j|jf|S)Nrr)rr specialized specialize_compute_thread_per_blockr{r)r@rSrrwrrLrLrM__call__s    zForAll.__call__cCsZ|j}|dkr|St}tt|j}t|jd|j dd}|j f|\}}|SdS)Nri)funcZb2d_funcZmemsizeZblocksizelimit) rrnextiter overloadsvaluesr[r<r]rZget_max_potential_block_size)r@rrrxrHkwargs_rLrLrMrsz ForAll._compute_thread_per_blockN)rrrr(rrrLrLrLrMrs rc@seZdZddZddZdS)_LaunchConfigurationc Csd||_||_||_||_||_tjr`d}|d|d|d}||kr`d|d}tt|dS)Nrrz Grid size zB will likely result in GPU under-utilization due to low occupancy.) rrrwr{rrZCUDA_LOW_OCCUPANCY_WARNINGSrr ) r@rrrwr{rZ min_grid_sizeZ grid_sizemsgrLrLrMr(s  z_LaunchConfiguration.__init__cGs|j||j|j|j|jSrN)rcallrrwr{rr@rSrLrLrMrsz_LaunchConfiguration.__call__N)rrrr(rrLrLrLrMrsrc@s$eZdZddZddZddZdS) CUDACacheImplcCs|SrN)r\)r@rHrLrLrMrvszCUDACacheImpl.reducecCs tjf|SrN)rrZ)r@r.payloadrLrLrMrebuildszCUDACacheImpl.rebuildcCsdS)NTrL)r@rCrLrLrMcheck_cachableszCUDACacheImpl.check_cachableN)rrrrvrrrLrLrLrMrsrc@seZdZdZeZdS) CUDACachezS Implements a cache that saves and loads CUDA kernels and compile results. N)rrrrrZ _impl_classrLrLrLrMrsrcseZdZdZdZeZeffdd Ze ddZ ddZ e j d d d9d d ZddZd:ddZe ddZddZddZddZddZddZe ddZd;d!d"Zd#d$Zd%d&Zd'd(Zd)d*Zdd/d0Zd?d1d2Z d3d4Z!e"d5d6Z#d7d8Z$Z%S)@CUDADispatchera CUDA Dispatcher object. When configured and called, the dispatcher will specialize itself for the given arguments (if no suitable specialized version already exists) & compute capability, and launch on the device associated with the current context. Dispatcher objects are not to be constructed by the user, but instead are created using the :func:`numba.cuda.jit` decorator. Fcs"tj|||dd|_i|_dS)N) targetoptionspipeline_classF)r'r( _specializedspecializations)r@r*rrrJrLrMr(s   zCUDADispatcher.__init__cCs t|SrN) cuda_typesrrOrLrLrM _numba_type_*szCUDADispatcher._numba_type_cCst|j|_dSrN)rr*_cacherOrLrLrMenable_caching.szCUDADispatcher.enable_cachingr)maxsizercCst||\}}t|||||SrN)rr)r@rrwr{rrLrLrM configure1szCUDADispatcher.configurecCst|dkrtd|j|S)N)rrz.must specify at least the griddim and blockdim)rrkrrrLrLrM __getitem__6s zCUDADispatcher.__getitem__cCst|||||dS)a3Returns a 1D-configured dispatcher for a given number of tasks. This assumes that: - the kernel maps the Global Thread ID ``cuda.grid(1)`` to tasks on a 1-1 basis. - the kernel checks that the Global Thread ID is upper-bounded by ``ntasks``, and does nothing if it is not. :param ntasks: The number of tasks. :param tpb: The size of a block. An appropriate value is chosen if this parameter is not supplied. :param stream: The stream on which the configured dispatcher will be launched. :param sharedmem: The number of bytes of dynamic shared memory required by the kernel. :return: A configured dispatcher, ready to launch on a set of arguments.)rr{r)r)r@rrr{rrLrLrMforall;szCUDADispatcher.forallcCs |jdS)aS A list of objects that must have a `prepare_args` function. When a specialized kernel is called, each argument will be passed through to the `prepare_args` (from the last object in this list to the first). The arguments to `prepare_args` are: - `ty` the numba type of the argument - `val` the argument value itself - `stream` the CUDA stream used for the current call to the kernel - `retr` a list of zero-arg functions that you may want to append post-call cleanup work to. The `prepare_args` function must return a tuple `(ty, val)`, which will be passed in turn to the next right-most `extension`. After all the extensions have been called, the resulting `(ty, val)` will be passed into Numba's default argument marshalling logic. r,)rgetrOrLrLrMr,QszCUDADispatcher.extensionscOs ttdSrN)rkr)r@rSrrLrLrMrfszCUDADispatcher.__call__cCsB|jrtt|j}ntjj|f|}||||||dS)zJ Compile if necessary and invoke this kernel with *args*. N) rrrrrrr Z _cuda_callr)r@rSrrwr{rrHrLrLrMrjszCUDADispatcher.callcs(|rtfdd|D}t|S)Ncsg|]}|qSrL) typeof_pyvalrarOrLrMrxsz4CUDADispatcher._compile_for_args..)rcompilerR)r@rSkwsr+rLrOrM_compile_for_argsusz CUDADispatcher._compile_for_argscCsNzt|tjWStk rHt|rBttj|ddtjYSYnXdS)NF)sync)rr argumentrkrZ is_cuda_arrayZ as_cuda_array)r@rrLrLrMr{s zCUDADispatcher.typeof_pyvalcstj}tfdd|D}jr,tdj||f}|rD|Sj}tj |d}| || d|_ |j||f<|S)zd Create a new instance of this dispatcher specialized for the given *args*. csg|]}j|qSrL)Z typingctxZresolve_argument_typerrOrLrMrsz-CUDADispatcher.specialize..zDispatcher already specialized)rT) rcompute_capabilityrRrr&rrrrr*rZdisable_compiler)r@rSrer+ZspecializationrrLrOrMrs$ zCUDADispatcher.specializecCs|jS)z> True if the Dispatcher has been specialized. )rrOrLrLrMrszCUDADispatcher.specializedNcCsH|dk r|j|jjS|jr0tt|jjSdd|jDSdS)a Returns the number of registers used by each thread in this kernel for the device in the current context. :param signature: The signature of the compiled kernel to get register usage for. This may be omitted for a specialized kernel. :return: The number of registers used by the compiled variant of the kernel for the given signature and current device. NcSsi|]\}}||jqSrL)rarsigoverloadrLrLrM sz6CUDADispatcher.get_regs_per_thread..)rrSrarrrritemsr@r9rLrLrMget_regs_per_threads z"CUDADispatcher.get_regs_per_threadcCsP|jr|t||jj}d|}tj|||jd}t |j}||||fS)z Get a typing.ConcreteTemplate for this dispatcher and the given *args* and *kws* types. This allows resolution of the return type. A (template, pysig, args, kws) tuple is returned. zCallTemplate({0}))keyZ signatures) _can_compilecompile_devicerRr*rformatrZmake_concrete_templateZnopython_signaturesrZ pysignature)r@rSr func_namer7Z call_templateZpysigrLrLrMget_call_templates  z CUDADispatcher.get_call_templatec Cs||jkr|j|jd}|jd}|jd}||jdrHdnd|d}t|jd|||||d }||j|<|j|j|j |j gW5QRXn |j|}|S) zCompile the device function for the given argument types. Each signature is compiled once by caching the compiled function inside this object. Returns the `CompileResult`. r r$r"r#rr)r r#r"N)r r$r"r%) rZ_compiling_counterrrrr*r.Zinsert_user_functionr)r3r2)r@rSr r$r"r%rCrLrLrMrs,        zCUDADispatcher.compile_devicecCs,dd|D}|j||dd||j|<dS)NcSsg|] }|jqSrL)_coderrLrLrMrsz/CUDADispatcher.add_overload..Tr)Z_insertr)r@rHr+Zc_sigrLrLrM add_overloadszCUDADispatcher.add_overloadcCst|\}}|dks$|tjks$t|jr.cSsi|]\}}||qSrL)rcrrLrLrMr7s)rrrr2rbrcrr@r9rBrLrLrMrc#s zCUDADispatcher.inspect_llvmcs|tj|jd}|dk rD|r2|j|jS|j|Sn4|r`fdd|jDSfdd|jDSdS)a+ Return this kernel's PTX assembly code for for the device in the current context. :param signature: A tuple of argument types. :return: The PTX code for the given signature, or a dict of PTX codes for all previously-encountered signatures. rBNcsi|]\}}||jqSrL)r2r4rrdrLrMrLsz.CUDADispatcher.inspect_asm..csi|]\}}||qSrL)rfrrdrLrMrOs) rrrrrr2r4rfrrrLrdrMrf:s    zCUDADispatcher.inspect_asmcCsB|jdrtd|dk r*|j|Sdd|jDSdS)a Return this kernel's SASS assembly code for for the device in the current context. :param signature: A tuple of argument types. :return: The SASS code for the given signature, or a dict of SASS codes for all previously-encountered signatures. SASS for the device in the current context is returned. Requires nvdisasm to be available on the PATH. rBz(Cannot inspect SASS of a device functionNcSsi|]\}}||qSrL)rg)rrdefnrLrLrMresz/CUDADispatcher.inspect_sass..)rrr&rrgrrrLrLrMrgRs zCUDADispatcher.inspect_sasscCs2|dkrtj}|jD]\}}|j|dqdS)rhNri)rlrmrrro)r@rjrrrLrLrMrohszCUDADispatcher.inspect_typescCs|jD] }|q dSrN)rrr^)r@rrLrLrMr^tszCUDADispatcher.bindcCs|||}|S)rUrL)rWr*rrYrLrLrMrZxs zCUDADispatcher._rebuildcCst|j|jdS)zd Reduce the instance for serialization. Compiled definitions are discarded. )r*r)r[r*rrOrLrLrMr\szCUDADispatcher._reduce_states)rr)rrr)N)N)N)N)N)&rrrrZ _fold_argsrZ targetdescrrr(rrrru lru_cacherrrr,rrrrrrrr rr rrcrfrgror^rrZr\rrLrLrJrMrsB         $$     r):Znumpyrrrlr~ruZ numba.corerrrrrrZnumba.core.cachingrr Znumba.core.compiler_lockr Znumba.core.dispatcherr Znumba.core.errorsr Znumba.core.typing.typeofr rZnumba.cuda.apirZnumba.cuda.argsrZnumba.cuda.compilerrrZnumba.cuda.cudadrvrZnumba.cuda.cudadrv.devicesrZnumba.cuda.cudadrv.libsrZnumba.cuda.descriptorrZnumba.cuda.errorsrrZ numba.cudarZnumbarrwarningsrZ ReduceMixinrobjectrrrrrrLrLrLrMs>              .