   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));

    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];
      }

  uint2 u32;
};

static __device__ uint64_t c_m2[64];
static __device__ uint32_t d_error[1];

static __global__ void
