shWSrSSKrSSKJr SSKrSSKrSSKJrJ r "SS\ 5r Sr "SS \ 5r "S S \ 5r"S S \ 5r"SS\ 5r"SS\ 5r"SS\ 5r"SS\ 5r"SS\ 5r"SS\ 5r"SS\ 5r"SS\ 5r"SS\ 5r"S S!\ 5r"S"S#\ 5r"S$S%\ 5r"S&S'\ 5r"S(S)\ 5r"S*S+\ 5r"S,S-\ 5r"S.S/\ 5r "S0S1\ 5r!"S2S3\ 5r""S4S5\ 5r#"S6S7\ 5r$"S8S9\ 5r%"S:S;\ 5r&"S<S=\ 5r'"S>S?\ 5r("S@SA\ 5r)"SBSC\ 5r*SDr+SEr,\+"5r-\,"\-5 g)Fz1 This scripts specifies all PTX special objects. N) defaultdict) Signature Parameterc,\rSrSrSrSrSrSrSrSr g)Stub zf A stub object to represent special objects that are meaningless outside the context of a CUDA kernel zc[SU-5e)Nz%s is not instantiableNotImplementedError)clss d/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/stubs.py__new__ Stub.__new__s!":S"@AAcUR$N) _description_selfs r__repr__ Stub.__repr__s!!!rN) __name__ __module__ __qualname____firstlineno____doc__r __slots__rr__static_attributes__r rrrr s*MIB"rrcF^[R"T5U4Sj5nU$)zj A stub function to represent special functions that are meaningless outside the context of a CUDA kernel c >[ST-5e)Nz"%s cannot be called from host coder )argskwargsfns rwrappedstub_function..wrappeds!"F"KLLr) functoolswraps)r$r%s` r stub_functionr)s' __RMM NrcL\rSrSrSrSr\S5r\S5r\S5r Sr g) Dim3)zA triple, (x, y, z)zcgrr rs rxDim3.x- rcgrr rs ryDim3.y1r0rcgrr rs rzDim3.z5r0rr N) rrrrrrpropertyr.r2r5rr rrr+r+)sDM       rr+c\rSrSrSrSrSrg) threadIdx:z The thread indices in the current thread block. Each index is an integer spanning the range from 0 inclusive to the corresponding value of the attribute in :attr:`numba.cuda.blockDim` exclusive. zr Nrrrrrrrr rrr9r9:s *Mrr9c\rSrSrSrSrSrg)blockIdxCz The block indices in the grid of thread blocks. Each index is an integer spanning the range from 0 inclusive to the corresponding value of the attribute in :attr:`numba.cuda.gridDim` exclusive. zr Nr;r rrr=r=C )Mrr=c\rSrSrSrSrSrg)blockDimLz The shape of a block of threads, as declared when instantiating the kernel. This value is the same for all threads in a given kernel launch, even if they belong to different blocks (i.e. each block is "full"). zr Nr;r rrrArALr?rrAc\rSrSrSrSrSrg)gridDimUzc The shape of the grid of blocks. This value is the same for all threads in a given kernel launch. zr Nr;r rrrDrDUs(MrrDc\rSrSrSrSrSrg)warpsize]zS The size of a warp. All architectures implemented to date have a warp size of 32. z r Nr;r rrrGrG]s!MrrGc\rSrSrSrSrSrg)laneidezU This thread's lane within a warp. Ranges from 0 to :attr:`numba.cuda.warpsize` - 1. zr Nr;r rrrJrJesMrrJc,\rSrSrSrSr\S5rSrg)sharedpz Shared memory namespace zcg)ad Allocate a shared array of the given *shape* and *type*. *shape* is either an integer or a tuple of integers representing the array's dimensions. *type* is a :ref:`Numba type ` of the elements needing to be stored in the array. The returned array-like object can be read and written to like any normal device array (e.g. through indexing). Nr shapedtypes rarray shared.arrayvrr N rrrrrrr)rSrr rrrMrMps M    rrMc,\rSrSrSrSr\S5rSrg)localz Local memory namespace zcg)z Allocate a local array of the given *shape* and *type*. The array is private to the current thread, and resides in global memory. An array-like object is returned which can be read and written to like any standard array (e.g. through indexing). Nr rPs rrS local.arrayrUrr NrVr rrrXrXs M  rrXc(\rSrSrSr\S5rSrg)constz Constant memory namespace cg)zx Create a const array from *ndarry*. The resulting const array will have the same shape, type, and values as *ndarray*. Nr )ndarrays r array_likeconst.array_likerUrr N)rrrrrr)rarr rrr]r]s  rr]c\rSrSrSrSrSrg)syncwarpzO syncwarp(mask=0xFFFFFFFF) Synchronizes a masked subset of threads in a warp. z r Nr;r rrrdrds $Mrrdc\rSrSrSrSrSrg)shfl_sync_intrinsicz shfl_sync_intrinsic(mask, mode, value, mode_offset, clamp) Nvvm intrinsic for shuffling data across a warp docs.nvidia.com/cuda/nvvm-ir-spec/index.html#nvvm-intrin-warp-level-datamove z r Nr;r rrrgrg $Mrrgc\rSrSrSrSrSrg)vote_sync_intrinsicz vote_sync_intrinsic(mask, mode, predictate) Nvvm intrinsic for performing a reduce and broadcast across a warp docs.nvidia.com/cuda/nvvm-ir-spec/index.html#nvvm-intrin-warp-level-vote z r Nr;r rrrkrkrirrkc\rSrSrSrSrSrg)match_any_syncz match_any_sync(mask, value) Nvvm intrinsic for performing a compare and broadcast across a warp. Returns a mask of threads that have same value as the given value from within the masked warp. zr Nr;r rrrnrns)Mrrnc\rSrSrSrSrSrg)match_all_syncar match_all_sync(mask, value) Nvvm intrinsic for performing a compare and broadcast across a warp. Returns a tuple of (mask, pred), where mask is a mask of threads that have same value as the given value from within the masked warp, if they all have the same value, otherwise it is 0. Pred is a boolean of whether or not all threads in the mask warp have the same warp. zr Nr;r rrrqrqs)Mrrqc\rSrSrSrSrSrg) activemaska= activemask() Returns a 32-bit integer mask of all currently active threads in the calling warp. The Nth bit is set if the Nth lane in the warp is active when activemask() is called. Inactive threads are represented by 0 bits in the returned mask. Threads which have exited the kernel are always marked as inactive. zr Nr;r rrrtrts%Mrrtc\rSrSrSrSrSrg) lanemask_ltzy lanemask_lt() Returns a 32-bit integer mask of all lanes (including inactive ones) with ID less than the current lane. zr Nr;r rrrwrws &Mrrwc\rSrSrSrSrSrg)threadfence_blockz& A memory fence at thread block level zr Nr;r rrrzrzs,Mrrzc\rSrSrSrSrSrg)threadfence_systemz0 A memory fence at system level: across devices zr Nr;r rrr}r}s-Mrr}c\rSrSrSrSrSrg) threadfencez A memory fence at device level zr Nr;r rrrrs&Mrrc\rSrSrSrSrg)popci z/ popc(x) Returns the number of set bits in x. r Nrrrrrrr rrrr rrc\rSrSrSrSrg)brevizc brev(x) Returns the reverse of the bit pattern of x. For example, 0b10110110 becomes 0b01101101. r Nrr rrrrrrc\rSrSrSrSrg)clziz3 clz(x) Returns the number of leading zeros in z. r Nrr rrrrrrrc\rSrSrSrSrg)ffsi#z ffs(x) Returns the position of the first (least significant) bit set to 1 in x, where the least significant bit position is 1. ffs(0) returns 0. r Nrr rrrr#rrrc\rSrSrSrSrg)selpi/zd selp(a, b, c) Select between source operands, based on the value of the predicate source operand. r Nrr rrrr/rrrc\rSrSrSrSrg)fmai;z9 fma(a, b, c) Perform the fused multiply-add operation. r Nrr rrrr;rrrc\rSrSrSrSrg)cbrtiCz," cbrt(a) Perform the cube root operation. r Nrr rrrrCrrrcP\rSrSrSrSr"SS\5r"SS\5r"SS \5r "S S \5r "S S \5r "SS\5r "SS\5r "SS\5r"SS\5r"SS\5r"SS\5r"SS\5r"SS\5r"SS\5rS rg!)"atomiciNz$Namespace for atomic operations zc\rSrSrSrSrg) atomic.addiSzadd(ary, idx, val) Perform atomic ``ary[idx] += val``. Supported on int32, float32, and float64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rraddrS rrc\rSrSrSrSrg) atomic.subi]zsub(ary, idx, val) Perform atomic ``ary[idx] -= val``. Supported on int32, float32, and float64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrsubr]rrrc\rSrSrSrSrg) atomic.and_igzand_(ary, idx, val) Perform atomic ``ary[idx] &= val``. Supported on int32, int64, uint32 and uint64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrand_rgrrrc\rSrSrSrSrg) atomic.or_iqzor_(ary, idx, val) Perform atomic ``ary[idx] |= val``. Supported on int32, int64, uint32 and uint64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rror_rqrrrc\rSrSrSrSrg) atomic.xori{zxor(ary, idx, val) Perform atomic ``ary[idx] ^= val``. Supported on int32, int64, uint32 and uint64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrxorr{rrrc\rSrSrSrSrg) atomic.incizinc(ary, idx, val) Perform atomic ``ary[idx] += 1`` up to val, then reset to 0. Supported on uint32, and uint64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrincrrrrc\rSrSrSrSrg) atomic.deciadec(ary, idx, val) Performs:: ary[idx] = (value if (array[idx] == 0) or (array[idx] > value) else array[idx] - 1) Supported on uint32, and uint64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrdecrs rrc\rSrSrSrSrg) atomic.exchizexch(ary, idx, val) Perform atomic ``ary[idx] = val``. Supported on int32, int64, uint32 and uint64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrexchrrrrc\rSrSrSrSrg) atomic.maxizmax(ary, idx, val) Perform atomic ``ary[idx] = max(ary[idx], val)``. Supported on int32, int64, uint32, uint64, float32, float64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrmaxr rrc\rSrSrSrSrg) atomic.minizmin(ary, idx, val) Perform atomic ``ary[idx] = min(ary[idx], val)``. Supported on int32, int64, uint32, uint64, float32, float64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrminrrrrc\rSrSrSrSrg) atomic.nanmaxia>nanmax(ary, idx, val) Perform atomic ``ary[idx] = max(ary[idx], val)``. NOTE: NaN is treated as a missing value such that: nanmax(NaN, n) == n, nanmax(n, NaN) == n Supported on int32, int64, uint32, uint64, float32, float64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrnanmaxr rrc\rSrSrSrSrg) atomic.nanminia?nanmin(ary, idx, val) Perform atomic ``ary[idx] = min(ary[idx], val)``. NOTE: NaN is treated as a missing value, such that: nanmin(NaN, n) == n, nanmin(n, NaN) == n Supported on int32, int64, uint32, uint64, float32, float64 operands only. Returns the old value at the index location as if it is loaded atomically. r Nrr rrnanminrrrrc\rSrSrSrSrg)atomic.compare_and_swapiacompare_and_swap(ary, old, val) Conditionally assign ``val`` to the first element of an 1D array ``ary`` if the current value matches ``old``. Supported on int32, int64, uint32, uint64 operands only. Returns the old value as if it is loaded atomically. r Nrr rrcompare_and_swapr rrc\rSrSrSrSrg) atomic.casiacas(ary, idx, old, val) Conditionally assign ``val`` to the element ``idx`` of an array ``ary`` if the current value of ``ary[idx]`` matches ``old``. Supported on int32, int64, uint32, uint64 operands only. Returns the old value as if it is loaded atomically. r Nrr rrcasrrrrr N)rrrrrrrrrrrrrrrrrrrrrrr rrrrNsM d  d  t  d  d  d   d   t   d    d            4    d  rrc\rSrSrSrSrSrg) nanosleepizz nanosleep(ns) Suspends the thread for a sleep duration approximately close to the delay `ns`, specified in nanoseconds. z r Nr;r rrrrs #Mrrc\rSrSrSrSr"SS\5r"SS\5r"SS \5r "S S \5r "S S \5r "SS\5r "SS\5r "SS\5r"SS\5r"SS\5r"SS\5r"SS\5r"SS\5r"SS\5r"S S!\5r"S"S#\5r"S$S%\5r"S&S'\5r"S(S)\5r"S*S+\5r"S,S-\5r"S.S/\5r"S0S1\5r"S2S3\5r"S4S5\5r"S6S7\5r "S8S9\5r!"S:S;\5r""S<S=\5r#"S>S?\5r$S@r%gA)Bfp16iz"Namespace for fp16 operations zc\rSrSrSrSrg) fp16.haddi zhadd(a, b) Perform fp16 addition, (a + b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the addition. r Nrr rrhaddr rrrc\rSrSrSrSrg) fp16.hsubizhsub(a, b) Perform fp16 subtraction, (a - b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the subtraction. r Nrr rrhsubrrrrc\rSrSrSrSrg) fp16.hmulizhmul(a, b) Perform fp16 multiplication, (a * b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the multiplication. r Nrr rrhmulrrrrc\rSrSrSrSrg) fp16.hdivi'zhdiv(a, b) Perform fp16 division, (a / b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the division r Nrr rrhdivr'rrrc\rSrSrSrSrg) fp16.hfmai1zhfma(a, b, c) Perform fp16 multiply and accumulate, (a * b) + c in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the multiplication. r Nrr rrhfmar1rrrc\rSrSrSrSrg) fp16.hnegi;zrhneg(a) Perform fp16 negation, -(a). Supported on fp16 operands only. Returns the fp16 result of the negation. r Nrr rrhnegr; rrc\rSrSrSrSrg) fp16.habsiDz}habs(a) Perform fp16 absolute value, |a|. Supported on fp16 operands only. Returns the fp16 result of the absolute value. r Nrr rrhabsrDrrrc\rSrSrSrSrg) fp16.hsiniMzshsin(a) Calculate sine in round to nearest even mode. Supported on fp16 operands only. Returns the sine result. r Nrr rrhsinrMrrrc\rSrSrSrSrg) fp16.hcosiWzwhsin(a) Calculate cosine in round to nearest even mode. Supported on fp16 operands only. Returns the cosine result. r Nrr rrhcosrWrrrc\rSrSrSrSrg) fp16.hlogiazhlog(a) Calculate natural logarithm in round to nearest even mode. Supported on fp16 operands only. Returns the natural logarithm result. r Nrr rrhlograrrrc\rSrSrSrSrg) fp16.hlog10ikzhlog10(a) Calculate logarithm base 10 in round to nearest even mode. Supported on fp16 operands only. Returns the logarithm base 10 result. r Nrr rrhlog10rkrrrc\rSrSrSrSrg) fp16.hlog2iuzhlog2(a) Calculate logarithm base 2 in round to nearest even mode. Supported on fp16 operands only. Returns the logarithm base 2 result. r Nrr rrhlog2rurrrc\rSrSrSrSrg) fp16.hexpizhexp(a) Calculate natural exponential, exp(a), in round to nearest mode. Supported on fp16 operands only. Returns the natural exponential result. r Nrr rrhexprrrrc\rSrSrSrSrg) fp16.hexp10izhexp10(a) Calculate exponential base 10 (10 ** a) in round to nearest mode. Supported on fp16 operands only. Returns the exponential base 10 result. r Nrr rrhexp10rrrrc\rSrSrSrSrg) fp16.hexp2izhexp2(a) Calculate exponential base 2 (2 ** a) in round to nearest mode. Supported on fp16 operands only. Returns the exponential base 2 result. r Nrr rrhexp2rrrrc\rSrSrSrSrg) fp16.hfloorizhfloor(a) Calculate the floor, the largest integer less than or equal to 'a'. Supported on fp16 operands only. Returns the floor result. r Nrr rrhfloorrrrrc\rSrSrSrSrg) fp16.hceilizhceil(a) Calculate the ceil, the smallest integer greater than or equal to 'a'. Supported on fp16 operands only. Returns the ceil result. r Nrr rrhceilrrrrc\rSrSrSrSrg) fp16.hsqrtizhsqrt(a) Calculate the square root of the input argument in round to nearest mode. Supported on fp16 operands only. Returns the square root result. r Nrr rrhsqrtrrrrc\rSrSrSrSrg) fp16.hrsqrtizhrsqrt(a) Calculate the reciprocal square root of the input argument in round to nearest even mode. Supported on fp16 operands only. Returns the reciprocal square root result. r Nrr rrhrsqrtrrrrc\rSrSrSrSrg) fp16.hrcpizhrcp(a) Calculate the reciprocal of the input argument in round to nearest even mode. Supported on fp16 operands only. Returns the reciprocal result. r Nrr rrhrcprrrrc\rSrSrSrSrg) fp16.hrintiz|hrint(a) Round the input argument to nearest integer value. Supported on fp16 operands only. Returns the rounded result. r Nrr rrhrintrrrrc\rSrSrSrSrg) fp16.htruncizhtrunc(a) Truncate the input argument to its integer portion. Supported on fp16 operands only. Returns the truncated result. r Nrr rrhtruncrrrr c\rSrSrSrSrg)fp16.heqizheq(a, b) Perform fp16 comparison, (a == b). Supported on fp16 operands only. Returns True if a and b are equal and False otherwise. r Nrr rrheqr rrr c\rSrSrSrSrg)fp16.hneizhne(a, b) Perform fp16 comparison, (a != b). Supported on fp16 operands only. Returns True if a and b are not equal and False otherwise. r Nrr rrhnerrrrc\rSrSrSrSrg)fp16.hgeizhge(a, b) Perform fp16 comparison, (a >= b). Supported on fp16 operands only. Returns True if a is >= b and False otherwise. r Nrr rrhgerrrrc\rSrSrSrSrg)fp16.hgtiz~hgt(a, b) Perform fp16 comparison, (a > b). Supported on fp16 operands only. Returns True if a is > b and False otherwise. r Nrr rrhgtrrrrc\rSrSrSrSrg)fp16.hlei zhle(a, b) Perform fp16 comparison, (a <= b). Supported on fp16 operands only. Returns True if a is <= b and False otherwise. r Nrr rrhler rrrc\rSrSrSrSrg)fp16.hltiz~hlt(a, b) Perform fp16 comparison, (a < b). Supported on fp16 operands only. Returns True if a is < b and False otherwise. r Nrr rrhltrrrrc\rSrSrSrSrg) fp16.hmaxizhmax(a, b) Perform fp16 maximum operation, max(a,b) Supported on fp16 operands only. Returns a if a is greater than b, returns b otherwise. r Nrr rrhmaxrrrrc\rSrSrSrSrg) fp16.hmini)zhmin(a, b) Perform fp16 minimum operation, min(a,b). Supported on fp16 operands only. Returns a if a is less than b, returns b otherwise. r Nrr rrhminr )rrr!r N)&rrrrrrrrrrrrrrrrrrrrrrrrrrrrr r rrrrrrr!rr rrrrstM t  t  t  t  t  t  t  t  t  t      t              t      d  d  d  d  d  d  t  t rrct/nSnSnSn[R"X5HupEUSU3nUSUn[U[40UVs0sHoS_M snESUS3[ USUV s/sHn [ U [ R S 9PM sn S 9S US 3S .ES/0E5n URU 5 M U$s snfs sn f)z)Make user facing objects for vector types) int8int16int32int64uint8uint16uint32uint64float32float64))r.r2r5wr.Ncgrr rs r(make_vector_type_stubs..Rsdr<>)namekind) parametersz A stub for z to be used in CUDA kernels.)r __signature__raliases) itertoolsproducttyperrrPOSITIONAL_ONLYappend) vector_type_stubsvector_type_prefixvector_type_element_countsvector_type_attribute_namesprefixnelem type_name attr_namesattr attr_namevector_type_stubs rmake_vector_type_stubsrL7s# "."6"** haw' 0%8  w 7ABzt**zB () 1%5%.,6fu+=;,>i"!*1J1J+=;& "-YK8$ $  b/  "   !12/0 C;s  B0""B5cS[R"[R5RS-3S[R"[R5RS-3S[R"[R 5RS-3S[R"[R 5RS-3S[R"[R5RS-3S[R"[R5RS-3S[R"[R5RS-3S[R"[R5RS-3S[R"[R5RS-3S[R"[R5RS-3S[R"[R5RS-3S[R"[R5RS-3S. n[[ 5nUH#nX#R"SSR%U5 M% UR'5H@upEX%nUH1nUR"SnUR(R%UU35 M3 MB g) zNFor each of the stubs, create its aliases. For example: float64x3 -> double3 intuintfloat) charshortrNlonglonglongucharushortrPulong ulonglongrQdoubleN)nprRbyteitemsizerSintcint_rUubyterWuintcrPrYsinglerZrlistrr@itemsr;)rAbase_type_to_aliasbase_type_to_vector_typestubalias base_typerFs rmap_vector_type_stubs_to_aliasrlds7bhhrww'001456rxx)22Q678RXXbgg&//!345bhhrww'001456"((2;;/881<=>*33a789")),559:;rxx)22Q678)22Q678BHHR\\2;;a?@A")),559:;"((299-66:;<  +40! s!34;;DA"/4464?%DMM"%E LL  5'% 1 2&7r).rnumpyr] collectionsrr'r<inspectrrobjectrr)r+r9r=rArDrGrJrMrXr]rdrgrkrnrqrtrwrzr}rrrrrrrrrrrrLrl_vector_type_stubsr rrrrs#( "6 " 4 "**)t))t)(d(!t!T T &  D   D  $t$$$$$$$)T) )T ) % %&$&,,--&$&44$$4$4c Tc R##m 4m f *Z3B,-12r