     uint32_t value[L];
   };*/
 
-  __device__
-  inline uint32_t processHistoElement(uint32_t& key, uint32_t offset, const uint4 element)
+  __device__ __forceinline__ uint32_t processHistoElement(uint32_t& key, uint32_t offset, const uint4 element)
   {
     assert(element.x <= element.y);
     assert(element.y <= element.z);

  __launch_bounds__(128)
  void extract(uint32_t* __restrict__ out_d,
               const uint4* __restrict__ hp_d,
               OffsetBlob offsetBlob)
  {

    __shared__ uint32_t offsets[12];

    return offset;
  }

  __device__ __forceinline__ uint32_t processDataElement(uint32_t& key, uint32_t offset, const uint4 element)
  {
    if (element.x <= key) {
      key -= element.x;
