 public:
   au32 *bits;
 
-  __device__ twice_set() {
-    assert(bits = (au32 *)calloc(TWICE_WORDS, sizeof(au32)));
+  twice_set() {
+    checkCudaErrors(cudaMalloc((void**)&bits, TWICE_WORDS * sizeof(au32)));
+    checkCudaErrors(cudaMemset(bits, 0, TWICE_WORDS * sizeof(au32)));
   }
   __device__ void reset() {
-    memset(bits, 0, TWICE_WORDS*sizeof(au32));
+    memset(bits, 0, TWICE_WORDS * sizeof(au32));
   }
   __device__ void set(node_t u) {
     node_t idx = u/16;

    }
  }
#endif
  // free(tp->buckets);

#if 0
  thread_ctx *threads = (thread_ctx *)calloc(nthreads, sizeof(thread_ctx));

     (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) {
