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

        }

        // Disabling this gains about .5 seconds on a 5 second run
        printf("thread %p fit chunk %d after %d iterations\n", pthread_self(), chunk_id, iteration);
        for (int m = 0; m < M; m++)
        {
            invgauss_params_t *p = &params_new[m];
            printf("\tcomp %d alpha=%lf mu=%lf lambda=%lf\n", m, p->alpha, p->mu, p->lambda);
        }

        // FIXME BODGE
        chunk_id = OSAtomicIncrement32(g_chunk_id);

    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
