code: def get_reduction_module(
      out_type,
      block_size,
      neutral,
      reduce_expr,
      map_expr,
      arguments,
      name="reduce_kernel",
      keep=False,
      options=None,
      preamble="",
  ):
      from pycuda.compiler import SourceModule
  
      src = """
          #include <pycuda-complex.hpp>
  
          #define BLOCK_SIZE %(block_size)d
          #define READ_AND_MAP(i) (%(map_expr)s)
          #define REDUCE(a, b) (%(reduce_expr)s)
  
          %(preamble)s
  
          typedef %(out_type)s out_type;
  
          extern "C"
          __global__
          void %(name)s(out_type *out, %(arguments)s,
            unsigned int seq_count, size_t n)
          {
            // Needs to be variable-size to prevent the braindead CUDA compiler from
            // running constructors on this array. Grrrr.
            extern __shared__ out_type sdata[];
  
            unsigned int tid = threadIdx.x;
  
            size_t i = blockIdx.x*BLOCK_SIZE*seq_count + tid;
  
            out_type acc = %(neutral)s;
            for (unsigned s = 0; s < seq_count; ++s)
            {
              if (i >= n)
                break;
              acc = REDUCE(acc, READ_AND_MAP(i));
  
              i += BLOCK_SIZE;
            }
  
            sdata[tid] = acc;
  
            __syncthreads();
  
            #if (BLOCK_SIZE >= 512)
              if (tid < 256) { sdata[tid] = REDUCE(sdata[tid], sdata[tid + 256]); }
              __syncthreads();
            #endif
  
            #if (BLOCK_SIZE >= 256)
              if (tid < 128) { sdata[tid] = REDUCE(sdata[tid], sdata[tid + 128]); }
              __syncthreads();
            #endif
  
            #if (BLOCK_SIZE >= 128)
              if (tid < 64) { sdata[tid] = REDUCE(sdata[tid], sdata[tid + 64]); }
              __syncthreads();
            #endif
  
            if (tid < 32)
            {
              // 'volatile' required according to Fermi compatibility guide 1.2.2
              volatile out_type *smem = sdata;
              if (BLOCK_SIZE >= 64) smem[tid] = REDUCE(smem[tid], smem[tid + 32]);
              if (BLOCK_SIZE >= 32) smem[tid] = REDUCE(smem[tid], smem[tid + 16]);
              if (BLOCK_SIZE >= 16) smem[tid] = REDUCE(smem[tid], smem[tid + 8]);
              if (BLOCK_SIZE >= 8)  smem[tid] = REDUCE(smem[tid], smem[tid + 4]);
              if (BLOCK_SIZE >= 4)  smem[tid] = REDUCE(smem[tid], smem[tid + 2]);
              if (BLOCK_SIZE >= 2)  smem[tid] = REDUCE(smem[tid], smem[tid + 1]);
            }
  
            if (tid == 0) out[blockIdx.x] = sdata[0];
          }
          """ % {
          "out_type": out_type,
          "arguments": arguments,
          "block_size": block_size,
          "neutral": neutral,
          "reduce_expr": reduce_expr,
          "map_expr": map_expr,
          "name": name,
          "preamble": preamble,
      }
      return SourceModule(
          src, options=options, keep=keep, no_extern_c=True
      )
  
file_path: pycuda/pycuda/reduction.py
imports: from pycuda.tools import context_dependent_memoize
  from pycuda.tools import dtype_to_ctype
  import numpy
label_desc: Use image node_id 1 for calling a global function with example usage: get_reduction_module(out_type, block_size, neutral, reduce_expr, map_expr, arguments, name, keep, options, preamble) and returns: SourceModule