J/PhWdZddlZddlmZddlZddlZddlmZm Z Gdde Z dZ Gdd e Z Gd d e ZGd d e ZGdde ZGdde ZGdde ZGdde ZGdde ZGdde ZGdde ZGdde ZGdde ZGd d!e ZGd"d#e ZGd$d%e ZGd&d'e ZGd(d)e ZGd*d+e ZGd,d-e ZGd.d/e Z Gd0d1e Z!Gd2d3e Z"Gd4d5e Z#Gd6d7e Z$Gd8d9e Z%Gd:d;e Z&Gd<d=e Z'Gd>d?e Z(Gd@dAe Z)GdBdCe Z*dDZ+dEZ,e+Z-e,e-dS)Fz1 This scripts specifies all PTX special objects. N) defaultdict) Signature Parameterc&eZdZdZdZdZdZdZdS)Stubzr A stub object to represent special objects that are meaningless outside the context of a CUDA kernel zc&td|z)Nz%s is not instantiableNotImplementedError)clss P/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/stubs.py__new__z Stub.__new__s!":S"@AAAc|jSN) _description_selfs r __repr__z Stub.__repr__s !!rN)__name__ __module__ __qualname____doc__r __slots__rrrrr rr sK*MIBBB"""""rrcFtjfd}|S)zv A stub function to represent special functions that are meaningless outside the context of a CUDA kernel c(tdz)Nz"%s cannot be called from host coder )argskwargsfns r wrappedzstub_function..wrappeds!"F"KLLLr) functoolswraps)rr s` r stub_functionr#s= _RMMMMM NrcXeZdZdZdZedZedZedZdS)Dim3zA triple, (x, y, z)zcdSrrrs r xzDim3.x- rcdSrrrs r yzDim3.y1r(rcdSrrrs r zzDim3.z5r(rN) rrrrrpropertyr'r*r,rrr r%r%)snM   X   X   X   rr%ceZdZdZdZdS) threadIdxz 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. zNrrrrrrrr r/r/:s *MMMrr/ceZdZdZdZdS)blockIdxz 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. zNr0rrr r2r2C )MMMrr2ceZdZdZdZdS)blockDimz 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"). zNr0rrr r5r5Lr3rr5ceZdZdZdZdS)gridDimzo The shape of the grid of blocks. This value is the same for all threads in a given kernel launch. zNr0rrr r7r7Us(MMMrr7ceZdZdZdZdS)warpsizez_ The size of a warp. All architectures implemented to date have a warp size of 32. z Nr0rrr r9r9]s!MMMrr9ceZdZdZdZdS)laneidza This thread's lane within a warp. Ranges from 0 to :attr:`numba.cuda.warpsize` - 1. zNr0rrr r;r;esMMMrr;c,eZdZdZdZedZdS)sharedz! Shared memory namespace zcdS)a 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). Nrshapedtypes r arrayz shared.arrayvrNrrrrrr#rBrrr r=r=ps>M   ]    rr=c,eZdZdZdZedZdS)localz Local memory namespace zcdS)a 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). Nrr?s r rBz local.arrayrCrNrDrrr rFrFs>M  ]   rrFc(eZdZdZedZdS)constz# Constant memory namespace cdS)z Create a const array from *ndarry*. The resulting const array will have the same shape, type, and values as *ndarray*. Nr)ndarrays r array_likezconst.array_likerCrN)rrrrr#rLrrr rIrIs9  ]   rrIceZdZdZdZdS)syncwarpz[ syncwarp(mask=0xFFFFFFFF) Synchronizes a masked subset of threads in a warp. z Nr0rrr rNrNs $MMMrrNceZdZdZdZdS)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 Nr0rrr rPrP $MMMrrPceZdZdZdZdS)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 Nr0rrr rSrSrQrrSceZdZdZdZdS)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. zNr0rrr rUrUs)MMMrrUceZdZdZdZdS)match_all_synca 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. zNr0rrr rWrWs)MMMrrWceZdZdZdZdS) activemaskaY 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. zNr0rrr rYrYs%MMMrrYceZdZdZdZdS) lanemask_ltz lanemask_lt() Returns a 32-bit integer mask of all lanes (including inactive ones) with ID less than the current lane. zNr0rrr r[r[s &MMMrr[ceZdZdZdZdS)threadfence_blockz. A memory fence at thread block level zNr0rrr r]r]s,MMMrr]ceZdZdZdZdS)threadfence_systemz8 A memory fence at system level: across devices zNr0rrr r_r_s-MMMrr_ceZdZdZdZdS) threadfencez( A memory fence at device level zNr0rrr raras&MMMrraceZdZdZdS)popcz; popc(x) Returns the number of set bits in x. Nrrrrrrr rcrc rrcceZdZdZdS)brevzs brev(x) Returns the reverse of the bit pattern of x. For example, 0b10110110 becomes 0b01101101. Nrdrrr rgrgrrgceZdZdZdS)clzz? clz(x) Returns the number of leading zeros in z. Nrdrrr rjrjrerrjceZdZdZdS)ffsz 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. Nrdrrr rlrl#rhrrlceZdZdZdS)selpzt selp(a, b, c) Select between source operands, based on the value of the predicate source operand. Nrdrrr rnrn/rhrrnceZdZdZdS)fmazE fma(a, b, c) Perform the fused multiply-add operation. Nrdrrr rprp;rerrpceZdZdZdS)cbrtz8" cbrt(a) Perform the cube root operation. Nrdrrr rrrrCrerrrceZdZdZdZGddeZGddeZGddeZGd d eZ Gd d eZ Gd deZ GddeZ GddeZ GddeZGddeZGddeZGddeZGddeZGddeZdS) atomicz$Namespace for atomic operations zceZdZdZdS) atomic.addzadd(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. Nrdrrr addrvS    rrwceZdZdZdS) atomic.subzsub(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. Nrdrrr subrz]rxrr{ceZdZdZdS) atomic.and_zand_(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. Nrdrrr and_r}grxrr~ceZdZdZdS) atomic.or_zor_(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. Nrdrrr or_rqrxrrceZdZdZdS) atomic.xorzxor(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. Nrdrrr xorr{rxrrceZdZdZdS) atomic.inczinc(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. Nrdrrr incrrxrrceZdZdZdS) atomic.deca:dec(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. Nrdrrr decrs rrceZdZdZdS) atomic.exchzexch(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. Nrdrrr exchrrxrrceZdZdZdS) atomic.maxamax(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. Nrdrrr maxr rrceZdZdZdS) atomic.minamin(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. Nrdrrr minrrrrceZdZdZdS) atomic.nanmaxa~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. Nrdrrr nanmaxr rrceZdZdZdS) atomic.nanminananmin(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. Nrdrrr nanminrrrrceZdZdZdS)atomic.compare_and_swapa(compare_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. Nrdrrr compare_and_swapr    rrceZdZdZdS) atomic.casa/cas(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. Nrdrrr casrrrrN)rrrrrrrwr{r~rrrrrrrrrrrrrr rtrtNsoM     d        d        t        d        d        d         d         t         d          d                              4          d      rrtceZdZdZdZdS) nanosleepz nanosleep(ns) Suspends the thread for a sleep duration approximately close to the delay `ns`, specified in nanoseconds. z Nr0rrr rrs #MMMrrc^eZdZdZdZGddeZGddeZGddeZGd d eZ Gd d eZ Gd deZ GddeZ GddeZ GddeZGddeZGddeZGddeZGddeZGddeZGdd eZGd!d"eZGd#d$eZGd%d&eZGd'd(eZGd)d*eZGd+d,eZGd-d.eZGd/d0eZGd1d2eZGd3d4eZGd5d6eZGd7d8eZ Gd9d:eZ!Gd;deZ#d?S)@fp16z"Namespace for fp16 operations zceZdZdZdS) fp16.haddzhadd(a, b) Perform fp16 addition, (a + b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the addition. Nrdrrr haddr rxrrceZdZdZdS) fp16.hsubzhsub(a, b) Perform fp16 subtraction, (a - b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the subtraction. Nrdrrr hsubrrxrrceZdZdZdS) fp16.hmulzhmul(a, b) Perform fp16 multiplication, (a * b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the multiplication. Nrdrrr hmulrrxrrceZdZdZdS) fp16.hdivzhdiv(a, b) Perform fp16 division, (a / b) in round to nearest mode. Supported on fp16 operands only. Returns the fp16 result of the division Nrdrrr hdivr'rxrrceZdZdZdS) fp16.hfmazhfma(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. Nrdrrr hfmar1rxrrceZdZdZdS) fp16.hnegzhneg(a) Perform fp16 negation, -(a). Supported on fp16 operands only. Returns the fp16 result of the negation. Nrdrrr hnegr;    rrceZdZdZdS) fp16.habszhabs(a) Perform fp16 absolute value, |a|. Supported on fp16 operands only. Returns the fp16 result of the absolute value. Nrdrrr habsrDrrrceZdZdZdS) fp16.hsinzhsin(a) Calculate sine in round to nearest even mode. Supported on fp16 operands only. Returns the sine result. Nrdrrr hsinrMrxrrceZdZdZdS) fp16.hcoszhsin(a) Calculate cosine in round to nearest even mode. Supported on fp16 operands only. Returns the cosine result. Nrdrrr hcosrWrxrrceZdZdZdS) fp16.hlogzhlog(a) Calculate natural logarithm in round to nearest even mode. Supported on fp16 operands only. Returns the natural logarithm result. Nrdrrr hlograrxrrceZdZdZdS) fp16.hlog10zhlog10(a) Calculate logarithm base 10 in round to nearest even mode. Supported on fp16 operands only. Returns the logarithm base 10 result. Nrdrrr hlog10rkrxrrceZdZdZdS) fp16.hlog2zhlog2(a) Calculate logarithm base 2 in round to nearest even mode. Supported on fp16 operands only. Returns the logarithm base 2 result. Nrdrrr hlog2rurxrrceZdZdZdS) fp16.hexpzhexp(a) Calculate natural exponential, exp(a), in round to nearest mode. Supported on fp16 operands only. Returns the natural exponential result. Nrdrrr hexprrxrrceZdZdZdS) fp16.hexp10zhexp10(a) Calculate exponential base 10 (10 ** a) in round to nearest mode. Supported on fp16 operands only. Returns the exponential base 10 result. Nrdrrr hexp10rrxrrceZdZdZdS) fp16.hexp2zhexp2(a) Calculate exponential base 2 (2 ** a) in round to nearest mode. Supported on fp16 operands only. Returns the exponential base 2 result. Nrdrrr hexp2rrxrrceZdZdZdS) fp16.hfloorzhfloor(a) Calculate the floor, the largest integer less than or equal to 'a'. Supported on fp16 operands only. Returns the floor result. Nrdrrr hfloorrrxrrceZdZdZdS) fp16.hceilzhceil(a) Calculate the ceil, the smallest integer greater than or equal to 'a'. Supported on fp16 operands only. Returns the ceil result. Nrdrrr hceilrrxrrceZdZdZdS) fp16.hsqrtzhsqrt(a) Calculate the square root of the input argument in round to nearest mode. Supported on fp16 operands only. Returns the square root result. Nrdrrr hsqrtrrxrrceZdZdZdS) fp16.hrsqrtzhrsqrt(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. Nrdrrr hrsqrtrrxrrceZdZdZdS) fp16.hrcpzhrcp(a) Calculate the reciprocal of the input argument in round to nearest even mode. Supported on fp16 operands only. Returns the reciprocal result. Nrdrrr hrcprrxrrceZdZdZdS) fp16.hrintzhrint(a) Round the input argument to nearest integer value. Supported on fp16 operands only. Returns the rounded result. Nrdrrr hrintrrxrrceZdZdZdS) fp16.htrunczhtrunc(a) Truncate the input argument to its integer portion. Supported on fp16 operands only. Returns the truncated result. Nrdrrr htruncrrxrrceZdZdZdS)fp16.heqzheq(a, b) Perform fp16 comparison, (a == b). Supported on fp16 operands only. Returns True if a and b are equal and False otherwise. Nrdrrr heqrrxrrceZdZdZdS)fp16.hnezhne(a, b) Perform fp16 comparison, (a != b). Supported on fp16 operands only. Returns True if a and b are not equal and False otherwise. Nrdrrr hnerrxrrceZdZdZdS)fp16.hgezhge(a, b) Perform fp16 comparison, (a >= b). Supported on fp16 operands only. Returns True if a is >= b and False otherwise. Nrdrrr hgerrxrrceZdZdZdS)fp16.hgtzhgt(a, b) Perform fp16 comparison, (a > b). Supported on fp16 operands only. Returns True if a is > b and False otherwise. Nrdrrr hgtrrxrrceZdZdZdS)fp16.hlezhle(a, b) Perform fp16 comparison, (a <= b). Supported on fp16 operands only. Returns True if a is <= b and False otherwise. Nrdrrr hler rxrrceZdZdZdS)fp16.hltzhlt(a, b) Perform fp16 comparison, (a < b). Supported on fp16 operands only. Returns True if a is < b and False otherwise. Nrdrrr hltrrxrrceZdZdZdS) fp16.hmaxzhmax(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. Nrdrrr hmaxrrxrrceZdZdZdS) fp16.hminzhmin(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. Nrdrrr hminr)rxrrN)$rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr rrsM     t        t        t        t        t        t        t        t        t        t                        t                                                        t                        d        d        d        d        d        d        t        t     rrc Bg}d}d}d}tj||D]\}}|d|}|d|}t|tfid|Dd|dt d |d|D d |d d dgi}|||S)z)Make user facing objects for vector types) int8int16int32int64uint8uint16uint32uint64float32float64))r'r*r,wr'Nci|]}|dS)cdSrrrs r z3make_vector_type_stubs...Rsdrr).0attrs r z*make_vector_type_stubs..RsBBBt4**BBBr<>cDg|]}t|tjS))namekind)rPOSITIONAL_ONLY)r attr_names r z*make_vector_type_stubs..UsA;;;("!*1J;;;r) parametersz A stub for z to be used in CUDA kernels.)r __signature__raliases) itertoolsproducttyperrappend) vector_type_stubsvector_type_prefixvector_type_element_countsvector_type_attribute_namesprefixnelem type_name attr_namesvector_type_stubs r make_vector_type_stubsr07s5 "."6"*633 '''' 0%8  w BBzBBB &6%5%5%5%.;;,6fuf+=;;;&&&  $Y $ $ $   b/   "   !12222 rcdtjtjjdzdtjtjjdzdtjtjjdzdtjtjjdzdtjtjjdzdtjtjjdzdtjtj jdzdtjtj jdzdtjtj jdzdtjtj jdzdtjtj jdzdtjtjjdzd }tt }|D]*}||jdd|+|D]=\}}||}|D]-}|jd}|j||.>dS) zVFor each of the stubs, create its aliases. For example: float64x3 -> double3 intuintfloat) charshortr2longlonglongucharushortr4ulong ulonglongr5doubleN)nprAbyteitemsizer7intcint_r9ubyter;uintcr4r=singler>rlistrr&itemsr")r'base_type_to_aliasbase_type_to_vector_typestubalias base_typer,s r map_vector_type_stubs_to_aliasrPdsF7bhrw''014668rx))2Q6885RXbg&&/!3556bhrw''01466>"(2;//81<>>9**3a799;"),,59;;8rx))2Q6888))2Q688ABHR\22;a?AA;"),,59;;<"(29--6:<<   +400!BB ss!34;;DAAAA.446633y4Y?% 3 3DM"%E L  5 1% 1 1 2 2 2 2 333r).rnumpyrA collectionsrr!r#inspectrrobjectrr#r%r/r2r5r7r9r;r=rFrIrNrPrSrUrWrYr[r]r_rarcrgrjrlrnrprrrtrrr0rP_vector_type_stubsrrr rVs######(((((((( " " " " "6 " " "     4   "********)))))t))))))))t)))(((((d(((!!!!!t!!!T     T   &      D         D    $$$$$t$$$$$$$$$$$$$$$$$$$$$)))))T))) ) ) ) ) )T ) ) ) % % % % % % % %&&&&&$&&&,,,,,,,,--------&&&&&$&&&44$$4$4c c c c c Tc c c R########m m m m m 4m m m f ***Z333B,+--122222r