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


    printf("The first few entries look like %f %f %f.\n", dataset[0], dataset[1], dataset[2]);

    // make space for output results
    output_t *outputs = (output_t *)calloc(NUM_CHUNKS, sizeof(output_t));

    struct timeval start_time, end_time;
    double elapsed_time;
// #if 0 // DISABLED to make my life easier
    printf("Processing chunks sequentially\n");

    double *device_chunk;
    output_t *device_output;
    // TODO check errors
    checkCudaErrors(cudaMalloc(&device_chunk, CHUNK_BYTES));
    checkCudaErrors(cudaMalloc(&device_output, OUTPUT_BYTES));

    gettimeofday(&start_time, 0);


// simultaneous kernel executions supported on the Kepler architecture.
// Adjusting this may or may not provide performance improvement; I'm guessing
// that it's best to keep it near the device configuration.
#define SIMULTANEOUS_KERNELS 32
//#define SIMULTANEOUS_KERNELS 4
//#define SIMULTANEOUS_KERNELS 2


