                                                                              counts.e0 + counts.e1,
                                                                              counts.e0 + counts.e1 + counts.e2,
                                                                              counts.e0 + counts.e1 + counts.e2 + counts.e3);
+          }
+          if (indexed && vsum) {
+            out_vertex_level0_d[32 * 5 * blockIdx.x + 32 * y + thread] = make_uint4(vc0,
+                                                                                    vc0 + vc1,
+                                                                                    vc0 + vc1 + vc2,
+                                                                                    vc0 + vc1 + vc2 + vc3);
+
           }
           
         }

      hp5_result r = traverseDown(pyramid, level_offset, level_offset[15], ix);

      uint8_t index_case = index_cases[r.offset];
      uint8_t index_code = index_table[16u * index_case + r.remainder];
      uint3 p0 = decodeCellPosition(chunks, r.offset);

      p0.x += (index_code >> 3) & 1;
      p0.y += (index_code >> 4) & 1;
      p0.z += (index_code >> 5) & 1;
      uint3 p1 = p0;
      p1.x += (index_code >> 0) & 1;
      p1.y += (index_code >> 1) & 1;
      p1.z += (index_code >> 2) & 1;

      emitVertexNormal(output + 6 * ix,
                       field,

        "  popc.b32 %0, t3;\n"          // Count number of 1's (should be 0..3).
        "}"
        : "=r"(n) :"r"(c));

    unsigned q = piercingAxesFromCase(c);
    assert(((q >> 1) & 1) + ((q >> 2) & 1) + ((q >> 4) & 1) == n);

    unsigned m =
      ((c & 1) ^ ((c >> 1) & 1)) +
