   printf("Elapsed time = %.6f seconds.\n", time);
   double gups = 4 * n / (double) ms * 1.0e-6;
   printf("Giga Updates per second = %.6f GUP/s.\n", gups);
-  d_bench<<<grid, thread>>>(n, d_t);
+  d_bench<ATOMICTYPE_CAS><<<grid, thread>>>(n, d_t);
   void *p_error;
   cudaGetSymbolAddress(&p_error, d_error);
   cudaMemset(d_error, 0, sizeof(uint32_t));

  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);
    atomicXor(&t[ran.u64 & (n - 1)].u32.x, ran.u32.x);
    atomicXor(&t[ran.u64 & (n - 1)].u32.y, ran.u32.y);
  }
}


  return ran;
}

enum atomictype_t {
  ATOMICTYPE_CAS,
  ATOMICTYPE_XOR,
};

template<atomictype_t ATOMICTYPE>
__global__ void
d_bench(size_t n, benchtype *t)
{
  size_t num_threads = gridDim.x * blockDim.x;
