 // grow with cube root of size, hardly affected by trimming
 #define MAXPATHLEN (8 << (SIZESHIFT/3))
 
+#define checkCudaErrors(ans) { gpuAssert((ans), __FILE__, __LINE__); }
+inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) {
+  if (code != cudaSuccess) {
+    fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
+    if (abort) exit(code);
+  }
+}
+
 // set that starts out full and gets reset by threads on disjoint words
 class shrinkingset {
 public:
-  std::vector<u32> bits;
-  std::vector<u64> cnt;
-
-  __device__ shrinkingset(u32 nthreads) {
+  u32 *bits;
+
+  shrinkingset() {
     nonce_t nwords = HALFSIZE/32;
-    bits.resize(nwords);
-    cnt.resize(nthreads);
-    cnt[0] = HALFSIZE;
-  }
-  __device__ u64 count() const {
-    u64 sum = 0L;
-    for (u32 i=0; i<cnt.size(); i++)
-      sum += cnt[i];
-    return sum;
-  }
-  __device__ void reset(nonce_t n, u32 thread) {
+    checkCudaErrors(cudaMalloc((void**)&bits, nwords * sizeof(u32)));
+    checkCudaErrors(cudaMemset(bits, 0, nwords * sizeof(u32)));
+  }
+  __device__ void reset(nonce_t n) {
     bits[n/32] |= 1 << (n%32);
-    cnt[thread]--;
   }
   __device__ bool test(node_t n) const {
     return !((bits[n/32] >> (n%32)) & 1);
   }
   __device__ u32 block(node_t n) const {
     return ~bits[n/32];
+  }
+  ~shrinkingset() {
+    checkCudaErrors(cudaFree(bits));
   }
 };
 

  for (nonce_t block = tp->id*32; block < HALFSIZE; block += ctx->nthreads*32) {
    for (nonce_t nonce = block; nonce < block+32 && nonce < HALFSIZE; nonce++) {
      if (alive->test(nonce)) {
        node_t u0, v0;
        sipedge(&ctx->sip_ctx, nonce, &u0, &v0);
        v0 += HALFSIZE;  // make v's different from u's
        node_t u = cuckoo[u0], v = cuckoo[v0];
        us[0] = u0;
        vs[0] = v0;
        u32 nu = path(cuckoo, u, us), nv = path(cuckoo, v, vs);
        if (us[nu] == vs[nv]) {
          u32 min = nu < nv ? nu : nv;

  __device__ u32 test(node_t u) const {
    return (bits[u/16] >> (2 * (u%16))) & 2;
  }
  ~twice_set() {
    checkCudaErrors(cudaFree(bits));
  }
};

