Using cufft and streams

[mnicely]

    PUSH_RANGE( "1st DST", 2 )
    CUDA_RT_CALL( cudaEventRecord( events[2], streams[2] ) );

    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[0], cudaEventWaitDefault ) );  // Wait for plan creation
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[1], cudaEventWaitDefault ) );  // Wait for plan2 creation
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[2], cudaEventWaitDefault ) );  // Wait for sys.rhs

    load_1st_DST_wrapper( streams[0], sys, dst, sys.rhs, in, in2 );

    CUDA_RT_CALL( cudaEventRecord( events[0], streams[0] ) );

    CUDA_RT_CALL( cudaStreamWaitEvent( streams[1], events[0], cudaEventWaitDefault ) );  // Wait for load_1st_DST
    CUDA_RT_CALL( cufftSetStream( plan, streams[0] ) );
    CUDA_RT_CALL( cufftSetWorkArea( plan, d_workspace[0] ) );
    CUDA_RT_CALL( cufftExecD2Z( plan, in, out ) );  // Running in streams[0]

    CUDA_RT_CALL( cufftSetStream( plan, streams[1] ) );
    CUDA_RT_CALL( cufftSetWorkArea( plan, d_workspace[1] ) );
    CUDA_RT_CALL( cufftExecD2Z( plan, in2, out2 ) );  // Running in streams[1]

    CUDA_RT_CALL( cudaEventRecord( events[0], streams[0] ) );
    CUDA_RT_CALL( cudaEventRecord( events[1], streams[1] ) );
    CUDA_RT_CALL( cudaEventRecord( events[2], streams[3] ) );
    POP_RANGE

    PUSH_RANGE( "Trig Solver", 3 )
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[0], cudaEventWaitDefault ) );  // Wait for plan execution
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[1], cudaEventWaitDefault ) );  // Wait for plan2 execution
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[2], cudaEventWaitDefault ) );  // Wait forsys.U, sys.L, sys.Up

    middle_stuff_ls_DST_wrapper( streams[0], sys, dst, out, out2, in, in2, d_y );

    CUDA_RT_CALL( cudaEventRecord( events[0], streams[0] ) );
    POP_RANGE

    PUSH_RANGE( "2nd DST", 4 )
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[1], events[0], cudaEventWaitDefault ) );  // Wait for middle_stuff_ls_DST
    CUDA_RT_CALL( cufftSetStream( plan, streams[0] ) );

    CUDA_RT_CALL( cufftSetWorkArea( plan, d_workspace[0] ) );
    CUDA_RT_CALL( cufftExecD2Z( plan, in, out ) );  // Running in streams[0]

    CUDA_RT_CALL( cufftSetStream( plan, streams[1] ) );
    CUDA_RT_CALL( cufftSetWorkArea( plan, d_workspace[1] ) );
    CUDA_RT_CALL( cufftExecD2Z( plan, in2, out2 ) );  // Running in streams[1]

    CUDA_RT_CALL( cudaEventRecord( events[0], streams[0] ) );
    CUDA_RT_CALL( cudaEventRecord( events[1], streams[1] ) );
    CUDA_RT_CALL( cudaEventRecord( events[2], streams[4] ) );

    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[0], cudaEventWaitDefault ) );  // Wait for plan execution
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[1], cudaEventWaitDefault ) );  // Wait for plan2 execution
    CUDA_RT_CALL( cudaStreamWaitEvent( streams[0], events[2], cudaEventWaitDefault ) );  // Wait for sys.sol
    store_2st_DST_wrapper( streams[0], sys, dst, out, out2, sys.sol );

    CUDA_RT_CALL( cudaStreamSynchronize( streams[0] ) );  // Wait for store_2st_DST
    POP_RANGE