                 }
             } else {
                 // MODE == Rollback
+                uint32 insert_idx_M = table.find_insertion_point(keys.M);
+                assert(insert_idx_M < LOCAL_TABLE_SIZE);
                 rollback(insert_idx_M, keys.M, fractional_error_M);
+
+                uint32 insert_idx_I = table.find_insertion_point(keys.I);
+                assert(insert_idx_I < LOCAL_TABLE_SIZE);
                 rollback(insert_idx_I, keys.I, fractional_error_I);
+
+                uint32 insert_idx_D = table.find_insertion_point(keys.D);
+                assert(insert_idx_D < LOCAL_TABLE_SIZE);
                 rollback(insert_idx_D, keys.D, fractional_error_D);
             }
         }

    // returns false if the table is full
    CUDA_HOST_DEVICE bool update(uint32 insert_idx, covariate_key key, float fractional_error)
    {
        if (!table.exists(insert_idx) || table.keys[insert_idx] != key)
        {
            //            printf("adding key %x (RG = %d)\n", key, covariate_chain::decode(key, 1));


                bool ret;

                uint32 insert_idx_M = table.find_insertion_point(keys.M);
                assert(insert_idx_M < LOCAL_TABLE_SIZE);

                ret = update(insert_idx_M, keys.M, fractional_error_M);
                if (ret == false)
                {
//                    printf("[B %d T %d] overflow! rollback at read index %u event index %u event M insert index %d\n", blockIdx.x, threadIdx.x, read_index, cigar_event_index, insert_idx_M);

                    // mark for rollback from here
                    last_cigar_event_index = cigar_event_index;
                    return false;
                }

                uint32 insert_idx_I = table.find_insertion_point(keys.I);
                assert(insert_idx_I < LOCAL_TABLE_SIZE);

                ret = update(insert_idx_I, keys.I, fractional_error_I);
                if (ret == false)
                {
//                    printf("[B %d T %d] overflow! rollback at read index %u event index %u event I insert index %d\n", blockIdx.x, threadIdx.x, read_index, cigar_event_index, insert_idx_I);

//                    printf("[B %d T %d] rolling back M (index %d)\n", blockIdx.x, threadIdx.x, insert_idx_M);
                    rollback(insert_idx_M, keys.M, fractional_error_M);

                    // mark for rollback from here
