   {
      // FIXME: no point in up-traversal if cell doesn't change.
      // FIXME: If extract 3 in a go, organize traversals to reduce up-traversal.
+     // FIXME: Try to let a thread do all three indices of a triangle with shared down-traversal
+     // FIXME: Try to recylce up-traversals, they might often be indentical (maybe check the table first to see if this is true)
 
 
     uint32_t ix = blockDim.x * blockIdx.x + threadIdx.x;

    uint3 cell = make_uint3(31 * chunk.x + thread,
                            5 * chunk.y,
                            5 * chunk.z);

    // TODO check if it is faster to only process 31 cells and don't rely on cache for fetching field.

    // One warp processes a 32 * 5 * 5 chunk, outputs 800 mc cases to base level and 800/5=160 sums to sideband.
    const FieldType* ptr = field_d

                            5 * chunk.y,
                            5 * chunk.z);

    // One warp processes a 31 * 5 * 5 chunk, outputs 800 mc cases to base level and 800/5=160 sums to sideband.
    const FieldType* ptr = field_d
                         + cell.z * field_slice_stride
                         + cell.y * field_row_stride
