Ticket #558: nvenc-hacked-win32.patch

xpra/codecs/cuda_common/BGRA_to_NV12.cu

@@ -7 +7 @@
 
 #include <stdint.h>
 
-extern "C" __global__ void BGRA_to_NV12(uint8_t *srcImage, int src_w, int src_h, int srcPitch,
-                          uint8_t *dstImage, int dst_w, int dst_h, int dstPitch,
+extern "C" __global__ void BGRA_to_NV12(uint8_t *srcImage, int src_w, int src_h, int srcHPitch,
+                          uint8_t *dstImage, int dst_w, int dst_h,
+                          int dstHPitch, int dstVPitch,
                           int w, int h)
 {
     uint32_t gx, gy;
@@ -21 +22 @@
     if ((src_x < w) & (src_y < h)) {
         //4 bytes per pixel, and 2 pixels width/height at a time:
         //byte index:
-        uint32_t si = (src_y * srcPitch) + src_x * 4;
+        uint32_t si = (src_y * srcHPitch) + src_x * 4;
 
         //we may read up to 4 32-bit RGB pixels:
         uint8_t R[4];
@@ -38 +39 @@
         }
 
         //write up to 4 Y pixels:
-        uint32_t di = (gy * 2 * dstPitch) + gx * 2;
+        uint32_t di = (gy * 2 * dstHPitch) + gx * 2;
         dstImage[di] = __float2int_rn(0.257 * R[0] + 0.504 * G[0] + 0.098 * B[0] + 16);
         if (gx*2 + 1 < src_w) {
             R[1] = srcImage[si+6];
@@ -47 +48 @@
             dstImage[di + 1] = __float2int_rn(0.257 * R[1] + 0.504 * G[1] + 0.098 * B[1] + 16);
         }
         if (gy*2 + 1 < src_h) {
-            si += srcPitch;
-            di += dstPitch;
+            si += srcHPitch;
+            di += dstHPitch;
             R[2] = srcImage[si+2];
             G[2] = srcImage[si+1];
             B[2] = srcImage[si];
@@ -68 +69 @@
             u += -0.148 * R[j] - 0.291 * G[j] + 0.439 * B[j] + 128;
             v +=  0.439 * R[j] - 0.368 * G[j] - 0.071 * B[j] + 128;
         }
-        di = (dst_h + gy) * dstPitch + gx * 2;
+        di = (dstVPitch + gy) * dstHPitch + gx * 2;
         dstImage[di]      = __float2int_rn(u / 4.0);
         dstImage[di + 1]  = __float2int_rn(v / 4.0);
     }

xpra/codecs/cuda_common/BGRA_to_YUV444.cu

@@ -7 +7 @@
 
 #include <stdint.h>
 
-extern "C" __global__ void BGRA_to_YUV444(uint8_t *srcImage, int src_w, int src_h, int srcPitch,
-                             uint8_t *dstImage, int dst_w, int dst_h, int dstPitch,
-                             int w, int h)
+extern "C" __global__ void BGRA_to_YUV444(uint8_t *srcImage, int src_w, int src_h, int srcHPitch,
+                          uint8_t *dstImage, int dst_w, int dst_h,
+                          int dstHPitch, int dstVPitch,
+                          int w, int h)
 {
     uint32_t gx, gy;
     gx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -23 +24 @@
         uint8_t G;
         uint8_t B;
         //one 32-bit RGB pixel at a time:
-        uint32_t si = (src_y * srcPitch) + src_x * 4;
+        uint32_t si = (src_y * srcHPitch) + src_x * 4;
         R = srcImage[si+2];
         G = srcImage[si+1];
         B = srcImage[si];
 
         uint32_t di;
-        di = (gy * dstPitch) + gx;
+        di = (gy * dstHPitch) + gx;
         dstImage[di] = __float2int_rn(0.257 * R + 0.504 * G + 0.098 * B + 16);
-        di += dstPitch*dst_h;
+        di += dstHPitch*dstVPitch;
         dstImage[di] = __float2int_rn(-0.148 * R - 0.291 * G + 0.439 * B + 128);
-        di += dstPitch*dst_h;
+        di += dstHPitch*dstVPitch;
         dstImage[di] = __float2int_rn(0.439 * R - 0.368 * G - 0.071 * B + 128);
     }
 }

xpra/codecs/nvenc4/encoder.pyx

