         else if (rem == 4) vertex_index += item.w;
       }
 
-      if (index_code & 0b001) {
-        axes = 0b0000;
+      if (index_code & 0b001) {     // Vertex is on x-axis
+        axes = 0b0000;              // That is the first vertex in that cell.
       }
-      if (index_code & 0b010) {       // Want Y-axis
-        axes &= 0b0010;                     // Add 1 for X-axis if present
+      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) {       // Want Z-axis
-        axes &= 0b0110;                     // Add 1 for X and Y-axes if present
-      }                                     // Else X-axis and we need no adjustment
+      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;

  {
     // FIXME: no point in up-traversal if cell doesn't change.
     // FIXME: If extract 3 in a go, organize traversals to reduce up-traversal.


    uint32_t ix = blockDim.x * blockIdx.x + threadIdx.x;


  __device__ __forceinline__ uint32_t piercingAxesFromCase(uint32_t c)
  {
    uint32_t a;
    asm("{\n"
        "  .reg .u32 t1, t2, t3;\n"
        "  bfe.s32 t1, %1, 0, 1;\n"     // Sign-extend bit 0 (0,0,0) over whole word
        "  xor.b32 t2, %1, t1;\n"       // XOR with (0,0,0) to see which corners change sign wrt (0,0,0).
        "  and.b32 %0, t2, 0b10110;\n"  // Mask out (1,0,0), (0,1,0) and (0,0,1).
        "}"
        : "=r"(a) : "r"(c));
    return a;
  }

  __device__ __forceinline__ uint32_t axisCountFromCase(const uint32_t c)
  {
    uint32_t a = piercingAxesFromCase(c);
    uint32_t n;
    asm("popc.b32 %0, %1;": "=r"(n) : "r"(a));         // Count number of 1's (should be 0..3).
    return n;
  }

