Ticket #370: csc-with-stride.patch

xpra/codecs/codec_constants.py

@@ -4 +4 @@
 # Xpra is released under the terms of the GNU GPL v2, or, at your option, any
 # later version. See the file COPYING for details.
 
+#this table is used to convert a dimension (width x height)
+#into a list of plane sizes
 PIXEL_SUBSAMPLING = {
+         "NV12"      : ((1, 1), (1, 2)),            #2 planes: "Y" and "UV" packed together (U and V are subsampled by 2 in both dimensions - but packed the width is "unchanged")
          "YUV420P"   : ((1, 1), (2, 2), (2, 2)),
          "YUV422P"   : ((1, 1), (2, 1), (2, 1)),
          "YUV444P"   : ((1, 1), (1, 1), (1, 1)),
@@ -48 +51 @@
                     quality=100, speed=100,
                     setup_cost=50, cpu_cost=100, gpu_cost=0,
                     min_w=1, min_h=1, max_w=4*1024, max_h=4*1024, max_pixels=4*1024*4*1024,
+                    width_stride_rounding = 1, height_stride_rounding = 1,
                     can_scale=False,
                     width_mask=0xFFFF, height_mask=0xFFFF):
         self.codec_class = codec_class          #ie: xpra.codecs.enc_x264.encoder.Encoder
@@ -61 +65 @@
         self.min_h = min_h
         self.max_w = max_w
         self.max_h = max_h
+        self.width_stride_rounding = width_stride_rounding
+        self.height_stride_rounding = height_stride_rounding
         self.max_pixels = max_pixels
         self.width_mask = width_mask
         self.height_mask = height_mask

xpra/codecs/csc_nvcuda/colorspace_converter.py

@@ -260 +260 @@
         self.dst_width = 0
         self.dst_height = 0
         self.dst_format = ""
+        self.width_stride_rounding = 0
+        self.height_stride_rounding = 0
         self.time = 0
         self.frames = 0
         self.kernel_function = None
         self.context = None
 
     def init_context(self, src_width, src_height, src_format,
-                           dst_width, dst_height, dst_format, speed=100):  #@DuplicatedSignature
+                           dst_width, dst_height, dst_format,
+                           width_stride_rounding, height_stride_rounding,
+                           speed=100):  #@DuplicatedSignature
         validate_in_out(src_format, dst_format)
         init_context()
         self.src_width = src_width
@@ -275 +279 @@
         self.dst_width = dst_width
         self.dst_height = dst_height
         self.dst_format = dst_format
+        self.width_stride_rounding = width_stride_rounding
+        self.height_stride_rounding = height_stride_rounding
         self.context = context
         k = (src_format, dst_format)
         npp_fn = COLORSPACES_MAP.get(k)
@@ -452 +458 @@
             driver.memcpy_dtoh_async(pixels, out_buf, stream)
         else:
             #we don't want the crazy large GPU padding, so we do it ourselves:
-            stride = width*4
+            stride = roundup(width*4, self.width_stride_rounding)
             pixels = driver.pagelocked_empty(stride*height, dtype=numpy.byte)
             copy = driver.Memcpy2D()
             copy.set_src_device(out_buf)
@@ -508 +514 @@
         out_sizes = []
         for i in range(3):
             x_div, y_div = divs[i]
-            out_stride = roundup(width/x_div, 4)
-            out_height = roundup(height/y_div, 2)
+            out_stride = roundup(width/x_div, self.width_stride_rounding)
+            out_height = roundup(height/y_div, self.height_stride_rounding)
             out_buf, out_stride = driver.mem_alloc_pitch(out_stride, out_height, 4)
             out_bufs.append(out_buf)
             out_strides.append(out_stride)

xpra/codecs/csc_opencl/colorspace_converter.py

@@ -159 +159 @@
     #log.info("has_same_channels(%s, %s)=%s (%s - %s)", src, dst, scheck, dcheck, len(scheck)==0 and len(dcheck)==0)
     return len(scheck)==0 and len(dcheck)==0
 
