 namespace {
 
 
-  __global__ 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)
+  __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;

        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]);
        for (uint32_t i = L; 1 < i; i--) {
          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]);
      }

    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;
