     return a;
   }
 
+  // Count the number of edge-intersections of edges belonging to that cell (0-3).
+  // We just count the number of set bits in a bitmask.
   __device__ __forceinline__ uint32_t axisCountFromCase(const uint32_t c)
   {
     uint32_t a = piercingAxesFromCase(c);

    const uint32_t thread = threadIdx.x % 32;
    const uint32_t chunk_ix = blockIdx.x + warp;

    uint3 chunk = make_uint3(chunk_ix % chunks.x,
                             (chunk_ix / chunks.x) % chunks.y,
                             (chunk_ix / chunks.x) / chunks.y);

    return n;
  }

  // Fetch samples from scalar-field, deduce MC cases for cells, and write out
  // first level ot the HistoPyramid and sideband.
  //
  // See comments for reduceBase
  //
  // In addition to calculating the MC cell case, it also calculates the number
  // of unique vertices belong to that cell (0-3). This grid is one cell-layer
  // larger along X,Y and Z than the unindexed grid to have the simple relation-
  // ship of each cell having 0-3 vertices.
  //
  // But this requires careful masking to avoid producing extra vertices, we
  // basically implement GL_CLAMP-behavior.
  template<class FieldType>
  __global__ __launch_bounds__(32) void reduceBaseIndexed(uint8_t* __restrict__           index_cases_d,
                                                          uint4* __restrict__             out_vertex_level0_d,
