         offset = processHistoElement(key, 5 * offset, ((const uint4*)hp_d)[2 + offset]);
         offset = processHistoElement(key, 5 * offset, ((const uint4*)hp_d)[7 + 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]))[offset]);
         }
         out_d[index] = processMaskElement(key, 32 * offset, hp_d[offsets[0] + offset]);
       }

namespace {


  __global__ __launch_bounds__(128) void reduceBase(uint32_t* __restrict__ hp_d,
                                                    uint32_t* __restrict__ sb_d,
                                                    const uint32_t n1,
                                                    const uint32_t* __restrict__ src,
                                                    const uint32_t n0)
  {
    const uint32_t offset0 = blockDim.x * blockIdx.x + threadIdx.x;
    const uint32_t lane = threadIdx.x %  HP5_WARP_SIZE;

                                                  N);*/
  }
  else {
    ::reduceBase<<<(levels[0] + 3)/4, 4*32, 0, stream>>>(scratch_d + offsets[0],
                                                         scratch_d + offsets[L + 1],
                                                         levels[0],
                                                         in_d,
                                                         N);

    //::reduce1<true><<<levels[0], 160, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d) + offsets[0],
    //                                               scratch_d + 4 * offsets[L + 1],
    //                                               in_d,
    //                                               N);
    for (size_t i = 1; i < levels.size(); i++) {
      ::reduce1<<<(levels[i] + 31)/32, 160, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d + offsets[i]),
                                                         scratch_d + offsets[L + 1 + ((i + 0) & 1)],
                                                         levels[i],
                                                         scratch_d + offsets[L + 1 + ((i + 1) & 1)],
                                                         levels[i - 1]);

    }
    ::reduceApex<false, true><<<1, 128, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d),
