     const uint32_t thread = threadIdx.x % 32;
     const uint32_t chunk_ix = blockIdx.x + warp;
 
+    // Figure out which chunk we are in.
+    // FIXME: Try to use grid extents to mach chunk grid to avoid all these modulo/divisions.
     uint3 chunk = make_uint3(chunk_ix % chunks.x,
                              (chunk_ix / chunks.x) % chunks.y,
                              (chunk_ix / chunks.x) / chunks.y);

    float t4 = ((ptr < end) && (ptr[0] < t) ? 1.f : 0.f); ptr += stride;
    float t5 = ((ptr < end) && (ptr[0] < t) ? 1.f : 0.f); ptr += stride;

    float r0 = __fmaf_rn(float(1 << 16), t2, __fmaf_rn(float(1 << 8), t1, t0));
    float r1 = __fmaf_rn(float(1 << 16), t5, __fmaf_rn(float(1 << 8), t4, t3));
    return make_float2(r0, r1);
  }

  template<typename T>
  __device__ float2 fetch2(const T* ptr, const size_t stride, const uint32_t y, const uint32_t max_y, const float t)
  {

  {
    const uint32_t offset0 = 5 * 32 * blockIdx.x + threadIdx.x;

    // FIXME: Test iddea, each warp reads 32 values. read instead 32/4 uint4's.
    __shared__ uint32_t sb[5 * 32];
    sb[threadIdx.x] = offset0 < n0 ? sb0_d[offset0] : 0;
    __syncthreads();
