                  every; ///< the number of frames to average
   double         samples_per_scan;
   bool           invert;
+  unsigned       downsample;
   unsigned       alignment;      ///< output rows are aligned to this target number of elements.
   void  * __restrict__ src,  ///< device buffer
         * __restrict__ dst;  ///< device buffer

    unsigned j;
    j = lut[i]   =   floor(v);
    fpart[i]     = v-floor(v);
    hit[j  ]+=fpart[i];
    hit[j+1]+=(1.0f-fpart[i]);
  }
  // compute norms
  for(unsigned i=0;i<inwidth;++i)

}

template<typename T,unsigned BX,unsigned BY,unsigned WORK,bool EMIT_FRAME_AVERAGE>
__device__ void accumulate(const pipeline_ctx_t & ctx,float* __restrict__ buf[BY],const unsigned ox, const unsigned oy)
{
  float *acc=ctx.accumulator+ox+oy*ctx.istride;
  #pragma unroll
  for(int i=0;i<WORK;++i)                // accumulate
    acc[i*BX]+=buf[threadIdx.y][threadIdx.x+i*BX];
