+ :i_,^RIHt^RIHtHt^RIHt^RIHtH t ^RI H t ^RI H t ^RIHt^RIHtR t]R 4t]R 4t]R 4t] !]P,!]4R RR7R4t]R4tRt]R4t]R4t]R4tR#))ir)cudatypes)cgutils)RequireLiteralValueNumbaValueError) signature)overload_attribute) nvvmutils) intrinsiccVPpV^8Xd\PpM8VR9d'\P!\PV4pM \ R4h\ V\P 4#)zargument can only be 1, 2, 3)) literal_valuerint64UniTuplerr int32)ndimvalrestypes& S/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/intrinsics.py_type_grid_functionrsU   C ax++ ..c2<== Wekk **cz\V\P4'g \V4h\ V4pRpW#3#)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&VPpV\P8Xd\P!V^R7#\ V\P 4'd8\P!WPR7p\P!W4#R#)r)dimN) return_typerrr get_global_id isinstancercountr pack_array)contextbuildersigargsridss&&&& rcodegengrid..codegen1se// ekk !**7: :  0 0))'}}EC%%g3 31rr rIntegerLiteralrr) typingctxrr%r(s&& rgridr-s;" dE00 1 1!$'' d #C4 <rca\V\P4'g \V4h\ V4pRoV3RlpW#3#)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. c\P!^@4p\P!VRV 24p\P!VRV 24pVP VP W24VP WB44#)@zntid.znctaid.)rIntTyper call_sregmulsext)r$ri64ntidnctaids&& r_nthreads_for_dim#gridsize.._nthreads_for_dimRsbjjn""7eC5M:$$Wuo>{{7<<2GLL4MNNrcl<VPpS!VR4pV\P8XdV#\V\P4'dgS!VR4pVP ^8Xd\ P!WV34#VP ^8Xd"S!VR4p\ P!WWg34#R#R#)xyzN)rrrr rr!rr") r#r$r%r&rnxnynzr8s &&&& rr(gridsize..codegenXs// w , ekk !I  0 0"7C0B}}!))'8<<!#&w4))'<@@$ 1rr*)r,rr%r(r8s&& @rgridsizerB<sC" dE00 1 1!$'' d #CO A <rc@\\P4pRpW3#)c0\P!VR4#)warpsize)r r2)r#r$r%r&s&&&&rr(_warpsize..codegenns""7J77r)r rrr,r%r(s& r _warpsizerHjs EKK C8 <rrEr)targetc RpV#)zS The size of a warp. All architectures implemented to date have a warp size of 32. c\4#)N)rH)mods&rgetcuda_warpsize..getzs {r)rLrMs& r cuda_warpsizerPts  Jrc@\\P4pRpW3#)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. cRpVPp\P!\P!4R4p\P !WVV4pVP VR4VP4#)zllvm.nvvm.barrier0rO)moduler FunctionTypeVoidTyperget_or_insert_functioncallget_dummy_value)r#r$r%r&fnamelmodfntysyncs&&&& rr(syncthreads..codegensU$~~r{{}b1--d%@ T2&&((r)r rnonerGs& r syncthreadsr_s! EJJ C) <rca\V\P4'gR#\\P\P4pV3RlpW43#)Nc<\P!\P!^ 4\P!^ 434p\P!VP VS4pVP WS4#) )rrTr1rrVrSrW)r#r$r%r&r[r\rYs&&&& rr('_syncthreads_predicate..codegensMrzz"~ 2/@A--gnndEJ||D''r)r rIntegerr i4)r, predicaterYr%r(s&&f r_syncthreads_predicatergs: i / / EHHehh 'C( <rcRp\WV4#)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.popcrgr,rfrYs&& rsyncthreads_countrks &E !) >>rcRp\WV4#)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.andrirjs&& rsyncthreads_andrms %E !) >>rcRp\WV4#)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.orrirjs&& rsyncthreads_orros $E !) >>rN)llvmliternumbarr numba.corernumba.core.errorsrrnumba.core.typingr numba.core.extendingr numba.cudar numba.cuda.extendingr rr-rBrHModulerPr_rgrkrmrorOrrrysB'3 * +  @ * *Z  ELL& 6BC  (  ? ? ? ? ? ?r