   { 
     #pragma unroll
     for(int i=0;i<WORK;++i) // in-place intensity scaling
-      buf[threadIdx.y][threadIdx.x+i*BX]=fma(buf[threadIdx.y][threadIdx.x+i*BX],m,b);
+      buf[threadIdx.y][threadIdx.x+i*BX]=fma(buf[threadIdx.y][threadIdx.x+i*BX],ctx.m,ctx.b);
   }
   warp<T,BX,BY,WORK>(ctx,buf,dst,ox,oy);
 }

}

static int pipeline_upload(pipeline_t self, pipeline_image_t dst, const pipeline_image_t src)
{ if(self->src && (self->ctx.w>src->width || self->ctx.h>src->height))
  { CUTRY(cudaFree(self->src)); self->src=0;
    CUTRY(cudaFree(self->dst)); self->dst=0;
    CUTRY(cudaFree(self->tmp)); self->tmp=0;
  }
  if(!self->src)
  { CUTRY(cudaMalloc(&self->src,pipeline_image_nbytes(src)));
    CUTRY(cudaMalloc(&self->src,pipeline_image_nbytes(dst)));
    CUTRY(cudaMalloc(&self->tmp,pipeline_image_nelem(dst)*sizeof(float)));
  }
  CUTRY(cudaMemset(self->tmp,0,pipeline_image_nelem(dst)*sizeof(float))));
  CUTRY(cudaMemcpy(self->src,src->data,pipeline_image_nbytes(src),cudaMemcpyHostToDevice));
  self->ctx.w=src->width;
  self->ctx.h=src->height;
  self->ctx.istride=src->stride;
  self->ctx.ostride=dst->stride;
  return 1;

}

int pipeline_exec(pipeline_t self, pipeline_image_t dst, const pipeline_image_t src)
{ if(src->w>self->ctx.w)
  { TRY(pipeline_alloc_accumulator(self,src)); // these will free if one already exists
    TRY(pipeline_alloc_lut(self,src->w));
    TRY(pipeline_fill_lut(self,src->w));
  }
  pipeline_image_conversion_params(dst,src,self->invert,&self->ctx.m,&self->ctx.b);
  TRY(pipeline_upload(self,dst,src)); // updates context size and stride as well
