         uint32_t offset = 0;
         uint32_t key = index;
         offset = processHistoElement(key, 5 * offset, T);
-        offset = processHistoElement(key, 5 * offset, ((const uint4*)hp_d)[2 + offset]);
-        offset = processHistoElement(key, 5 * offset, ((const uint4*)hp_d)[7 + offset]);
+        offset = processHistoElement(key, 5 * offset, *(const uint4*)(hp_d + 8 + 4 * offset));
+        offset = processHistoElement(key, 5 * offset, *(const uint4*)(hp_d + 28 + 4 * offset));
         for (uint32_t i = L; 1 < i; i--) {
-          offset = processDataElement(key, 5 * offset, ((const uint4*)(hp_d + offsets[i - 1]))[offset]);
+          offset = processDataElement(key, 5 * offset, *(const uint4*)(hp_d + offsets[i - 1] + 4 * offset));
         }
         out_d[index] = processMaskElement(key, 32 * offset, hp_d[offsets[0] + offset]);
       }

  auto L = levels.size();
  if (L == 0) {
    assert(false);
/*    reduceApex<true, true><<<1, 128, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d),
                                                  sum_d,
                                                  in_d,
                                                  N);*/
  }
  else {
    ::reduceBase<<<(levels[0] + 3)/4, 4*32, 0, stream>>>(scratch_d + offsets[0],

    if (11 < L) offsets[11] = offsetBlob.dataB;

    uint32_t N = hp_d[0];
    uint4 T = *(const uint4*)(hp_d + 4);

    for (uint32_t k = HP5_EXTRACT_BLOCKSIZE * blockIdx.x; k < N; k += HP5_EXTRACT_BLOCKS * HP5_EXTRACT_BLOCKSIZE) {
      const uint32_t index = k + threadIdx.x;
