   {
     const uint32_t offset0 = 5 * 32 * blockIdx.x + threadIdx.x;
 
-    // Idea, each warp reads 32 values. read instead 32/4 uint4's.
-
+    // 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();

    }
  }

  __global__ __launch_bounds__(128) void reduceApex(uint4* __restrict__ apex_d,
                                                    uint32_t* sum_d,
                                                    const uint32_t* in_d,


          vsum = vc_y0 + vc_y1 + vc_y2 + vc_y3 + vc_y4;
          isum = ic_y0 + ic_y1 + ic_y2 + ic_y3 + ic_y4;
          if (isum | vsum) {
            // 8-bit arithmetic should suffice here
            out_index_level0_d[32 * 5 * blockIdx.x + 32 * z + thread] = make_uint4(vc_y0,
