  
             gupdate = 0.5*dtdx*bpasdq[mq];
             atomicAdd(&gm[I_q - 1 + ys],-gupdate);
-            atomicAdd(&gp[I_q - 1 + ys], -gupdate);
+            atomicAdd(&gp[I_q - 1 + ys],-gupdate);
 #endif            
         }            
     }

                               double *const amdq_trans, double *const apdq_trans, 
                               double *const bmdq_trans, double *const bpdq_trans,
                               double *const waves,      double *const speeds,
                               cudaclaw_cuda_rpn2_t rpn2,
                               cudaclaw_cuda_rpt2_t rpt2,
                               cudaclaw_cuda_b4step2_t b4step2,
                               double t,double dt)
{
    //typedef cub::BlockReduce<double,FC2D_CUDACLAW_BLOCK_SIZE> BlockReduce;
    
    //__shared__ typename BlockReduce::TempStorage temp_storage;

    extern __shared__ double shared_mem[];



            for(int mw = 0; mw < mwaves; mw++)
            {
                maxcfl = max(maxcfl,fabs(s[mw]*dtdx));

                if (order[0] == 2)
                {                    
                    int I_speeds = I + mw*zs;
                    speeds[I_speeds] = s[mw];
                    for(int mq = 0; mq < meqn; mq++)
                    {
                        int k = mw*meqn + mq;
