                                 double* waves, double *speeds,
                                 cudaclaw_cuda_rpn2_t rpn2)
 {
+    typedef cub::BlockReduce<int,128> BlockReduce;
+    __shared__ typename BlockReduce::TempStorage temp_storage;
     int mq, mw, m;
     int xs, ys, zs;
     int I, I_q, I_aux, I_waves, I_speeds;

#include "cudaclaw_flux2.h"
#include <fclaw_base.h>  /* Needed for SC_MIN, SC_MAX */
#include "cudaclaw_allocate.h"

#include <math.h>



            {
                I_speeds = I + (mwaves+mw)*zs;
                speeds[I_speeds] = s[mw];
                if(fabs(s[mw])*dtdy>max_speed){
                    max_speed=fabs(s[mw])*dtdy;
                }
            } 

#if 0        
