i_*ddlmZddlmZmZddlmZddlmZm Z ddl m Z ddl m Z ddlmZddlmZd Zed Zed Zed Ze ej,ed ddZedZdZedZedZedZy))ir)cudatypes)cgutils)RequireLiteralValueNumbaValueError) signature)overload_attribute) nvvmutils) intrinsicc|j}|dk(rtj}n4|dvr%tjtj|}n t dt |tj S)N)zargument can only be 1, 2, 3) literal_valuerint64UniTuplerr int32)ndimvalrestypes j/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/intrinsics.py_type_grid_functionrsU   C ax++ ..c2<== Wekk **cpt|tjs t|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}|tjk(rtj|dSt |tj r7tj||j}tj||Sy)Nr)dim) return_typerrr get_global_id isinstancercountr pack_array)contextbuildersigargsridss rcodegenzgrid..codegen1sf// ekk !**7: :  0))'w}}EC%%gs3 31rr rIntegerLiteralrr) typingctxrr%r(s rgridr,s:" dE00 1!$'' d #C4 <rc|t|tjs t|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|}|j |j |||j ||S)N@zntid.znctaid.)rIntTyper call_sregmulsext)r$ri64ntidnctaids r_nthreads_for_dimz#gridsize.._nthreads_for_dimRsfjjn""7eC5M:$$Wuo>{{7<<c2GLL4MNNrcR|j}|d}|tjk(r|St|tjrb|d}|j dk(rt j|||fS|j dk(r"|d}t j||||fSyy)Nxyrrz)rrrr rr!rr") r#r$r%r&rnxnynzr7s rr(zgridsize..codegenXs// w , ekk !I  0"7C0B}}!))'B8<<!#&w4))'BB<@@$ 1rr))r+rr%r(r7s @rgridsizer?<sB" dE00 1!$'' d #CO A <rcBttj}d}||fS)Nc.tj|dS)Nwarpsize)r r1)r#r$r%r&s rr(z_warpsize..codegenns""7J77r)r rrr+r%r(s r _warpsizerDjs! EKK C8 <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 |||}|j |d|jS)Nzllvm.nvvm.barrier0rJ)moduler FunctionTypeVoidTyperget_or_insert_functioncallget_dummy_value)r#r$r%r&fnamelmodfntysyncs rr(zsyncthreads..codegensW$~~r{{}b1--dD%@ T2&&((r)r rnonerCs r syncthreadsrYs# EJJ C) <rct|tjsyttjtj}fd}||fS)Nctjtjdtjdf}tj|j |}|j ||S)N )rrOr0rrQrNrR)r#r$r%r&rVrWrTs rr(z'_syncthreads_predicate..codegensOrzz"~ 2/@A--gnndEJ||D$''r)r rIntegerr i4)r+ predicaterTr%r(s ` r_syncthreads_predicater`s9 i / EHHehh '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`rdrfrhrJrrrrsB'3 * +  @ * *Z  LELL& 6BC  (  ? ? ? ? ? ?r