     d_output[ix] = min(v1,v2);
     d_output[ix + stride] = max(v1,v2);
   
+}
+__global__ void cSwaps(int *d_input,
+                       int *d_output,
+		       uint stride) 
+{
+    uint global_comparatorI = blockIdx.x * blockDim.x + threadIdx.x;
+   
+    uint pos = 2 * global_comparatorI - (global_comparatorI & (stride - 1));
+
+    int keyA = d_input[pos +      0];
+    int keyB = d_input[pos + stride];
+
+    int newA = min(keyA,keyB);
+    int newB = max(keyA,keyB);
+
+    d_input[pos +      0] = newA;   
+    d_input[pos + stride] = newB;
+
 }
 
 //* Assume input and output arrays have length 2^lenLog 

    } 
  }

  printf("\n%s",passed ? "Passed!" : "Failed!");

  return 0;

  tsortSmall<<<blocks, threads,4096>>>(d_data, d_data);
  // cudaThreadSynchronize();
  
  for(int i = 0 ; i < diff ; i += 1){ 
  
    for(int j = i; j >= 0; j -= 1){ 
      cSwaps<<<blocks,threads/2,0>>>(d_data, d_data,(1<<j)*SMALL_SIZE);
      
      //  cudaThreadSynchronize();
    }
