 //
 // For more details on the reduction algorithm (notably the multi-pass approach), see
 // the "reduction" sample in the CUDA SDK.
-template <unsigned int blockSize, bool nIsPow2>
+template <
+    unsigned int blockSize
+  , bool nIsPow2
+  , typename typ
+  , typename styp
+  , class task
+  >
 __device__  __inline__
-void reduceSinglePass_dev(const cuDoubleComplex *g_idata, double *g_odata, unsigned int n)
+void reduceSinglePass_devGen(const styp *g_idata, typ *g_odata, unsigned int n)
 {
 
     //
     // PHASE 1: Process all inputs assigned to this block
     //
 
-    reduceBlocks<blockSize, nIsPow2>(g_idata, g_odata, n);
+    reduceBlocksGen<blockSize, nIsPow2, typ, styp, task>(g_idata, g_odata, n);
 
     //
     // PHASE 2: Last block finished will process all partial sums

#ifndef _REDUCE_KERNEL_H_
#define _REDUCE_KERNEL_H_

#include <device_functions.h>
#include <cuComplex.h>

/*
    Parallel sum reduction using shared memory
    - takes log(n) steps for n input elements
    - uses n/2 threads
    - only works for power-of-2 arrays

    {
        const unsigned int tid = threadIdx.x;
        __shared__ bool amLast;
        extern typ __shared__ smem[];

        // wait until all outstanding memory instructions in this thread are finished
        __threadfence();
