             twod1);
         cudaThreadSynchronize();
     }
-
     cudaMemset(device_result + length - 1, 0, sizeof(int));
 
+    // downsweep phase    
     for (int twod = length >> 1; twod >= 1; twod >>= 1) {
         int twod1 = twod << 1;
         int partitions = length / twod1;

     * power of 2 larger than the input.
     */
    
    for (int twod = 1; twod < (length >> 1); twod <<= 1) {
        int twod1 = twod << 1;
        int partitions = length / twod1;


__global__ void check_repeats_kernel(
    int *device_input,
    int length) {

    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= length - 1) {
        device_input[index] = 0;
    } else {
        device_input[index] = (device_input[index] == device_input[index+1]);
    }
}

__global__ void find_repeats_kernel(
    int *device_input,
    int length,
    int *device_output) {

    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < length - 1 &&
        device_input[index] != device_input[index + 1]) {
        device_output[device_input[index]] = index;
    }
}

int find_repeats(int *device_input, int length, int *device_output) {
