sh_$SSKJr SSKJrJr SSKJr SSKJrJ r SSK J r SSK J r SSKJr SSKJr S r\S 5r\S 5r\S 5r\ "\R,"\5S SS9S5r\S5rSr\S5r\S5r\S5rg))ir)cudatypes)cgutils)RequireLiteralValueNumbaValueError) signature)overload_attribute) nvvmutils) intrinsiccURnUS:Xa[RnO7US;a&[R"[RU5nO [ S5e[ U[R 5$)N)zargument can only be 1, 2, 3) literal_valuerint64UniTuplerr int32)ndimvalrestypes i/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/intrinsics.py_type_grid_functionrsU   C ax++ ..c2<== Wekk **cx[U[R5(d [U5e[ U5nSnX#4$)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. cURnU[R:Xa[R"USS9$[ U[R 5(a4[R"XRS9n[R"X5$g)Nr)dim) 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 <rc^[U[R5(d [U5e[ U5nSmU4SjnX#4$)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[R"S5n[R"USU35n[R"USU35nUR UR X25UR XB55$)N@zntid.znctaid.)rIntTyper call_sregmulsext)r$ri64ntidnctaids r_nthreads_for_dim#gridsize.._nthreads_for_dimRsbjjn""7eC5M:$$Wuo>{{7<<2GLL4MNNrc`>URnT"US5nU[R:XaU$[U[R5(acT"US5nUR S:Xa[ R"XU45$UR S:Xa!T"US5n[ R"XXg45$gg)Nxyrrz)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@[[R5nSnX4$)Nc0[R"US5$)Nwarpsize)r r2)r#r$r%r&s rr(_warpsize..codegenns""7J77r)r rrr,r%r(s r _warpsizerHjs EKK C8 <rrEr)targetc SnU$)zS The size of a warp. All architectures implemented to date have a warp size of 32. c[5$)N)rH)mods rgetcuda_warpsize..getzs {r)rLrMs r cuda_warpsizerPts  Jrc@[[R5nSnX4$)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. cSnURn[R"[R"5S5n[R "XVU5nUR US5 UR5$)Nzllvm.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) <rc^[U[R5(dg[[R[R5nU4SjnX44$)Nc>[R"[R"S5[R"S545n[R"UR UT5nUR XS5$)N )rrTr1rrVrSrW)r#r$r%r&r[r\rYs rr('_syncthreads_predicate..codegensMrzz"~ 2/@A--gnndEJ||D''r)r rIntegerr i4)r, predicaterYr%r(s ` r_syncthreads_predicatergs: i / / EHHehh 'C( <rcSn[XU5$)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 !) >>rcSn[XU5$)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 !) >>rcSn[XU5$)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