   d_init<<<grid, thread>>>(n, d_t);
   cudaEventRecord(begin);
   cudaEventSynchronize(begin);
-  d_bench<<<grid, thread>>>(n, d_t);
+  d_bench<ATOMICTYPE_CAS><<<grid, thread>>>(n, d_t);
   cudaEventRecord(end);
   cudaEventSynchronize(end);
 

    return 1;
  }

  ptrdiff_t i;
  for (i = 63; i >= 0; --i) {
    if ((n >> i) & 1) {
      break;
    }
  }

  uint64_t ran = 2;
  while (i > 0) {
    uint64_t temp = 0;
    for (ptrdiff_t j = 0; j < 64; j++) {
      if ((ran >> j) & 1) {
        temp ^= c_m2[j];
      }

  ran.u64 = d_starts(start);
  for (ptrdiff_t i = start; i < end; ++i) {
    ran.u64 = (ran.u64 << 1) ^ ((int64_t) ran.u64 < 0 ? POLY : 0);
    switch (ATOMICTYPE) {
    case ATOMICTYPE_CAS:
      unsigned long long int *address, old, assumed;
      address = (unsigned long long int *)&t[ran.u64 & (n - 1)].u64;
      old = *address;
      do {
        assumed = old;
        old = atomicCAS(address, assumed, assumed ^ ran.u64);
      } while  (assumed != old);
      break;
    case ATOMICTYPE_XOR:
      atomicXor(&t[ran.u64 & (n - 1)].u32.x, ran.u32.x);
      atomicXor(&t[ran.u64 & (n - 1)].u32.y, ran.u32.y);
      break;
    }
  }
}

