
/*
    Parallel reduction kernels
*/

#ifndef _REDUCE_KERNEL_H_
#define _REDUCE_KERNEL_H_

#include "my_cutil.h"

#ifdef __DEVICE_EMULATION__
#define EMUSYNC __syncthreads()
#else
#define EMUSYNC
#endif

////////////////////////////////////////////////////////////////////////////////
//   Notes for running in PUG:
//   Bitvector size: 24 bits
////////////////////////////////////////////////////////////////////////////////

// Macros to append an SM version identifier to a function name
// This allows us to compile a file multiple times for different architecture
// versions
// The second macro is necessary to evaluate the value of the SMVERSION macro
// rather than appending "SMVERSION" itself
#define FUNCVERSION(x, y) x ## _ ## y
#define XFUNCVERSION(x, y) FUNCVERSION(x, y)
#define FUNC(NAME) XFUNCVERSION(NAME, SMVERSION) 

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

extern __shared__ int sdata[];

/* This version uses contiguous threads, but its interleaved 
   addressing results in many shared memory bank conflicts. */

__global__ void
FUNC(reduce1kernel)(int *g_idata, int *g_odata, unsigned int n)
{
  // SharedMemory<T> smem;
  //  T *sdata = smem.getPointer();

    // load shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    
    sdata[tid] = (i < n) ? g_idata[i] : 0;
    
    __syncthreads();

    // do reduction in shared mem
    for(unsigned int s=1; s < blockDim.x; s *= 2) 
    {
        int index = 2 * s * tid;
      //  int index;

        if (index < blockDim.x) 
        {
            sdata[index] += sdata[index + s];
        }
	 __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) {
#ifdef MUTATION
      __syncthreads();
#endif
      g_odata[blockIdx.x] = sdata[0];
    }
}

#endif // #ifndef _REDUCE_KERNEL_H_
