code: def inline_reduce(N, buf, pos, count, manner_fn):
      """
      Return C++ code for a function that reduces a contiguous buffer.
  
      Parameters
      ----------
      N
          Length of the buffer.
      buf
          buffer pointer.
      pos
          Index of executing thread.
      count
          Number of executing threads.
      manner_fn
          A function that accepts strings of arguments a and b, and
          returns c code for their reduction.
  
            return "%(a)s + %(b)s"
  
          for a sum reduction.
  
      Notes
      -----
      `buf` should be in gpu shared memory, we access it many times.
  
      This function leaves the answer in position 0 of the buffer. The
      rest of the buffer is trashed by this function.
  
      """
      loop_line = manner_fn("%s[%s]" % (buf, pos), "%s[i]" % (buf))
      r_n = manner_fn("%s[%s]" % (buf, pos), "%s[%s+_n]" % (buf, pos))
  
      return (
          """
      {
          // This function trashes buf[1..warpSize],
          // leaving the reduction result in buf[0].
  
          if (%(pos)s < warpSize) {
              for (int i = %(pos)s + warpSize; i < %(N)s; i += warpSize)
              {
                  %(buf)s[%(pos)s] = %(loop_line)s;
              }
          }
          __syncthreads();
          //reduce so that %(pos)s 0 has the reduction of everything
          for (unsigned int _n = warpSize / 2; _n > 0; _n /= 2) {
            if (%(pos)s < _n && %(pos)s + _n < %(N)s)
              %(buf)s[%(pos)s] = %(r_n)s;
            __syncthreads();
          }
      }
      """
          % locals()
      )
  
file_path: Theano/theano/gpuarray/kernel_codegen.py
imports: from __future__ import absolute_import, print_function, division
label_desc: Use image node_id 3 for calling a global function with example usage: inline_reduce(N, buf, pos, count, manner_fn) and returns: unknown