             cudaMemcpy(left,d_left,tpb*sizeof(REAL),cudaMemcpyDeviceToHost);
 
             CPU_diamond(right, left, tpb);
+            wholeDiamond <<< bks-1,tpb,smem2 >>>(d_right,d_left,1);
 
             cudaMemcpy(d_right, right, tpb*sizeof(REAL), cudaMemcpyHostToDevice);
             cudaMemcpy(d_left, left, tpb*sizeof(REAL), cudaMemcpyHostToDevice);

        while(t_eq < t_end)
        {

            splitDiamond <<< bks,tpb,smem2 >>>(d_right,d_left,-1);

            wholeDiamond <<< bks,tpb,smem2 >>>(d_right,d_left,1);
            //So it always ends on a left pass since the down triangle is a right pass.

	double t_eq = t_fullstep;

	// Call the kernels until you reach the iteration limit.
    if (cpu)
    {
    	while(t_eq < t_end)
    	{
