                                  OutType*   d_out,
                                  cudaStream_t stream=0) {
     //cout << "LAUNCH for " << nelement << endl;
+    // TODO: Is this really the best setup to use?
     dim3 block(8,1);
     dim3 grid(nbatch,1);
     /*

    cuda::child_stream stream(g_cuda_stream);
    BF_TRACE_STREAM(stream);
    
    /*
    dim3 block(16);
    dim3 grid(nbatch);
    
    void* args[] = {&dptr,
                    &uvgridptr, 
                    &illumptr,
                    &xloc,
                    &yloc,
                    &zloc,
                    &max_support,
                    &grid_size,
                    &data_size};
    
    BF_CHECK_CUDA_EXCEPTION(cudaLaunchKernel((void*)scatter_grid_kernel, 
                                             grid, block,
                                             &args[0], 0, stream), 
                            BF_STATUS_INTERNAL_ERROR);
    */
    
    
    scatter_grid_kernel <<< nbatch, 8, 0, stream >>> ((cuComplex*)dptr,
						  (cuComplex*)uvgridptr,
						  (cuComplex*)illumptr,

    BF_ASSERT( in->ndim == 4,          BF_STATUS_INVALID_SHAPE);
    BF_ASSERT(out->ndim == in->ndim+1, BF_STATUS_INVALID_SHAPE);
    
//     TODO: What all should we support here?
//     BFarray in_flattened;
//     if( in->ndim > 4 ) {
//         // Keep the last three dim but attempt to flatten all others
