#define MUTATION

#include "my_cutil.h"

#define NUM    32


////////////////////////////////////////////////////////////////////////////////
//    Notes for running in PUG:
//   Bitvector size: 16 bits
//   The ASSUME_NO_OVFLO flag can be turned off to obtain substantial speedups
////////////////////////////////////////////////////////////////////////////////


/*
__device__ inline void swap(int & a, int & b)
{
	// Alternative swap doesn't use a temporary register:
	// a ^= b;
	// b ^= a;
	// a ^= b;
	
    int tmp = a;
    a = b;
    b = tmp;
}
*/

__global__ void BitonicKernel(int * values)
{
  extern __shared__ int shared[];
  
  unsigned int tid = threadIdx.x;
  
  // Copy input to shared mem.
  shared[tid] = values[tid];
  
#ifdef MUTATION
  if (threadIdx.x == 0) {
#endif
  __syncthreads();
#ifdef MUTATION
   /* BUGINJECT: NON_UNIFORM_CONTROL_FLOW, UP */
  }
#endif
  
  // Parallel bitonic sort.
  for (unsigned int k = 2; k <= NUM; k *= 2)
    {
      // Bitonic merge:
      for (unsigned int j = k / 2; j>0; j /= 2)
        {
	  unsigned int ixj = tid ^ j;
          
	  if (ixj > tid)
            {
	      if ((tid & k) == 0)
                {
		  if (shared[tid] > shared[ixj])
                    {
		      unsigned int tmp = shared[tid];
		      shared[tid] = shared[ixj];
		      shared[ixj] = shared[tid];
                    }
                }
	      else
                {
		  if (shared[tid] < shared[ixj])
                    {
		      unsigned int tmp = shared[tid];
		      shared[tid] = shared[ixj];
		      shared[ixj] = shared[tid];
                    }
                }
            }
	  
	  __syncthreads();
        }
    }
  
  // Write result.
  values[tid] = shared[tid];
}

