ó \ K]c@`s‚dZddlmZmZmZddlZddlZddljj Z ddl m Z m Z mZmZddlmZdefd„ƒYZeje jƒZd efd „ƒYZd efd „ƒYZd efd„ƒYZdefd„ƒYZejdeƒZejdeƒZd„Zd„Zd„Z ejde de!ƒZ"d„Z#ejde#de!ƒZ$defd„ƒYZ%defd„ƒYZ&defd„ƒYZ'd efd!„ƒYZ(d"efd#„ƒYZ)d$efd%„ƒYZ*d&efd'„ƒYZ+d(efd)„ƒYZ,d*efd+„ƒYZ-d,efd-„ƒYZ.d.efd/„ƒYZ/d0efd1„ƒYZ0d2„Z1d3„Z2d4efd5„ƒYZ3d6„Z4d7efd8„ƒYZ5d9„Z6d:efd;„ƒYZ7d<efd=„ƒYZ8d>efd?„ƒYZ9d@efdA„ƒYZ:dBefdC„ƒYZ;dDefdE„ƒYZ<dFefdG„ƒYZ=dHefdI„ƒYZ>dS(Js1 This scripts specifies all PTX special objects. i(tprint_functiontabsolute_importtdivisionN(ttypestirttypingtmacroi(tnvvmtStubcB`s,eZdZdZdZd„Zd„ZRS(slA stub object to represent special objects which is meaningless outside the context of CUDA-python. scC`std|ƒ‚dS(Ns%s is not instantiable(tNotImplementedError(tcls((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt__new__scC`s|jS(N(t _description_(tself((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt__repr__s((t__name__t __module__t__doc__R t __slots__R R(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR s  t threadIdxcB`sJeZdZdZejdeƒZejdeƒZejdeƒZ RS(s 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. sstid.xstid.ystid.z( RRRR RtMacrotSREG_SIGNATUREtxtytz(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRs tblockIdxcB`sJeZdZdZejdeƒZejdeƒZejdeƒZ RS(s 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. ssctaid.xsctaid.ysctaid.z( RRRR RRRRRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR-s tblockDimcB`sDeZdZejdeƒZejdeƒZejdeƒZRS(sÝ 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"). sntid.xsntid.ysntid.z( RRRRRRRRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR;stgridDimcB`sJeZdZdZejdeƒZejdeƒZejdeƒZ RS(sh The shape of the grid of blocks, accressed through the attributes ``x``, ``y``, and ``z``. ssnctaid.xsnctaid.ysnctaid.z( RRRR RRRRRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRFs twarpsizetlaneidcC`sdS(N((((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt _ptx_grid1dWtcC`sdS(N((((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt _ptx_grid2dZRcC`s¦|dkrd}tj}n`|dkrHd}tjtjdƒ}n6|dkrrd}tjtjdƒ}n tdƒ‚tj|tj|tjƒd|gƒS( sgrid(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. is ptx.grid.1dis ptx.grid.2dis ptx.grid.3dsargument can only be 1, 2, 3targs( Rtint32tUniTuplet ValueErrorRt IntrinsicRt signaturetintp(tndimtfnametrestype((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt grid_expand]s     sptx.gridtcallablecC`s¦|dkrd}tj}n`|dkrHd}tjtjdƒ}n6|dkrrd}tjtjdƒ}n tdƒ‚tj|tj|tjƒd|gƒS( sx 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. isptx.gridsize.1disptx.gridsize.2disptx.gridsize.3dsargument can only be 1, 2 or 3R!( RR"R#R$RR%RR&R'(R(R)R*((s/lib/python2.7/site-packages/numba/cuda/stubs.pytgridsize_expands     s ptx.gridsizet syncthreadscB`seZdZdZRS(s 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. s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR.£stsyncthreads_countcB`seZdZdZRS(s¡ syncthreads_count(predictate) An extension to numba.cuda.syncthreads where the return value is a count of the threads where predicate is true. s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR/­stsyncthreads_andcB`seZdZdZRS(sž syncthreads_and(predictate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for all threads or 0 otherwise. s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR0·stsyncthreads_orcB`seZdZdZRS(sœ syncthreads_or(predictate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for any thread or 0 otherwise. s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR1ÁstsyncwarpcB`seZdZdZRS(sP syncwarp(mask) Synchronizes a masked subset of threads in a warp. s (RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR2Îstshfl_sync_intrinsiccB`seZdZdZRS(sÊ 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 s (RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR3×stvote_sync_intrinsiccB`seZdZdZRS(sÊ 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 s (RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR4ástmatch_any_synccB`seZdZdZRS(sÒ 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. s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR5ëstmatch_all_synccB`seZdZdZRS(sŽ 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. s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR6õs tthreadfence_blockcB`seZdZdZRS(s. A memory fence at thread block level s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR7stthreadfence_systemcB`seZdZdZRS(s8 A memory fence at system level: across devices s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR8 st threadfencecB`seZdZdZRS(s( A memory fence at device level s(RRRR (((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR9scC`sHt|tƒr|St|tƒr)|fStdjt|ƒƒƒ‚dS(Nsinvalid type for shape; got {0}(t isinstancettupletintt TypeErrortformatttype(tshape((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt_legalize_shapes cC`svt|ƒ}t|ƒ}d}tj||dƒ}tj|tjtj|ƒtjƒ}t j ||d||fƒS(Nsptx.smem.alloctCR!( RAtlenRtArrayRR&R#R'tAnyRR%(R@tdtypeR(R)R*tsig((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt shared_array&s   'tsharedcB`s8eZdZdZejdededddgƒZRS(s" Shared memory namespace. ss shared.arrayR,targnamesR@RF( RRRR RRRHtTruetarray(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRI/s  cC`svt|ƒ}t|ƒ}d}tj||dƒ}tj|tjtj|ƒtjƒ}t j ||d||fƒS(Nsptx.lmem.allocRBR!( RARCRRDRR&R#R'RERR%(R@RFR(R)R*RG((s/lib/python2.7/site-packages/numba/cuda/stubs.pyt local_arrayFs   'tlocalcB`s8eZdZdZejdededddgƒZRS(s! Local memory namespace. ss local.arrayR,RJR@RF( RRRR RRRMRKRL(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRNOs cC`sSd}ddlm}|jj|ƒ}tj||ƒ}tj||d|gƒS(Nsptx.cmem.arylikei(tCUDATargetDescR!(t descriptorROt typingctxtresolve_argument_typeRR&RR%(tndarrayR)ROtarytyRG((s/lib/python2.7/site-packages/numba/cuda/stubs.pytconst_array_likebs tconstcB`s5eZdZdZejdededdgƒZRS(s$ Constant memory namespace. ssconst.array_likeR,RJtary( RRRR RRRURKt array_like(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRVls tpopccB`seZdZRS(sK popc(val) Returns the number of set bits in the given value. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRY|stbrevcB`seZdZRS(ss brev(val) Reverse the bitpattern of an integer value; for example 0b10110110 becomes 0b01101101. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRZ„stclzcB`seZdZRS(sF clz(val) Counts the number of leading zeros in a value. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR[stffscB`seZdZRS(s^ ffs(val) Find the position of the least significant bit set to 1 in an integer. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR\•stselpcB`seZdZRS(sp selp(a, b, c) Select between source operands, based on the value of the predicate source operand. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR]ŸstfmacB`seZdZRS(sE fma(a, b, c) Perform the fused multiply-add operation. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR^©statomiccB`sleZdZdZdefd„ƒYZdefd„ƒYZdefd„ƒYZdefd „ƒYZRS( s$Namespace for atomic operations staddcB`seZdZRS(sàadd(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. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR`¸stmaxcB`seZdZRS(sÒmax(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. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRaÂs tmincB`seZdZRS(svmin(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. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRbÐs tcompare_and_swapcB`seZdZRS(sêcompare_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. (RRR(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyRcÛs( RRRR RR`RaRbRc(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyR_³s   (?Rt __future__RRRtoperatortnumpytllvmlite.llvmpy.coretllvmpytcoretlctnumbaRRRRtcudadrvRtobjectRR&R"RRRRRRRRRR R+RKtgridR-tgridsizeR.R/R0R1R2R3R4R5R6R7R8R9RARHRIRMRNRURVRYRZR[R\R]R^R_(((s/lib/python2.7/site-packages/numba/cuda/stubs.pyts\  "