     cudaStream_t stream;
     // cudaSetDevice(0); // TODO: adjust when we have multiple GPUs; probably assign a new group of threads to each GPU
     checkCudaErrors(cudaStreamCreate(&stream));
+    checkCudaErrors(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
 
     // persistent device memory
     // allocated out here to avoid cudaMalloc in main loop

        {
            iteration++;
            checkCudaErrors(cudaMemcpyAsync(dev_params, params_new, sizeof(invgauss_params_t) * M, cudaMemcpyHostToDevice, stream));
            checkCudaErrors(cudaStreamSynchronize(stream));

            //////// PROCESS CHUNK


            // The remaining operations are all summations over various
            // outputs from this kernel.

            // checkCudaErrors(cudaStreamSynchronize(stream));

            // have we converged?
            // TODO(perf): we could probably save some time by only doing this check once every few iterations - it's slow relative to the rest of the iteration time
