   }
   dst->h++; // pad by a line
   if(!self->src)
-  { CUTRY(cudaMalloc((void**)&self->src,pipeline_image_nbytes(src)));
+  { CUTRY(cudaMalloc((void**)&self->src,pipeline_image_nbytes(src)+1024));
     CUTRY(cudaMalloc((void**)&self->dst,pipeline_image_nbytes(dst)));
     CUTRY(cudaMalloc((void**)&self->tmp,self->nbytes_tmp=pipeline_image_nelem(dst)*sizeof(float)));
     CUTRY(cudaMemset(self->tmp,0,pipeline_image_nelem(dst)*sizeof(float)));

    { *emit=1;
      warp_kernel<Tsrc,BX,BY,WORK><<<blocks,threads>>>(self->ctx,(Tsrc*)self->src,self->tmp);
      cast_kernel<Tdst,BX,BY,WORK><<<blocks,threads>>>((Tdst*)self->dst,self->tmp,self->ctx.ostride*2,self->m*self->norm,self->b);
      CUTRY(cudaMemset(self->tmp,0,self->nbytes_tmp));
    } else
    { *emit=0;
      warp_kernel<Tsrc,BX,BY,WORK><<<blocks,threads>>>(self->ctx,(Tsrc*)self->src,self->tmp);

    for(unsigned i=0;i<halfw;++i)
      if(last!=lut[i])
        ilut[last=lut[i]]=i;
    ilut[ow  ]=inwidth/2; // add elements to deal with discontinuity
    ilut[ow+1]=inwidth; // subtract one to prevent reading off end
    ilut+=2;
    for(unsigned i=halfw;i<inwidth;++i)
      if(last!=lut[i])
