o i_@sddlmZddlmZmZddlmZddlmZm Z ddl m Z ddl m Z ddlmZddlmZd d Zed d Zed dZeddZe eedddddZeddZddZeddZeddZeddZd S)!)ir)cudatypes)cgutils)RequireLiteralValueNumbaValueError) signature)overload_attribute) nvvmutils) intrinsiccCsB|j}|dkr tj}n|dvrttj|}ntdt|tjS)N)zargument can only be 1, 2, 3) literal_valuerint64UniTuplerrint32)ndimvalrestyperR/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/numba/cuda/intrinsics.py_type_grid_functions rcCs,t|tjs t|t|}dd}||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. cSsJ|j}|tjkrtj|ddSt|tjr#tj||jd}t ||SdS)Nr )dim) return_typerrr get_global_id isinstancercountr pack_array)contextbuildersigargsridsrrrcodegen1s   zgrid..codegenrrIntegerLiteralrr typingctxrr!r$rrrgrids r)cs8t|tjs t|t|}ddfdd}||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. cSsJtd}t|d|}t|d|}|||||||S)N@zntid.znctaid.)rIntTyper call_sregmulsext)r ri64ntidnctaidrrr_nthreads_for_dimRs z#gridsize.._nthreads_for_dimcs||j}|d}|tjkr|St|tjr:|d}|jdkr't|||fS|jdkr<|d}t||||fSdSdS)Nxyr rz)rrrrrrrr)rr r!r"rnxnynzr2rrr$Xs       zgridsize..codegenr%r'rr9rgridsize<s  r:cCttj}dd}||fS)NcSs t|dS)Nwarpsize)r r,)rr r!r"rrrr$ns z_warpsize..codegen)rrrr(r!r$rrr _warpsizejs r>r<r)targetcCs dd}|S)z_ The size of a warp. All architectures implemented to date have a warp size of 32. cSstS)N)r>)modrrrgetzszcuda_warpsize..getr)r@rArrr cuda_warpsizetsrBcCr;)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. cSs<d}|j}ttd}t|||}||d|S)Nzllvm.nvvm.barrier0r)moduler FunctionTypeVoidTyperget_or_insert_functioncallget_dummy_value)rr r!r"fnamelmodfntysyncrrrr$s  zsyncthreads..codegen)rrnoner=rrr syncthreadss rNcs2t|tjsdSttjtj}fdd}||fS)Ncs6ttdtdf}t|j|}|||S)N )rrDr+rrFrCrG)rr r!r"rKrLrIrrr$s z'_syncthreads_predicate..codegen)rrIntegerri4)r( predicaterIr!r$rrPr_syncthreads_predicates  rTcCd}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.popcrTr(rSrIrrrsyncthreads_count rXcCrU)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.andrVrWrrrsyncthreads_andrYrZcCrU)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.orrVrWrrrsyncthreads_orrYr[N)llvmliternumbarr numba.corernumba.core.errorsrrnumba.core.typingrnumba.core.extendingr numba.cudar numba.cuda.extendingr rr)r:r>ModulerBrNrTrXrZr[rrrrs4        -