 
         // printf("thread %d chunk %d host_ptr %p\n", args->thread_id, chunk_id, host_chunk_ptr);
 
-
-        // printf("%d: pre copy\n", args->thread_id);
         cudaMemcpyAsync(args->device_chunk, host_chunk_ptr, CHUNK_BYTES, cudaMemcpyHostToDevice, stream);
-        // printf("%d: pre-sync\n", args->thread_id);
         cudaStreamSynchronize(stream);
+
         // do work on gpu
-        // printf("%d: launch\n", args->thread_id);
-        // 2048 simultaneous launches?
-        simple_kernel<<<GRID, BLOCK, 0, stream>>>(args->device_chunk);//, args->device_output);
-        // simple_kernel<<< grid, threads, mem_size, ??? >>>(data, left, nright, depth+1);
-        // kernel_call<<<dimGrid, dimBlock, 0>>>();
+        simple_kernel<<<args->gridSize, args->blockSize, 0, stream>>>(args->device_chunk, args->device_output);
         cudaStreamSynchronize(stream);
-        // printf("%d: sync\n", args->thread_id);
-        // These small (16 byte) copies are relatively inefficient. Eventually we'll have all of the posterior probabilities to copy back, so it might not matter.
+
+        // These small (24 byte) copies are relatively inefficient. Eventually we'll have all of the posterior probabilities to copy back, so it might not matter.
         cudaMemcpyAsync(&(args->host_outputs[chunk_id]), args->device_output, sizeof(output_t), cudaMemcpyDeviceToHost, stream);
         cudaStreamSynchronize(stream);
 

    cudaStream_t stream;

    cudaSetDevice(0); // TODO: adjust when we have multiple GPUs; probably assign a new group of threads to each GPU

    error = cudaStreamCreate(&stream);


        args[i].current_chunk_id = &current_chunk_id;
        args[i].host_dataset = dataset;
        args[i].host_outputs = host_outputs;
        args[i].gridSize = gridSize;
        args[i].blockSize = blockSize;

        int rc = pthread_create(&threads[i], NULL, thread, (void *)&args[i]);
        assert(rc == 0);
