ܙdZddlmZddlmZmZddlmZddlmZddl m Z ddl m Z ddl mZddlmZd Zed Zed Zed Ze ejed ddZedZdZedZedZedZdS))ir)cudatypes)cgutils)RequireLiteralValue) signature)overload_attribute) nvvmutils) intrinsicc|j}|dkr tj}n3|dvr tjtj|}nt dt |tjS)N)zargument can only be 1, 2, 3) literal_valuerint32UniTuple ValueErrorr)ndimvalrestypes 5lib/python3.11/site-packages/numba/cuda/intrinsics.py_type_grid_functionrs[  C axx+ .c227888 Wek * **ct|tjst|t |}d}||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. c|j}|tjkrtj|dSt |tjr0tj||j}tj ||SdS)Nr )dim) return_typerrr get_global_id isinstancercountr pack_array)contextbuildersigargsridss rcodegenzgrid..codegen1st/ ek ! !*7::: :  0 0 4)'w}EEEC%gs33 3 4 4rrrIntegerLiteralrr) typingctxrr$r's rgridr+sO" dE0 1 1(!$''' d # #C444 <rct|tjst|t |}dfd}||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. ctj|d|}tj|d|}|||S)Nzntid.znctaid.)r call_sregmul)r#rntidnctaids r_nthreads_for_dimz#gridsize.._nthreads_for_dimRsF"7MCMM::$Wooo>>{{4(((rcD|j}|d}|tjkr|St|tjr]|d}|jdkrt j|||fS|jdkr&|d}t j||||fSdSdS)Nxyrrz)rrrrrr rr!) r"r#r$r%rnxnynzr2s rr'zgridsize..codegenWs/  w , , ek ! !I  0 0 A""7C00B}!!)'B8<<<!##&&w44)'BB<@@@ A A $#rr()r*rr$r'r2s @rgridsizer:<sn" dE0 1 1(!$''' d # #C))) A A A A A <rcBttj}d}||fS)Nc,tj|dS)Nwarpsize)r r.)r"r#r$r%s rr'z_warpsize..codegenms"7J777r)rrrr*r$r's r _warpsizer?is, EK C888 <rr=r)targetc d}|S)z_ The size of a warp. All architectures implemented to date have a warp size of 32. ctS)N)r?)mods rgetzcuda_warpsize..getys {{r)rCrDs r cuda_warpsizerFss  JrcBttj}d}||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. cd}|j}tjtjd}t j|||}||d|S)Nzllvm.nvvm.barrier0rE)moduler FunctionTypeVoidTyperget_or_insert_functioncallget_dummy_value)r"r#r$r%fnamelmodfntysyncs rr'zsyncthreads..codegens_$~r{}}b11-dD%@@ T2&&(((r)rrnoner>s r syncthreadsrTs. EJ  C))) <rct|tjsdSttjtj}fd}||fS)Nctjtjdtjdf}tj|j|}|||S)N )rrJIntTyperrLrIrM)r"r#r$r%rQrRrOs rr'z'_syncthreads_predicate..codegensQrz"~~ 2/@AA-gndEJJ||D$'''r)rrIntegerri4)r* predicaterOr$r's ` r_syncthreads_predicater\sT i / /t EHeh ' 'C((((( <rc(d}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.popcr\r*r[rOs rsyncthreads_countr`s &E !)Y > >>rc(d}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.andr^r_s rsyncthreads_andrbs %E !)Y > >>rc(d}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.orr^r_s rsyncthreads_orrds $E !)Y > >>rN)llvmliternumbarr numba.corernumba.core.errorsrnumba.core.typingrnumba.core.extendingr numba.cudar numba.cuda.extendingr rr+r:r?ModulerFrTr\r`rbrdrErrrns111111''''''333333 ****** + + +  @ )) )X  LEL&& 6BBBCB  (    ?? ? ?? ? ?? ???r