1@j%^ddlmZddlmZmZddlmZddlmZm Z ddl m Z ddl m Z ddlmZddlmZd Zed Zed Zed Ze ejed ddZedZdZedZedZedZdS))ir)cudatypes)cgutils)RequireLiteralValueNumbaValueError) signature)overload_attribute) nvvmutils) intrinsicc|j}|dkr tj}n3|dvr tjtj|}nt dt |tjS)N)zargument can only be 1, 2, 3) literal_valuerint64UniTuplerr int32)ndimvalrestypes FC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/intrinsics.py_type_grid_functionrs[  C axx+ .c22<=== Wek * **ct|tjst|t |}d}||fS)agrid(ndim) Return the absolute position of the current thread in the entire grid of blocks. *ndim* should correspond to the number of dimensions declared when instantiating the kernel. If *ndim* is 1, a single integer is returned. If *ndim* is 2 or 3, a tuple of the given number of integers is returned. Computation of the first integer is as follows:: cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x and is similar for the other two indices, but using the ``y`` and ``z`` attributes. c|j}|tjkrtj|dSt |tjr0tj||j}tj ||SdS)Nr)dim) return_typerrr get_global_id isinstancercountr pack_array)contextbuildersigargsridss rcodegenzgrid..codegen1st/ ek ! !*7::: :  0 0 4)'w}EEEC%gs33 3 4 4rr rIntegerLiteralrr) typingctxrr%r(s rgridr,sO" dE0 1 1(!$''' d # #C444 <rct|tjst|t |}dfd}||fS)agridsize(ndim) Return the absolute size (or shape) in threads of the entire grid of blocks. *ndim* should correspond to the number of dimensions declared when instantiating the kernel. If *ndim* is 1, a single integer is returned. If *ndim* is 2 or 3, a tuple of the given number of integers is returned. Computation of the first integer is as follows:: cuda.blockDim.x * cuda.gridDim.x and is similar for the other two indices, but using the ``y`` and ``z`` attributes. ctjd}tj|d|}tj|d|}|||||||S)N@zntid.znctaid.)rIntTyper call_sregmulsext)r$ri64ntidnctaids r_nthreads_for_dimz#gridsize.._nthreads_for_dimRsojnn"7MCMM::$Wooo>>{{7<<c22GLL4M4MNNNrcD|j}|d}|tjkr|St|tjr]|d}|jdkrt j|||fS|jdkr&|d}t j||||fSdSdS)Nxyrrz)rrrr rr!rr") r#r$r%r&rnxnynzr7s rr(zgridsize..codegenXs/  w , , ek ! !I  0 0 A""7C00B}!!)'B8<<<!##&&w44)'BB<@@@ A A $#rr))r+rr%r(r7s @rgridsizer?<sq" dE0 1 1(!$''' d # #COOO A A A A A <rcBttj}d}||fS)Nc,tj|dS)Nwarpsize)r r1)r#r$r%r&s rr(z_warpsize..codegenns"7J777r)r rrr+r%r(s r _warpsizerDjs, EK C888 <rrBr)targetc d}|S)z_ The size of a warp. All architectures implemented to date have a warp size of 32. ctS)N)rD)mods rgetzcuda_warpsize..getzs {{r)rHrIs r cuda_warpsizerKts  JrcBttj}d}||fS)a Synchronize all threads in the same thread block. This function implements the same pattern as barriers in traditional multi-threaded programming: this function waits until all threads in the block call it, at which point it returns control to all its callers. cd}|j}tjtjd}t j|||}||d|S)Nzllvm.nvvm.barrier0rJ)moduler FunctionTypeVoidTyperget_or_insert_functioncallget_dummy_value)r#r$r%r&fnamelmodfntysyncs rr(zsyncthreads..codegens_$~r{}}b11-dD%@@ T2&&(((r)r rnonerCs r syncthreadsrYs. EJ  C))) <rct|tjsdSttjtj}fd}||fS)Nctjtjdtjdf}tj|j|}|||S)N )rrOr0rrQrNrR)r#r$r%r&rVrWrTs rr(z'_syncthreads_predicate..codegensQrz"~~ 2/@AA-gndEJJ||D$'''r)r rIntegerr i4)r+ predicaterTr%r(s ` r_syncthreads_predicater`sT i / /t EHeh ' 'C((((( <rc(d}t|||S)z syncthreads_count(predicate) An extension to numba.cuda.syncthreads where the return value is a count of the threads where predicate is true. zllvm.nvvm.barrier0.popcr`r+r_rTs rsyncthreads_countrds &E !)Y > >>rc(d}t|||S)z syncthreads_and(predicate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for all threads or 0 otherwise. zllvm.nvvm.barrier0.andrbrcs rsyncthreads_andrfs %E !)Y > >>rc(d}t|||S)z syncthreads_or(predicate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for any thread or 0 otherwise. zllvm.nvvm.barrier0.orrbrcs rsyncthreads_orrhs $E !)Y > >>rN)llvmliternumbarr numba.corernumba.core.errorsrrnumba.core.typingr numba.core.extendingr numba.cudar numba.cuda.extendingr rr,r?rDModulerKrYr`rdrfrhrJrrrrsBBBBBBBB''''''333333 ****** + + +  @ ** *Z  LEL&& 6BBBCB  (    ?? ? ?? ? ?? ???r