     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

                                 OutType*   d_out,
                                 cudaStream_t stream=0) {
    //cout << "LAUNCH for " << nelement << endl;
    dim3 block(8,1);
    dim3 grid(nbatch,1);
    /*

    cuda::child_stream stream(g_cuda_stream);
    BF_TRACE_STREAM(stream);
    
    scatter_grid_kernel <<< nbatch, 8, 0, stream >>> ((cuComplex*)dptr,
						  (cuComplex*)uvgridptr,
						  (cuComplex*)illumptr,