-KERNELS_DEFS = {}
+
 def gen_yuv_to_rgb():
     global context
     from xpra.codecs.csc_opencl.opencl_kernels import gen_yuv_to_rgb_kernels, rgb_mode_to_indexes, indexes_to_rgb_mode
@@ -316 +316 @@
     return RGB_to_YUV_KERNELS
 
 
+KERNELS_DEFS = {}
 def gen_kernels():
     """
     The code here is complicated by the fact that we don't know
@@ -412 +413 @@
         self.dst_width = 0
         self.dst_height = 0
         self.dst_format = ""
+        self.width_stride_rounding = 0
+        self.height_stride_rounding = 0
         self.time = 0
         self.frames = 0
         self.queue = None
@@ -421 +424 @@
         self.kernel_function_name = None
 
     def init_context(self, src_width, src_height, src_format,
-                           dst_width, dst_height, dst_format, csc_speed=100):  #@DuplicatedSignature
+                           dst_width, dst_height, dst_format,
+                           width_stride_rounding, height_stride_rounding,
+                           csc_speed=100):  #@DuplicatedSignature
         global context
         debug("init_context%s", (src_width, src_height, src_format, dst_width, dst_height, dst_format, csc_speed))
         validate_in_out(src_format, dst_format)
@@ -431 +436 @@
         self.dst_width = dst_width
         self.dst_height = dst_height
         self.dst_format = dst_format
+        self.width_stride_rounding = width_stride_rounding
+        self.height_stride_rounding = height_stride_rounding
         self.queue = pyopencl.CommandQueue(context)
         #sampling type:
         if self.src_width>self.dst_width and self.src_height>self.dst_height:
@@ -557 +564 @@
 
         #output image:
         oformat = pyopencl.ImageFormat(self.channel_order, pyopencl.channel_type.UNORM_INT8)
-        oimage = pyopencl.Image(context, mem_flags.WRITE_ONLY, oformat, shape=(self.dst_width, self.dst_height))
+        outstride = roundup(self.dst_width, self.width_stride_rounding)
+        outheight = roundup(self.dst_height, self.height_stride_rounding)
+        oimage = pyopencl.Image(context, mem_flags.WRITE_ONLY, oformat, shape=(outstride, outheight))
 
         iformat = pyopencl.ImageFormat(pyopencl.channel_order.R, pyopencl.channel_type.UNSIGNED_INT8)
         for i in range(3):
@@ -571 +580 @@
             iimage = pyopencl.Image(context, flags, iformat, shape=shape, hostbuf=plane)
             kernelargs.append(iimage)
 
+        #FIXME: pass out stride!
         kernelargs += [numpy.int32(self.src_width), numpy.int32(self.src_height),
                        numpy.int32(self.dst_width), numpy.int32(self.dst_height),
                        self.sampler, oimage]
@@ -582 +592 @@
         kend = time.time()
         debug("%s took %.1fms", self.kernel_function, 1000.0*(kend-kstart))
 
-        out_array = numpy.empty(self.dst_width*self.dst_height*4, dtype=numpy.byte)
+        out_array = numpy.empty(outstride*outheight*4, dtype=numpy.byte)
         pyopencl.enqueue_read_image(self.queue, oimage, origin=(0, 0), region=(self.dst_width, self.dst_height), hostbuf=out_array, is_blocking=True)
         self.queue.finish()
         debug("readback using %s took %.1fms", CHANNEL_ORDER_TO_STR.get(self.channel_order), 1000.0*(time.time()-kend))
@@ -607 +617 @@
         #adjust work dimensions for subsampling:
         #(we process N pixels at a time in each dimension)
         divs = get_subsampling_divs(self.dst_format)
+        assert len(divs) in ImageWrapper.PLANE_OPTIONS, "invalid number of planes: %s" % len(divs)
         wwidth = dimdiv(self.dst_width, max([x_div for x_div, _ in divs]))
         wheight = dimdiv(self.dst_height, max([y_div for _, y_div in divs]))
         globalWorkSize, localWorkSize  = self.get_work_sizes(wwidth, wheight)
@@ -635 +646 @@
         strides = []
         out_buffers = []
         out_sizes = []
-        for i in range(3):
+        nplanes = len(divs)
+        for i in range(nplanes):
             x_div, y_div = divs[i]
-            p_stride = roundup(self.dst_width / x_div, max(2, localWorkSize[0]))
-            p_height = roundup(self.dst_height / y_div, 2)
+            p_stride = roundup(self.dst_width / x_div, max(2, localWorkSize[0], self.width_stride_rounding))
+            p_height = roundup(self.dst_height / y_div, 2, self.height_stride_rounding)
             p_size = p_stride * p_height
             #debug("output buffer for channel %s: stride=%s, height=%s, size=%s", i, p_stride, p_height, p_size)
             out_buf = pyopencl.Buffer(context, mem_flags.WRITE_ONLY, p_size)
@@ -656 +668 @@
         #read back:
         pixels = []
         read_events = []
-        for i in range(3):
+        for i in range(nplanes):
             out_array = numpy.empty(out_sizes[i], dtype=numpy.byte)
             pixels.append(out_array.data)
             read = pyopencl.enqueue_read_buffer(self.queue, out_buffers[i], out_array, is_blocking=False)
@@ -667 +679 @@
         self.queue.finish()
         readend = time.time()
         debug("wait for read events took %.1fms", 1000.0*(readend-readstart))
-        return ImageWrapper(0, 0, self.dst_width, self.dst_height, pixels, self.dst_format, 24, strides, planes=ImageWrapper._3_PLANES)
+        return ImageWrapper(0, 0, self.dst_width, self.dst_height, pixels, self.dst_format, 24, strides, planes=nplanes)

xpra/codecs/csc_opencl/opencl_kernels.py

@@ -393 +393 @@
     return kname, kstr % args
 
 
+def gen_rgb_to_nv12_kernel(rgb_mode):
+    RGB_args = rgb_indexes(rgb_mode)
+    #kernel args: R, G, B are used 12 times each:
+    kname = "%s_to_nv12" % indexes_to_rgb_mode(RGB_args)
+    args = tuple([kname]+RGB_args*12)
+
+    kstr = """
+__kernel void %s(read_only image2d_t src,
+              uint srcw, uint srch, uint w, uint h,
+              const sampler_t sampler,
+              global uchar *dstY, uint strideY,
+              global uchar *dstUV, uint strideUV) {
+    uint gx = get_global_id(0);
+    uint gy = get_global_id(1);
+
+    if ((gx*2 < w) & (gy*2 < h)) {
+        uint srcx = gx*2*srcw/w;
+        uint srcy = gy*2*srch/h;
+        uint4 p1 = read_imageui(src, sampler, (int2)( srcx, srcy ));
+        uint4 p2 = p1;
+        uint4 p3 = p1;
+        uint4 p4 = p1;
+
+        //write up to 4 Y pixels:
+        float Y1 =  (0.257 * p1.s%s + 0.504 * p1.s%s + 0.098 * p1.s%s + 16);
+        //same logic as 422P for missing pixels:
+        uint i = gx*2 + gy*2*strideY;
+        dstY[i] = convert_uchar_rte(Y1);
+        if (gx*2+1 < w) {
+            srcx = (gx*2+1)*srcw/w;
+            p2 = read_imageui(src, sampler, (int2)( srcx, srcy ));
+            float Y2 =  (0.257 * p2.s%s + 0.504 * p2.s%s + 0.098 * p2.s%s + 16);
+            dstY[i+1] = convert_uchar_rte(Y2);
+        }
+        if (gy*2+1 < h) {
+            i += strideY;
+            srcx = gx*2*srcw/w;
+            srcy = (gy*2+1)*srch/h;
+            p3 = read_imageui(src, sampler, (int2)( srcx, srcy ));
+            float Y3 =  (0.257 * p3.s%s + 0.504 * p3.s%s + 0.098 * p3.s%s + 16);
+            dstY[i] = convert_uchar_rte(Y3);
+            if (gx*2+1 < w) {
+                srcx = (gx*2+1)*srcw/w;
+                p4 = read_imageui(src, sampler, (int2)( srcx, srcy ));
+                float Y4 =  (0.257 * p4.s%s + 0.504 * p4.s%s + 0.098 * p4.s%s + 16);
+                dstY[i+1] = convert_uchar_rte(Y4);
+            }
+        }
+        uint UVpos = gy*strideUV + gx*2;
+
+        //write 1 U pixel:
+        float U1 = (-0.148 * p1.s%s - 0.291 * p1.s%s + 0.439 * p1.s%s + 128);
+        float U2 = (-0.148 * p2.s%s - 0.291 * p2.s%s + 0.439 * p2.s%s + 128);
+        float U3 = (-0.148 * p3.s%s - 0.291 * p3.s%s + 0.439 * p3.s%s + 128);
+        float U4 = (-0.148 * p4.s%s - 0.291 * p4.s%s + 0.439 * p4.s%s + 128);
+        dstUV[UVpos] = convert_uchar_rte((U1+U2+U3+U4)/4.0);
+
+        //write 1 V pixel:
+        float V1 =  (0.439 * p1.s%s - 0.368 * p1.s%s - 0.071 * p1.s%s + 128);
+        float V2 =  (0.439 * p2.s%s - 0.368 * p2.s%s - 0.071 * p2.s%s + 128);
+        float V3 =  (0.439 * p3.s%s - 0.368 * p3.s%s - 0.071 * p3.s%s + 128);
+        float V4 =  (0.439 * p4.s%s - 0.368 * p4.s%s - 0.071 * p4.s%s + 128);
+        dstUV[UVpos + 1] = convert_uchar_rte((V1+V2+V3+V4)/4.0);
+    }
+}
+"""
+    return kname, kstr % args
+
+
 RGB_to_YUV_generators = {
                     "YUV444P"   : gen_rgb_to_yuv444p_kernel,
                     "YUV422P"   : gen_rgb_to_yuv422p_kernel,
                     "YUV420P"   : gen_rgb_to_yuv420p_kernel,
+                    "NV12"      : gen_rgb_to_nv12_kernel,
                     }
 
 def gen_rgb_to_yuv_kernels(rgb_mode="RGBX", yuv_modes=YUV_FORMATS):

xpra/codecs/csc_swscale/colorspace_converter.pyx

@@ -81 +81 @@
 COLORSPACES = []
 #keeping this array in scope ensures the strings don't go away!
 FORMAT_OPTIONS = [
-#    ("AV_PIX_FMT_NV12",     (1, 1, 0, 0),       (1, 0.5, 0, 0),     "NV12"),
+    ("AV_PIX_FMT_NV12",     (1, 1, 0, 0),       (1, 0.5, 0, 0),     "NV12"),
     ("AV_PIX_FMT_RGB24",    (3, 0, 0, 0),       (1, 0, 0, 0),       "RGB"   ),
     ("AV_PIX_FMT_BGR24",    (3, 0, 0, 0),       (1, 0, 0, 0),       "BGR"   ),
     ("AV_PIX_FMT_0RGB",     (4, 0, 0, 0),       (1, 0, 0, 0),       "XRGB"  ),
@@ -238 +238 @@
     cdef int buffer_size
 
     def init_context(self, int src_width, int src_height, src_format,
-                           int dst_width, int dst_height, dst_format, int speed=100):    #@DuplicatedSignature
+                           int dst_width, int dst_height, dst_format,
+                           width_stride_rounding, height_stride_rounding,
+                           int speed=100):    #@DuplicatedSignature
         debug("swscale.ColorspaceConverter.init_context%s", (src_width, src_height, src_format, dst_width, dst_height, dst_format, speed))
         cdef CSCPixelFormat src
         cdef CSCPixelFormat dst
@@ -258 +260 @@
         self.buffer_size = 0
         for i in range(4):
             self.out_height[i] = (int) (dst_height * dst.height_mult[i])
-            self.out_stride[i] = roundup((int) (dst_width * dst.width_mult[i]), 4)
-            #add one extra line to height so we can read a full rowstride
-            #no matter where we start to read on the last line.
-            #MEMALIGN may be redundant here but it is very cheap
-            if dst_format=="NV12" and i==0:
-                #no padding: packed UV plane follows Y plane
-                self.out_size[i] = self.out_stride[i] * self.out_height[i]
-            else:
-                self.out_size[i] = pad(self.out_stride[i] * (self.out_height[i]+1))
+            self.out_stride[i] = roundup((int) (dst_width * dst.width_mult[i]), self.width_stride_rounding)
+            #ensure we always add at least one extra line to height via roundup
+            #so we can safely read a full rowstride at a time
+            #no matter where we start to read on the last line of the output image
+            padded_height = roundup(self.out_height[i]+1, self.height_stride_rounding)
+            self.out_size[i] = pad(self.out_stride[i] * padded_height)
             self.buffer_size += self.out_size[i]
         debug("buffer size=%s", self.buffer_size)
 
@@ -274 +273 @@
         self.src_height = src_height
         self.dst_width = dst_width
         self.dst_height = dst_height
+        self.width_stride_rounding = width_stride_rounding
+        self.height_stride_rounding = height_stride_rounding
 
         self.flags = get_swscale_flags(speed)
         self.time = 0
@@ -382 +383 @@
         csci = CSCImage()           #keep a reference to memory for cleanup
         for i in range(4):
             csci.set_plane(i, NULL)
-        if self.dst_format.endswith("P"):
+        if self.dst_format.endswith("P") or str(self.dst_format)=="NV12":
             #planar mode, assume 3 planes:
             oplanes = ImageWrapper._3_PLANES
+            if str(self.dst_format)=="NV12":
+                oplanes = ImageWrapper._2_PLANES
             out = []
             strides = []
-            for i in range(3):
+            for i in range(oplanes):
                 if self.out_stride[i]>0 and output_image[i]!=NULL:
                     stride = self.out_stride[i]
-                    plane = PyBuffer_FromMemory(<void *>output_image[i], self.out_height[i] * self.out_stride[i])
+                    plane = PyBuffer_FromMemory(<void *>output_image[i], self.out_size[i])
                 else:
                     stride = 0
                     plane = None
                 csci.set_plane(i, output_image[i])
                 out.append(plane)
                 strides.append(stride)
-        elif str(self.dst_format)=="NV12":
-            #Y plane, followed by U and V packed
-            oplanes = ImageWrapper.PACKED
-            strides = self.out_stride[0]
-            out = PyBuffer_FromMemory(<void *>output_image[0], self.buffer_size)
-            csci.set_plane(0, output_image[0])
         else:
             #assume no planes, plain RGB packed pixels:
             oplanes = ImageWrapper.PACKED
             strides = self.out_stride[0]
-            out = PyBuffer_FromMemory(<void *>output_image[0], self.out_height[0] * self.out_stride[0])
+            out = PyBuffer_FromMemory(<void *>output_image[0], self.out_size[0])
             csci.set_plane(0, output_image[0])
         elapsed = time.time()-start
         debug("%s took %.1fms", self, 1000.0*elapsed)

xpra/codecs/image_wrapper.py

@@ -8 +8 @@
 class ImageWrapper(object):
 
     PACKED = 0
+    _2_PLANES = 2
     _3_PLANES = 3
     _4_PLANES = 4
     PLANE_OPTIONS = (PACKED, _3_PLANES, _4_PLANES)
     PLANE_NAMES = {PACKED       : "PACKED",
+                   _2_PLANES    : "2_PLANES",
                    _3_PLANES    : "3_PLANES",
                    _4_PLANES    : "4_PLANES"}
 

xpra/codecs/nvenc/encoder.pyx

@@ -1024 +1024 @@
     return codec_spec(Encoder, codec_type=get_type(), encoding=encoding,
                       quality=60, setup_cost=100, cpu_cost=10, gpu_cost=100,
                       min_w=2, min_h=2, max_w=4096, max_h=4096,
-                      width_mask=0xFFFE, height_mask=0xFFFE)
+                      width_mask=0xFFFE, height_mask=0xFFFE,
+                      width_stride_rounding=32, height_stride_rounding=32)
 
 
 def get_version():