@@ -40 +40 @@
 cdef int YUV444_THRESHOLD = int(os.environ.get("XPRA_NVENC_YUV444_THRESHOLD", "85"))
 cdef int LOSSLESS_THRESHOLD = int(os.environ.get("XPRA_NVENC_LOSSLESS_THRESHOLD", "100"))
 cdef int DEBUG_API = int(os.environ.get("XPRA_NVENC_DEBUG_API", "0")=="1")
+cdef int VADJUST = int(os.environ.get("XPRA_NVENC_VADJUST", "0"))
+cdef int VPAD = int(os.environ.get("XPRA_NVENC_VPAD", "0"))
+cdef int VROUND = int(os.environ.get("XPRA_NVENC_VROUND", "0"))
 
 cdef int QP_MAX_VALUE = 51   #newer versions of ffmpeg can decode up to 63
 
@@ -1295 +1298 @@
         self.quality = quality
         self.scaling = scaling or (1, 1)
         v, u = self.scaling
-        self.input_width = roundup(width, 32)
-        self.input_height = roundup(height, 32)
-        self.encoder_width = roundup(width*v//u, 32)
-        self.encoder_height = roundup(height*v//u, 32)
+        self.input_width = roundup(width, 64)
+        self.input_height = roundup(height+VPAD, VROUND)
+        self.encoder_width = roundup(width*v//u, 64)
+        self.encoder_height = roundup((height+VPAD)*v//u, VROUND)
         self.src_format = src_format
         self.dst_formats = dst_formats
         self.codec_name = "H264"
@@ -1401 +1404 @@
             self.cudaInputBuffer, self.inputPitch = driver.mem_alloc_pitch(max_input_stride, self.input_height, 16)
             log("CUDA Input Buffer=%#x, pitch=%s", int(self.cudaInputBuffer), self.inputPitch)
             #allocate CUDA output buffer (on device):
-            self.cudaOutputBuffer, self.outputPitch = driver.mem_alloc_pitch(self.encoder_width, self.encoder_height*3//plane_size_div, 16)
+            self.cudaOutputBuffer, self.outputPitch = driver.mem_alloc_pitch(self.encoder_width, roundup(self.encoder_height+MAX(0, VADJUST), 64)*3//plane_size_div, 16)
             log("CUDA Output Buffer=%#x, pitch=%s", int(self.cudaOutputBuffer), self.outputPitch)
             #allocate input buffer on host:
             self.inputBuffer = driver.pagelocked_zeros(self.inputPitch*self.input_height, dtype=numpy.byte)
@@ -1936 +1939 @@
 
         csc_start = time.time()
         args = (self.cudaInputBuffer, numpy.int32(in_w), numpy.int32(in_h), numpy.int32(stride),
-               self.cudaOutputBuffer, numpy.int32(self.encoder_width), numpy.int32(self.encoder_height), numpy.int32(self.outputPitch),
+               self.cudaOutputBuffer, numpy.int32(self.encoder_width), numpy.int32(self.encoder_height),
+               numpy.int32(self.outputPitch), numpy.int32(MAX(1, self.encoder_height+VADJUST)),
                numpy.int32(w), numpy.int32(h))
-        log("calling %s%s with block=%s, grid=%s", self.kernel, args, (blockw,blockh,1), (gridw, gridh))
+        log.info("calling %s%s with block=%s, grid=%s (VADJUST=%s)", self.kernel, args, (blockw,blockh,1), (gridw, gridh), VADJUST)
         self.kernel(*args, block=(blockw,blockh,1), grid=(gridw, gridh))
         csc_end = time.time()
         log("compress_image(..) kernel %s executed - CSC took %.1f ms", self.kernel_name, (csc_end - csc_start)*1000.0)
@@ -2001 +2005 @@
                 r = self.functionList.nvEncEncodePicture(self.context, &picParams)
             raiseNVENC(r, "error during picture encoding")
             encode_end = time.time()
-            log("compress_image(..) encoded in %.1f ms, info=%s", (encode_end-csc_end)*1000.0, self.get_info())
+            log.info("compress_image(..) encoded in %.1f ms, info=%s", (encode_end-csc_end)*1000.0, self.get_info())
 
             #lock output buffer:
             memset(&lockOutputBuffer, 0, sizeof(NV_ENC_LOCK_BITSTREAM))