parameters: (func, argnames, extras)
return_statement: return kernel
docstring: Make a kernel that does a stride loop over the input chunks.
  
      Each block is responsible for a chunk in each iteration.
      Several iteration may be needed to handling a large number of chunks.
  
      The user function *func* will have all threads in the block for its
      computation.
  
      The resulting kernel can be used with any 1D grid size and 1D block size.
docstring_summary: Make a kernel that does a stride loop over the input chunks.
docstring_tokens: Make a kernel that does a stride loop over the input chunks .
function: def _make_chunk_wise_kernel(func, argnames, extras):
      """
      Make a kernel that does a stride loop over the input chunks.
  
      Each block is responsible for a chunk in each iteration.
      Several iteration may be needed to handling a large number of chunks.
  
      The user function *func* will have all threads in the block for its
      computation.
  
      The resulting kernel can be used with any 1D grid size and 1D block size.
      """
  
      # Build kernel source
      argnames = list(map(_mangle_user, argnames))
      extras = list(map(_mangle_user, extras))
      source = """
  def chunk_wise_kernel(nrows, chunks, {args}):
  {body}
  """
  
      args = ", ".join(argnames)
      body = []
  
      body.append("blkid = cuda.blockIdx.x")
      body.append("nblkid = cuda.gridDim.x")
      body.append("tid = cuda.threadIdx.x")
      body.append("ntid = cuda.blockDim.x")
  
      # Stride loop over the block
      body.append("for curblk in range(blkid, chunks.size, nblkid):")
      indent = " " * 4
  
      body.append(indent + "start = chunks[curblk]")
      body.append(
          indent
          + "stop = chunks[curblk + 1]"
          + " if curblk + 1 < chunks.size else nrows"
      )
  
      slicedargs = {}
      for a in argnames:
          if a not in extras:
              slicedargs[a] = f"{a}[start:stop]"
          else:
              slicedargs[a] = str(a)
      body.append(
          "{}inner({})".format(
              indent, ", ".join(slicedargs[k] for k in argnames)
          )
      )
  
      indented = ["{}{}".format(" " * 4, ln) for ln in body]
      # Finalize source
      concrete = source.format(args=args, body="\n".join(indented))
      # Get bytecode
      glbs = {"inner": cuda.jit(device=True)(func), "cuda": cuda}
      exec(concrete, glbs)
      # Compile as CUDA kernel
      kernel = cuda.jit(glbs["chunk_wise_kernel"])
      return kernel
function_tokens: def _make_chunk_wise_kernel ( func , argnames , extras ) : # Build kernel source argnames = list ( map ( _mangle_user , argnames ) ) extras = list ( map ( _mangle_user , extras ) ) source = """
  def chunk_wise_kernel(nrows, chunks, {args}):
  {body}
  """ args = ", " . join ( argnames ) body = [ ] body . append ( "blkid = cuda.blockIdx.x" ) body . append ( "nblkid = cuda.gridDim.x" ) body . append ( "tid = cuda.threadIdx.x" ) body . append ( "ntid = cuda.blockDim.x" ) # Stride loop over the block body . append ( "for curblk in range(blkid, chunks.size, nblkid):" ) indent = " " * 4 body . append ( indent + "start = chunks[curblk]" ) body . append ( indent + "stop = chunks[curblk + 1]" + " if curblk + 1 < chunks.size else nrows" ) slicedargs = { } for a in argnames : if a not in extras : slicedargs [ a ] = f"{a}[start:stop]" else : slicedargs [ a ] = str ( a ) body . append ( "{}inner({})" . format ( indent , ", " . join ( slicedargs [ k ] for k in argnames ) ) ) indented = [ "{}{}" . format ( " " * 4 , ln ) for ln in body ] # Finalize source concrete = source . format ( args = args , body = "\n" . join ( indented ) ) # Get bytecode glbs = { "inner" : cuda . jit ( device = True ) ( func ) , "cuda" : cuda } exec ( concrete , glbs ) # Compile as CUDA kernel kernel = cuda . jit ( glbs [ "chunk_wise_kernel" ] ) return kernel