         cases_d[5 * (32 * 5 * blockIdx.x + 32 * y + thread) + 2] = cases.e2;
         cases_d[5 * (32 * 5 * blockIdx.x + 32 * y + thread) + 3] = cases.e3;
         cases_d[5 * (32 * 5 * blockIdx.x + 32 * y + thread) + 4] = cases.e4;
-        out_level0_d[32 * 5 * blockIdx.x + 32 * y + thread] = make_uint4(cases.e0, cases.e1, cases.e2, cases.e3);
+        out_level0_d[32 * 5 * blockIdx.x + 32 * y + thread] = make_uint4(counts.e0,
+                                                                         counts.e0 + counts.e1,
+                                                                         counts.e0 + counts.e1 + counts.e2,
+                                                                         counts.e0 + counts.e1 + counts.e2 + counts.e3);
       }
       cell.z++;
       out_sideband_d[32 * (5 * blockIdx.x + y) + threadIdx.x] = sum;// t.x + t.y + t.z + t.w + sideband[5 * thread + 4];

                              sb[5 * threadIdx.x + 1],
                              sb[5 * threadIdx.x + 2],
                              sb[5 * threadIdx.x + 3]);
        hp1_d[offset1] = hp;
        sb1_d[offset1] = hp.x + hp.y + hp.z + hp.w + sb[5 * threadIdx.x + 4];
      }
    }

  if (output_buffer) {
    ::runExtractionP<<<1, 32, 0, stream>>>(reinterpret_cast<float*>(output_buffer),
                                           static_cast<uint32_t>(output_buffer_size / (3 * sizeof(float))),
                                           ctx->pyramid,
                                           make_uint2(ctx->chunks.x, ctx->chunks.y),
                                           make_float3(1.f/ctx->cells.x,
                                                       1.f/ctx->cells.y,
                                                       1.f/ctx->cells.z));
  }

  assert(output_buffer);
