B \z5@sdZddlmZmZmZddlZddlZddlmm Z ddl m Z m Z mZmZddlmZGdddeZee jZGd d d eZGd d d eZGd ddeZGdddeZedeZedeZddZddZddZ ejde ddZ!ddZ"ejde"ddZ#Gdd d eZ$Gd!d"d"eZ%Gd#d$d$eZ&Gd%d&d&eZ'Gd'd(d(eZ(Gd)d*d*eZ)Gd+d,d,eZ*Gd-d.d.eZ+Gd/d0d0eZ,Gd1d2d2eZ-Gd3d4d4eZ.Gd5d6d6eZ/d7d8Z0d9d:Z1Gd;d<dZ3Gd?d@d@eZ4dAdBZ5GdCdDdDeZ6GdEdFdFeZ7GdGdHdHeZ8GdIdJdJeZ9GdKdLdLeZ:GdMdNdNeZ;GdOdPdPeZcCstd|dS)Nz%s is not instantiable)NotImplementedError)clsr r /lib/python3.7/site-packages/numba/cuda/stubs.py__new__sz Stub.__new__cCs|jS)N) _description_)selfr r r__repr__sz Stub.__repr__N)__name__ __module__ __qualname____doc__r __slots__rrr r r rr s r c@s8eZdZdZdZedeZedeZ edeZ dS) threadIdxa The thread indices in the current thread block, accessed through the attributes ``x``, ``y``, and ``z``. Each index is an integer spanning the range from 0 inclusive to the corresponding value of the attribute in :attr:`numba.cuda.blockDim` exclusive. zztid.xztid.yztid.zN) rrrrrrMacroSREG_SIGNATURExyzr r r rrs   rc@s8eZdZdZdZedeZedeZ edeZ dS)blockIdxa The block indices in the grid of thread blocks, accessed through the attributes ``x``, ``y``, and ``z``. Each index is an integer spanning the range from 0 inclusive to the corresponding value of the attribute in :attr:`numba.cuda.gridDim` exclusive. zzctaid.xzctaid.yzctaid.zN) rrrrrrrrrrrr r r rr-s   rc@s4eZdZdZedeZedeZedeZ dS)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, even if they belong to different blocks (i.e. each block is "full"). zntid.xzntid.yzntid.zN) rrrrrrrrrrr r r rr ;s  r c@s8eZdZdZdZedeZedeZ edeZ dS)gridDimzh The shape of the grid of blocks, accressed through the attributes ``x``, ``y``, and ``z``. zznctaid.xznctaid.yznctaid.zN) rrrrrrrrrrrr r r rr!Fs   r!warpsizelaneidcCsdS)Nr r r r r _ptx_grid1dWsr$cCsdS)Nr r r r r _ptx_grid2dZsr%cCsp|dkrd}tj}n@|dkr0d}ttjd}n$|dkrLd}ttjd}ntdtj|t|tj|gdS) 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. r z ptx.grid.1dz ptx.grid.2dz ptx.grid.3dzargument can only be 1, 2, 3)args) rint32UniTuple ValueErrorr Intrinsicr signatureintp)ndimfnamerestyper r r grid_expand]sr2zptx.gridT)callablecCsp|dkrd}tj}n@|dkr0d}ttjd}n$|dkrLd}ttjd}ntdtj|t|tj|gdS) ax 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. 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. r zptx.gridsize.1dr&zptx.gridsize.2dr'zptx.gridsize.3dzargument can only be 1, 2 or 3)r() rr)r*r+rr,rr-r.)r/r0r1r r rgridsize_expands r4z ptx.gridsizec@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. zN)rrrrrr r r rr5sr5c@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. zN)rrrrrr r r rr6sr6c@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. zN)rrrrrr r r rr7sr7c@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. zN)rrrrrr r r rr8sr8c@seZdZdZdZdS)syncwarpzP syncwarp(mask) Synchronizes a masked subset of threads in a warp. z N)rrrrrr r r rr9sr9c@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 N)rrrrrr r r rr: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 N)rrrrrr r r rr;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. zN)rrrrrr r r rr<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. zN)rrrrrr r r rr=s r=c@seZdZdZdZdS)threadfence_blockz. A memory fence at thread block level zN)rrrrrr r r rr>sr>c@seZdZdZdZdS)threadfence_systemz8 A memory fence at system level: across devices zN)rrrrrr r r rr? sr?c@seZdZdZdZdS) threadfencez( A memory fence at device level zN)rrrrrr r r rr@sr@cCs4t|tr|St|tr|fStdt|dS)Nzinvalid type for shape; got {0}) isinstancetupleint TypeErrorformattype)shaper r r_legalize_shapes   rHcCsPt|}t|}d}t||d}t|ttj|tj}t j ||||fdS)Nzptx.smem.allocC)r() rHlenrArrayrr-r*r.Anyrr,)rGdtyper/r0r1sigr r r shared_array&s rOc@s*eZdZdZdZejdedddgdZdS) sharedz" Shared memory namespace. zz shared.arrayTrGrM)r3argnamesN) rrrrrrrrOarrayr r r rrP/s   rPcCsPt|}t|}d}t||d}t|ttj|tj}t j ||||fdS)Nzptx.lmem.allocrI)r() rHrJrrKrr-r*r.rLrr,)rGrMr/r0r1rNr r r local_arrayFs rSc@s*eZdZdZdZejdedddgdZdS) localz! Local memory namespace. zz local.arrayTrGrM)r3rQN) rrrrrrrrSrRr r r rrTOs   rTcCs:d}ddlm}|j|}t||}tj|||gdS)Nzptx.cmem.aryliker )CUDATargetDesc)r()Z descriptorrUZ typingctxZresolve_argument_typerr-rr,)Zndarrayr0rUZarytyrNr r rconst_array_likebs    rVc@s(eZdZdZdZejdeddgdZdS)constz$ Constant memory namespace. zzconst.array_likeTZary)r3rQN) rrrrrrrrVZ array_liker r r rrWls  rWc@seZdZdZdS)popczK popc(val) Returns the number of set bits in the given value. N)rrrrr r r rrX|srXc@seZdZdZdS)brevzs brev(val) Reverse the bitpattern of an integer value; for example 0b10110110 becomes 0b01101101. N)rrrrr r r rrYsrYc@seZdZdZdS)clzzF clz(val) Counts the number of leading zeros in a value. N)rrrrr r r rrZsrZc@seZdZdZdS)ffsz^ ffs(val) Find the position of the least significant bit set to 1 in an integer. N)rrrrr r r rr[sr[c@seZdZdZdS)selpzp selp(a, b, c) Select between source operands, based on the value of the predicate source operand. N)rrrrr r r rr\sr\c@seZdZdZdS)fmazE fma(a, b, c) Perform the fused multiply-add operation. N)rrrrr r r rr]sr]c@sTeZdZdZdZGdddeZGdddeZGdddeZGd d d eZ d S) 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. N)rrrrr r r raddsr_c@seZdZdZdS)z atomic.maxamax(ary, idx, val) Perform atomic ary[idx] = max(ary[idx], val). NaN is treated as a missing value, so max(NaN, n) == max(n, NaN) == n. Note that this differs from Python and Numpy behaviour, where max(a, b) is always a when either a or b is a NaN. Supported on int32, int64, uint32, uint64, float32, float64 operands only. Returns the old value at the index location as if it is loaded atomically. N)rrrrr r r rmaxs r`c@seZdZdZdS)z atomic.minavmin(ary, idx, val) Perform atomic ary[idx] = min(ary[idx], val). NaN is treated as a missing value, so min(NaN, n) == min(n, NaN) == n. Note that this differs from Python and Numpy behaviour, where min(a, b) is always a when either a or b is a NaN. Supported on int32, int64, uint32, uint64, float32, float64 operands only. N)rrrrr r r rmins rac@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. N)rrrrr r r rcompare_and_swapsrbN) rrrrrr r_r`rarbr r r rr^s   r^)>rZ __future__rrroperatorZnumpyZllvmlite.llvmpy.coreZllvmpyZcoreZlcZnumbarrrrZcudadrvr objectr r-r)rrrr r!rr"r#r$r%r2Zgridr4Zgridsizer5r6r7r8r9r:r;r<r=r>r?r@rHrOrPrSrTrVrWrXrYrZr[r\r]r^r r r rs\