     cycle.insert(edge(us[(nu+1)&~1], us[nu|1])); // u's in even position; v's in odd
   while (nv--)
     cycle.insert(edge(vs[nv|1], vs[(nv+1)&~1])); // u's in odd position; v's in even
-  u32 soli = atomicAdd(&ctx->nsols, 1U);
   for (nonce_t nonce = n = 0; nonce < HALFSIZE; nonce++)
     if (ctx->alive->test(nonce)) {
       edge e(sipnode(&ctx->sip_ctx, nonce, 0), HALFSIZE+sipnode(&ctx->sip_ctx, nonce, 1));
       if (cycle.find(e) != cycle.end()) {
-        ctx->sols[soli][n++] = nonce;
+        printf(" %lx", (long)nonce);
         if (PROOFSIZE > 2)
           cycle.erase(e);
       }
     }
   assert(n==PROOFSIZE);
-}
+  printf("\n");
+}
+#endif
 
 #include <unistd.h>
 

public:
  au32 *bits;

  __device__ twice_set() {
    assert(bits = (au32 *)calloc(TWICE_WORDS, sizeof(au32)));
  }
  __device__ void reset() {
    memset(bits, 0, TWICE_WORDS*sizeof(au32));
  }
  __device__ void set(node_t u) {
    node_t idx = u/16;

     (int)edgeBytes, " KMGT"[edgeUnit], (int)nodeBytes, " KMGT"[nodeUnit]);
  cuckoo_ctx ctx(header, nthreads, ntrims, maxsols);

  cuckoo_ctx *device_ctx;
  checkCudaErrors(cudaMalloc((void**)&device_ctx, sizeof(cuckoo_ctx)));
  cudaMemcpy(device_ctx, &ctx, sizeof(cuckoo_ctx), cudaMemcpyHostToDevice);

  trim_edges<<<nthreads,1>>>(device_ctx);

  u32 *bits;
  assert(bits = (u32 *)calloc(HALFSIZE/32, sizeof(u32)));
  cudaMemcpy(bits, ctx.alive->bits, (HALFSIZE/32) * sizeof(u32), cudaMemcpyDeviceToHost);

  u32 cnt = 0;
  for (nonce_t nonce = 0; nonce < HALFSIZE; nonce++)
    cnt += ((bits[nonce/32] >> (nonce%32)) & 1) ^ 1;

  u32 load = (u32)(100 * cnt / CUCKOO_SIZE);
  printf("final load %d%%\n", load);

#if 0
  if (tp->id == 0) {
