 
 __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_output[index] = 0;
-    } else {
-        device_output[index] = (device_input[index] == device_input[index+1]);
-    }
-}
-
-__global__ void find_repeats_kernel(
-    int *) {
-
+    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) {

            twod1);
        cudaThreadSynchronize();
    }

    cudaMemset(device_result + length - 1, 0, sizeof(int));

    for (int twod = length >> 1; twod >= 1; twod >>= 1) {
        int twod1 = twod << 1;
        int partitions = length / twod1;

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