U ,dR@sdZddlZddlmZddlZddlZddlmZm Z Gddde Z ddZ Gd d d e Z Gd d d e ZGd dde ZGddde ZGddde ZGddde ZGddde ZGddde ZGddde ZGddde ZGddde ZGdd d e ZGd!d"d"e ZGd#d$d$e ZGd%d&d&e ZGd'd(d(e ZGd)d*d*e ZGd+d,d,e ZGd-d.d.e ZGd/d0d0e Z Gd1d2d2e Z!Gd3d4d4e Z"Gd5d6d6e Z#Gd7d8d8e Z$Gd9d:d:e Z%Gd;d<dd>e Z'Gd?d@d@e Z(GdAdBdBe Z)GdCdDdDe Z*GdEdFdFe Z+GdGdHdHe Z,GdIdJdJe Z-GdKdLdLe Z.GdMdNdNe Z/GdOdPdPe Z0GdQdRdRe Z1dSdTZ2dUdVZ3e2Z4e3e4dS)Wz1 This scripts specifies all PTX special objects. N) defaultdict) Signature Parameterc@s(eZdZdZdZdZddZddZdS) Stubzr A stub object to represent special objects that are meaningless outside the context of a CUDA kernel zcCstd|dS)Nz%s is not instantiableNotImplementedError)clsrr4/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/stubs.py__new__sz Stub.__new__cCs|jSN) _description_selfrrr __repr__sz Stub.__repr__N)__name__ __module__ __qualname____doc__r __slots__r rrrrr r s rcstfdd}|S)zv A stub function to represent special functions that are meaningless outside the context of a CUDA kernel cstddS)Nz"%s cannot be called from host coder)argskwargsfnrr wrappedszstub_function..wrapped) functoolswraps)rrrrr stub_functionsrc@s8eZdZdZdZeddZeddZeddZd S) Dim3zA triple, (x, y, z)zcCsdSr rrrrr x-szDim3.xcCsdSr rrrrr y1szDim3.ycCsdSr rrrrr z5szDim3.zN) rrrrr propertyrr r!rrrr r)s  rc@seZdZdZdZdS) 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. zNrrrrr rrrr r#:sr#c@seZdZdZdZdS)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. zNr$rrrr r%Csr%c@seZdZdZdZdS)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"). zNr$rrrr r&Lsr&c@seZdZdZdZdS)gridDimzo The shape of the grid of blocks. This value is the same for all threads in a given kernel launch. zNr$rrrr r'Usr'c@seZdZdZdZdS)warpsizez_ The size of a warp. All architectures implemented to date have a warp size of 32. z Nr$rrrr r(]sr(c@seZdZdZdZdS)laneidza This thread's lane within a warp. Ranges from 0 to :attr:`numba.cuda.warpsize` - 1. zNr$rrrr r)esr)c@seZdZdZdZdS)gridagrid(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. z Nr$rrrr r*msr*c@seZdZdZdZdS)gridsizeagridsize(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. zNr$rrrr r+sr+c@s eZdZdZdZeddZdS)sharedz! Shared memory namespace zcCsdS)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). Nrshapedtyperrr arraysz shared.arrayNrrrrr rr0rrrr r,sr,c@s eZdZdZdZeddZdS)localz Local memory namespace zcCsdS)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-rrr r0sz local.arrayNr1rrrr r2sr2c@seZdZdZeddZdS)constz# Constant memory namespace cCsdS)z Create a const array from *ndarry*. The resulting const array will have the same shape, type, and values as *ndarray*. Nr)Zndarrayrrr array_likeszconst.array_likeN)rrrrrr4rrrr r3sr3c@s,eZdZdZeddZGdddeZdS)cgz Cooperative groups cCsdS)z- Get the current grid group. Nrrrrr this_gridsz cg.this_gridc@seZdZddZdS)z cg.GridGroupcCsdS)z= Synchronize the current grid group. Nrrrrr syncszcg.GridGroup.syncN)rrrr7rrrr GridGroupsr8N)rrrrrr6rr8rrrr r5s r5c@seZdZdZdZdS) syncthreadsa 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. zNr$rrrr r9sr9c@seZdZdZdZdS)syncthreads_countz syncthreads_count(predictate) An extension to numba.cuda.syncthreads where the return value is a count of the threads where predicate is true. zNr$rrrr r:sr:c@seZdZdZdZdS)syncthreads_andz syncthreads_and(predictate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for all threads or 0 otherwise. zNr$rrrr r;sr;c@seZdZdZdZdS)syncthreads_orz syncthreads_or(predictate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for any thread or 0 otherwise. zNr$rrrr r<sr<c@seZdZdZdZdS)syncwarpz[ syncwarp(mask=0xFFFFFFFF) Synchronizes a masked subset of threads in a warp. z Nr$rrrr r=sr=c@seZdZdZdZdS)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 Nr$rrrr r>sr>c@seZdZdZdZdS)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 Nr$rrrr r?sr?c@seZdZdZdZdS)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. zNr$rrrr r@$sr@c@seZdZdZdZdS)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. zNr$rrrr rA/s rAc@seZdZdZdZdS) 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. zNr$rrrr rB<s rBc@seZdZdZdZdS) lanemask_ltz lanemask_lt() Returns a 32-bit integer mask of all lanes (including inactive ones) with ID less than the current lane. zNr$rrrr rCIsrCc@seZdZdZdZdS)threadfence_blockz. A memory fence at thread block level zNr$rrrr rDVsrDc@seZdZdZdZdS)threadfence_systemz8 A memory fence at system level: across devices zNr$rrrr rE]srEc@seZdZdZdZdS) threadfencez( A memory fence at device level zNr$rrrr rFdsrFc@seZdZdZdS)popcz; popc(x) Returns the number of set bits in x. Nrrrrrrrr rGnsrGc@seZdZdZdS)brevzs brev(x) Returns the reverse of the bit pattern of x. For example, 0b10110110 becomes 0b01101101. NrHrrrr rIvsrIc@seZdZdZdS)clzz? clz(x) Returns the number of leading zeros in z. NrHrrrr rJsrJc@seZdZdZdS)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. NrHrrrr rKsrKc@seZdZdZdS)selpzt selp(a, b, c) Select between source operands, based on the value of the predicate source operand. NrHrrrr rLsrLc@seZdZdZdS)fmazE fma(a, b, c) Perform the fused multiply-add operation. NrHrrrr rMsrMc@seZdZdZdS)cbrtz8" cbrt(a) Perform the cube root operation. NrHrrrr rNsrNc@seZdZdZdZGdddeZGdddeZGdddeZGd d d eZ Gd d d eZ Gd ddeZ GdddeZ GdddeZ GdddeZGdddeZGdddeZGdddeZGdddeZdS)atomicz$Namespace for atomic operations zc@seZdZdZdS)z 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. NrHrrrr addsrPc@seZdZdZdS)z 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. NrHrrrr subsrQc@seZdZdZdS)z 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. NrHrrrr and_srRc@seZdZdZdS)z 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. NrHrrrr or_srSc@seZdZdZdS)z 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. NrHrrrr xorsrTc@seZdZdZdS)z 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. NrHrrrr incsrUc@seZdZdZdS)z atomic.decadec(ary, idx, val) Perform 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. NrHrrrr decsrVc@seZdZdZdS)z 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. NrHrrrr exchsrWc@seZdZdZdS)z atomic.maxa max(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. NrHrrrr max srXc@seZdZdZdS)z atomic.mina min(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. NrHrrrr minsrYc@seZdZdZdS)z atomic.nanmaxaznanmax(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. NrHrrrr nanmax!srZc@seZdZdZdS)z atomic.nanmina{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. NrHrrrr nanmin0sr[c@seZdZdZdS)zatomic.compare_and_swapzcompare_and_swap(ary, old, val) Conditionally assign ``val`` to the first element of an 1D array ``ary`` if the current value matches ``old``. Returns the current value as if it is loaded atomically. NrHrrrr compare_and_swap?sr\N)rrrrr rrPrQrRrSrTrUrVrWrXrYrZr[r\rrrr rOs          rOc@seZdZdZdZdS) nanosleepz nanosleep(ns) Suspends the thread for a sleep duration approximately close to the delay `ns`, specified in nanoseconds. z Nr$rrrr r]Lsr]c@seZdZdZdZGdddeZGdddeZGdddeZGd d d eZ Gd d d eZ Gd ddeZ GdddeZ GdddeZ GdddeZGdddeZGdddeZGdddeZGdddeZGdddeZdS) fp16z"Namespace for fp16 operations zc@seZdZdZdS)z 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. NrHrrrr hadd^sr_c@seZdZdZdS)z 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. NrHrrrr hsubhsr`c@seZdZdZdS)z 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. NrHrrrr hmulrsrac@seZdZdZdS)z 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. NrHrrrr hfma|srbc@seZdZdZdS)z fp16.hnegzhneg(a) Perform fp16 negation, -(a). Supported on fp16 operands only. Returns the fp16 result of the negation. NrHrrrr hnegsrcc@seZdZdZdS)z fp16.habszhabs(a) Perform fp16 absolute value, |a|. Supported on fp16 operands only. Returns the fp16 result of the absolute value. NrHrrrr habssrdc@seZdZdZdS)zfp16.heqzheq(a, b) Perform fp16 comparison, (a == b). Supported on fp16 operands only. Returns True if a and b are equal and False otherwise. NrHrrrr heqsrec@seZdZdZdS)zfp16.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. NrHrrrr hnesrfc@seZdZdZdS)zfp16.hgezhge(a, b) Perform fp16 comparison, (a >= b). Supported on fp16 operands only. Returns True if a is >= b and False otherwise. NrHrrrr hgesrgc@seZdZdZdS)zfp16.hgtzhgt(a, b) Perform fp16 comparison, (a > b). Supported on fp16 operands only. Returns True if a is > b and False otherwise. NrHrrrr hgtsrhc@seZdZdZdS)zfp16.hlezhle(a, b) Perform fp16 comparison, (a <= b). Supported on fp16 operands only. Returns True if a is <= b and False otherwise. NrHrrrr hlesric@seZdZdZdS)zfp16.hltzhlt(a, b) Perform fp16 comparison, (a < b). Supported on fp16 operands only. Returns True if a is < b and False otherwise. NrHrrrr hltsrjc@seZdZdZdS)z 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. NrHrrrr hmaxsrkc@seZdZdZdS)z 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. NrHrrrr hminsrlN)rrrrr rr_r`rarbrcrdrerfrgrhrirjrkrlrrrr r^Ys              r^c Csg}d}d}d}t||D]~\}}|d|}|d|}t|tfdd|Dd|d td d |d|Dd d |dddgi}||q|S)z)Make user facing objects for vector types) Zint8Zint16Zint32Zint64Zuint8Zuint16Zuint32Zuint64Zfloat32Zfloat64))rr r!wrNcSsi|]}|ddqS)cSsdSr rrrrr z3make_vector_type_stubs...r).0attrrrr sz*make_vector_type_stubs..<>cSsg|]}t|tjdqS))namekind)rPOSITIONAL_ONLY)rt attr_namerrr s z*make_vector_type_stubs..) parametersz A stub for z to be used in CUDA kernels.)r __signature__raliases) itertoolsproducttyperrappend) vector_type_stubsZvector_type_prefixZvector_type_element_countsZvector_type_attribute_namesprefixnelem type_nameZ attr_namesZvector_type_stubrrr make_vector_type_stubss2         rcCsdttjjddttjjddttjjddttjjddttjjddttjjddttj jddttj jddttj jddttj jddttj jddttjjdd }tt}|D]}||jdd|q|D]<\}}||}|D]$}|jd}|j||qVqBdS) zVFor each of the stubs, create its aliases. For example: float64x3 -> double3 intuintfloat) charshortrlonglonglongZucharushortrZulong ulonglongrdoubleN)npr/byteitemsizerZintcint_rZubyterZuintcrrZsinglerrlistrritemsr)rZbase_type_to_aliasZbase_type_to_vector_typeZstubaliasZ base_typerrrr map_vector_type_stubs_to_aliass* r)5rZnumpyr collectionsrrrinspectrrobjectrrrr#r%r&r'r(r)r*r+r,r2r3r5r9r:r;r<r=r>r?r@rArBrCrDrErFrGrIrJrKrLrMrNrOr]r^rrZ_vector_type_stubsrrrr sd                     -!