                             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.
+    // 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

    if (threadIdx.x == 0) {
      if (indexed) {
        uint32_t vertex_count_clamped = min(vertex_capacity, vertex_pyramid[8].x);
        if (vertex_count_clamped) {
          extractIndexedVertexPN<<<(vertex_count_clamped + 255) / 256, 256>>>(vertex_output,
                                                                              vertex_pyramid,

        else if (rem == 4) vertex_index += item.w;
      }

      if (index_code & 0b001) {     // Vertex is on x-axis
        axes = 0b0000;              // That is the first vertex in that cell.
      }
      if (index_code & 0b010) {     // Vertex is on y-axis
        axes &= 0b0010;             // If cell has a vertex on the x-axis, that comes before the one on the y-axis.
      }
      if (index_code & 0b100) {     // Vertex is on z-axis
        axes &= 0b0110;             // If cell has vertices on either the x or y axes, those come before the one on the z-axis.
      }
      uint32_t n;
      asm("popc.b32 %0, %1;": "=r"(n) : "r"(axes));
      vertex_index += n;
