 	cudaMemset(d_count, 0, sizeof(int));
 
   // Try uncommenting one kernel call at a time
-	//xyzw_frequency<<<8, 256>>>(d_count, d_text, len);
-  xyzw_frequency_thrust_device<<<1, 1>>>(d_count, d_text, len);
+	xyzw_frequency<<<8, 256>>>(d_count, d_text, len);
+  //xyzw_frequency_thrust_device<<<1, 1>>>(d_count, d_text, len);
   cudaMemcpy(&count, d_count, sizeof(int), cudaMemcpyDeviceToHost);
   
   //xyzw_frequency_thrust_host(&count, h_text, len);

    return false;
  });
}

// a bug in Thrust 1.8 causes warnings when this is uncommented
// so commented out by default -- fixed in Thrust master branch

  });
}

// A bug in CUDA 7.0 causes errors when this is called
// Comment out by default, but will work in CUDA 7.5
#if 0
__global__
void xyzw_frequency_thrust_device(int *count, char *text, int n)
{
