 
                     if (reference_bp != read_bp)
                     {
-                        ctx.cigar.is_snp[idx.read_start + current_bp_idx] = 1;
+                        snp_vector[idx.read_start + current_bp_idx] = 1;
 
                         // if we are inside the clipped read window, count this error
                         if (current_bp_idx >= read_window_clipped.x && current_bp_idx <= read_window_clipped.y)

    ctx.is_deletion.resize(batch.device.reads.size());
    ctx.num_errors.resize(batch.device.num_reads);

    // the following require zero initialization
    thrust::fill(ctx.is_snp.m_storage.begin(), ctx.is_snp.m_storage.end(), 0);
    thrust::fill(ctx.is_insertion.m_storage.begin(), ctx.is_insertion.m_storage.end(), 0);
    thrust::fill(ctx.is_deletion.m_storage.begin(), ctx.is_deletion.m_storage.end(), 0);
    thrust::fill(ctx.num_errors.begin(), ctx.num_errors.end(), 0);

    // cigar_events_read_index is initialized to -1; this means that all reads are considered inactive

template <target_system system>
struct compute_error_vectors : public lambda<system>
{
    LAMBDA_INHERIT_MEMBERS;

    typename vector<system, uint8>::view snp_vector;
    typename vector<system, uint8>::view ins_vector;
    typename vector<system, uint8>::view del_vector;

    compute_error_vectors(typename firepony_context<system>::view ctx,
                          const typename alignment_batch_device<system>::const_view batch,
                          typename vector<system, uint8>::view snp_vector,
                          typename vector<system, uint8>::view ins_vector,
                          typename vector<system, uint8>::view del_vector)
        : lambda<system>(ctx, batch),
          snp_vector(snp_vector),
          ins_vector(ins_vector),
          del_vector(del_vector)
    { }

    CUDA_HOST_DEVICE void operator() (const uint32 read_index)
    {
