U hd@sddlmZddlmZmZddlmZddlmZddl m Z ddl m Z ddl mZddlmZd d Zed d Zed dZeddZe eedddddZeddZddZeddZeddZeddZd S)!)ir)cudatypes)cgutils)RequireLiteralValue) signature)overload_attribute) nvvmutils) intrinsiccCsB|j}|dkrtj}n |dkr.ttj|}ntdt|tjS)N)zargument can only be 1, 2, 3)Z literal_valuerint32UniTuple ValueErrorr)ndimvalrestyperN/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/intrinsics.py_type_grid_functionsrcCs,t|tjst|t|}dd}||fS)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. cSsJ|j}|tjkrtj|ddSt|tjrFtj||jd}t ||SdS)Nr )dim) return_typerrr Z get_global_id isinstancercountr pack_array)contextbuildersigargsrZidsrrrcodegen1s   zgrid..codegenrrZIntegerLiteralrr typingctxrrr rrrgrids  r$cs8t|tjst|t|}ddfdd}||fS)agridsize(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. cSs0t|d|}t|d|}|||S)Nzntid.znctaid.)r call_sregmul)rrZntidZnctaidrrr_nthreads_for_dimRsz#gridsize.._nthreads_for_dimcsx|j}|d}|tjkr|St|tjrt|d}|jdkrNt|||fS|jdkrt|d}t||||fSdS)Nxyr r z)rrrrrrrr)rrrrrnxnyZnzr'rrr Ws       zgridsize..codegenr!r"rr-rgridsize<s   r.cCsttj}dd}||fS)NcSs t|dS)Nwarpsize)r r%)rrrrrrrr msz_warpsize..codegen)rrrr#rr rrr _warpsizeis r1r/r)targetcCs dd}|S)z_ The size of a warp. All architectures implemented to date have a warp size of 32. cSstS)N)r1)modrrrgetyszcuda_warpsize..getr)r3r4rrr cuda_warpsizessr5cCsttj}dd}||fS)a 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. cSs<d}|j}ttd}t|||}||d|S)Nzllvm.nvvm.barrier0r)moduler FunctionTypeZVoidTyperget_or_insert_functioncallZget_dummy_value)rrrrfnameZlmodfntysyncrrrr s  zsyncthreads..codegen)rrnoner0rrr syncthreadss r>cs2t|tjsdSttjtj}fdd}||fS)Ncs6ttdtdf}t|j|}|||S)N )rr7ZIntTyperr8r6r9)rrrrr;r<r:rrr sz'_syncthreads_predicate..codegen)rrIntegerri4)r# predicater:rr rr@r_syncthreads_predicates   rDcCsd}t|||S)z syncthreads_count(predicate) An extension to numba.cuda.syncthreads where the return value is a count of the threads where predicate is true. zllvm.nvvm.barrier0.popcrDr#rCr:rrrsyncthreads_countsrGcCsd}t|||S)z syncthreads_and(predicate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for all threads or 0 otherwise. zllvm.nvvm.barrier0.andrErFrrrsyncthreads_andsrHcCsd}t|||S)z syncthreads_or(predicate) An extension to numba.cuda.syncthreads where 1 is returned if predicate is true for any thread or 0 otherwise. zllvm.nvvm.barrier0.orrErFrrrsyncthreads_orsrIN)ZllvmliterZnumbarrZ numba.corerZnumba.core.errorsrZnumba.core.typingrZnumba.core.extendingrZ numba.cudar Znumba.cuda.extendingr rr$r.r1Moduler5r>rDrGrHrIrrrrs2          ,