--- /dev/null
+/* GStreamer
+ * Copyright (C) 2010 David Schleef <ds@schleef.org>
+ * Copyright (C) 2010 Sebastian Dröge <sebastian.droege@collabora.co.uk>
+ * Copyright (C) 2019 Seungha Yang <seungha.yang@navercorp.com>
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Library General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Library General Public License for more details.
+ *
+ * You should have received a copy of the GNU Library General Public
+ * License along with this library; if not, write to the
+ * Free Software Foundation, Inc., 51 Franklin St, Fifth Floor,
+ * Boston, MA 02110-1301, USA.
+ */
+
+/**
+ * SECTION:cudaconverter
+ * @title: GstCudaConverter
+ * @short_description: Generic video conversion using CUDA
+ *
+ * This object is used to convert video frames from one format to another.
+ * The object can perform conversion of:
+ *
+ * * video format
+ * * video colorspace
+ * * video size
+ */
+
+/**
+ * TODO:
+ * * Add more interpolation method and make it selectable,
+ * currently default bi-linear interpolation only
+ * * Add fast-path for conversion like videoconvert
+ * * Full colorimetiry and chroma-siting support
+ * * cropping, and x, y position support
+ */
+
+#ifdef HAVE_CONFIG_H
+#include "config.h"
+#endif
+
+#include "cuda-converter.h"
+#include "gstcudautils.h"
+#include "gstcudaloader.h"
+#include "gstcudanvrtc.h"
+#include <string.h>
+
+#define CUDA_BLOCK_X 16
+#define CUDA_BLOCK_Y 16
+#define DIV_UP(size,block) (((size) + ((block) - 1)) / (block))
+
+static gboolean cuda_converter_lookup_path (GstCudaConverter * convert);
+
+#ifndef GST_DISABLE_GST_DEBUG
+#define GST_CAT_DEFAULT ensure_debug_category()
+static GstDebugCategory *
+ensure_debug_category (void)
+{
+ static gsize cat_gonce = 0;
+
+ if (g_once_init_enter (&cat_gonce)) {
+ gsize cat_done;
+
+ cat_done = (gsize) _gst_debug_category_new ("cuda-converter", 0,
+ "cuda-converter object");
+
+ g_once_init_leave (&cat_gonce, cat_done);
+ }
+
+ return (GstDebugCategory *) cat_gonce;
+}
+#else
+#define ensure_debug_category()
+#endif
+
+#define GST_CUDA_KERNEL_FUNC "gst_cuda_kernel_func"
+
+#define GST_CUDA_KERNEL_FUNC_TO_Y444 "gst_cuda_kernel_func_to_y444"
+
+#define GST_CUDA_KERNEL_FUNC_Y444_TO_YUV "gst_cuda_kernel_func_y444_to_yuv"
+
+#define GST_CUDA_KERNEL_FUNC_TO_ARGB "gst_cuda_kernel_func_to_argb"
+
+#define GST_CUDA_KERNEL_FUNC_SCALE_RGB "gst_cuda_kernel_func_scale_rgb"
+
+/* *INDENT-OFF* */
+/**
+ * read_chroma:
+ * @tex1: a CUDA texture object representing a semi-planar chroma plane
+ * @tex2: dummy object
+ * @x: the x coordinate to read data from @tex1
+ * @y: the y coordinate to read data from @tex1
+ *
+ * Returns: a #ushort2 vector representing both chroma pixel values
+ */
+static const gchar READ_CHROMA_FROM_SEMI_PLANAR[] =
+"__device__ ushort2\n"
+"read_chroma (cudaTextureObject_t tex1, cudaTextureObject_t tex2, \n"
+" float x, float y)\n"
+"{\n"
+" return tex2D<ushort2>(tex1, x, y);\n"
+"}";
+
+/**
+ * read_chroma:
+ * @tex1: a CUDA texture object representing a chroma planar plane
+ * @tex2: a CUDA texture object representing the orher planar plane
+ * @x: the x coordinate to read data from @tex1 and @tex2
+ * @y: the y coordinate to read data from @tex1 and @tex2
+ *
+ * Returns: a #ushort2 vector representing both chroma pixel values
+ */
+static const gchar READ_CHROMA_FROM_PLANAR[] =
+"__device__ ushort2\n"
+"read_chroma (cudaTextureObject_t tex1, cudaTextureObject_t tex2, \n"
+" float x, float y)\n"
+"{\n"
+" unsigned short u, v;\n"
+" u = tex2D<unsigned short>(tex1, x, y);\n"
+" v = tex2D<unsigned short>(tex2, x, y);\n"
+" return make_ushort2(u, v);\n"
+"}";
+
+/**
+ * write_chroma:
+ * @dst1: a CUDA global memory pointing to a semi-planar chroma plane
+ * @dst2: dummy
+ * @u: a pixel value to write @dst1
+ * @v: a pixel value to write @dst1
+ * @x: the x coordinate to wrtie data into @tex1
+ * @x: the y coordinate to wrtie data into @tex1
+ * @pstride: the pixel stride of @dst1
+ * @mask: bitmask to be applied to high bitdepth plane
+ *
+ * Write @u and @v pixel value to @dst1 semi-planar plane
+ */
+static const gchar WRITE_CHROMA_TO_SEMI_PLANAR[] =
+"__device__ void\n"
+"write_chroma (unsigned char *dst1, unsigned char *dst2, unsigned short u,\n"
+" unsigned short v, int x, int y, int pstride, int stride, int mask)\n"
+"{\n"
+" if (OUT_DEPTH > 8) {\n"
+" *(unsigned short *)&dst1[x * pstride + y * stride] = (u & mask);\n"
+" *(unsigned short *)&dst1[x * pstride + 2 + y * stride] = (v & mask);\n"
+" } else {\n"
+" dst1[x * pstride + y * stride] = u;\n"
+" dst1[x * pstride + 1 + y * stride] = v;\n"
+" }\n"
+"}";
+
+/**
+ * write_chroma:
+ * @dst1: a CUDA global memory pointing to a planar chroma plane
+ * @dst2: a CUDA global memory pointing to a the other planar chroma plane
+ * @u: a pixel value to write @dst1
+ * @v: a pixel value to write @dst1
+ * @x: the x coordinate to wrtie data into @tex1
+ * @x: the y coordinate to wrtie data into @tex1
+ * @pstride: the pixel stride of @dst1
+ * @mask: bitmask to be applied to high bitdepth plane
+ *
+ * Write @u and @v pixel value into @dst1 and @dst2 planar planes
+ */
+static const gchar WRITE_CHROMA_TO_PLANAR[] =
+"__device__ void\n"
+"write_chroma (unsigned char *dst1, unsigned char *dst2, unsigned short u,\n"
+" unsigned short v, int x, int y, int pstride, int stride, int mask)\n"
+"{\n"
+" if (OUT_DEPTH > 8) {\n"
+" *(unsigned short *)&dst1[x * pstride + y * stride] = (u & mask);\n"
+" *(unsigned short *)&dst2[x * pstride + y * stride] = (v & mask);\n"
+" } else {\n"
+" dst1[x * pstride + y * stride] = u;\n"
+" dst2[x * pstride + y * stride] = v;\n"
+" }\n"
+"}";
+
+/* CUDA kernel source for from YUV to YUV conversion and scale */
+static const gchar templ_YUV_TO_YUV[] =
+"extern \"C\"{\n"
+"__constant__ float SCALE_H = %f;\n"
+"__constant__ float SCALE_V = %f;\n"
+"__constant__ float CHROMA_SCALE_H = %f;\n"
+"__constant__ float CHROMA_SCALE_V = %f;\n"
+"__constant__ int WIDTH = %d;\n"
+"__constant__ int HEIGHT = %d;\n"
+"__constant__ int CHROMA_WIDTH = %d;\n"
+"__constant__ int CHROMA_HEIGHT = %d;\n"
+"__constant__ int IN_DEPTH = %d;\n"
+"__constant__ int OUT_DEPTH = %d;\n"
+"__constant__ int PSTRIDE = %d;\n"
+"__constant__ int CHROMA_PSTRIDE = %d;\n"
+"__constant__ int IN_SHIFT = %d;\n"
+"__constant__ int OUT_SHIFT = %d;\n"
+"__constant__ int MASK = %d;\n"
+"__constant__ int SWAP_UV = %d;\n"
+"\n"
+"__device__ unsigned short\n"
+"do_scale_pixel (unsigned short val) \n"
+"{\n"
+" unsigned int diff;\n"
+" if (OUT_DEPTH > IN_DEPTH) {\n"
+" diff = OUT_DEPTH - IN_DEPTH;\n"
+" return (val << diff) | (val >> (IN_DEPTH - diff));\n"
+" } else if (IN_DEPTH > OUT_DEPTH) {\n"
+" return val >> (IN_DEPTH - OUT_DEPTH);\n"
+" }\n"
+" return val;\n"
+"}\n"
+"\n"
+/* __device__ ushort2
+ * read_chroma (cudaTextureObject_t tex1, cudaTextureObject_t tex2, float x, float y);
+ */
+"%s\n"
+"\n"
+/* __device__ void
+ * write_chroma (unsigned char *dst1, unsigned char *dst2, unsigned short u,
+ * unsigned short v, int x, int y, int pstride, int stride, int mask);
+ */
+"%s\n"
+"\n"
+"__global__ void\n"
+GST_CUDA_KERNEL_FUNC
+"(cudaTextureObject_t tex0, cudaTextureObject_t tex1, cudaTextureObject_t tex2,\n"
+" unsigned char *dst0, unsigned char *dst1, unsigned char *dst2,\n"
+" int stride)\n"
+"{\n"
+" int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
+" int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+" if (x_pos < WIDTH && y_pos < HEIGHT) {\n"
+" float src_xpos = SCALE_H * x_pos;\n"
+" float src_ypos = SCALE_V * y_pos;\n"
+" unsigned short y = tex2D<unsigned short>(tex0, src_xpos, src_ypos);\n"
+" y = y >> IN_SHIFT;\n"
+" y = do_scale_pixel (y);\n"
+" y = y << OUT_SHIFT;\n"
+" if (OUT_DEPTH > 8) {\n"
+" *(unsigned short *)&dst0[x_pos * PSTRIDE + y_pos * stride] = (y & MASK);\n"
+" } else {\n"
+" dst0[x_pos * PSTRIDE + y_pos * stride] = y;\n"
+" }\n"
+" }\n"
+" if (x_pos < CHROMA_WIDTH && y_pos < CHROMA_HEIGHT) {\n"
+" float src_xpos = CHROMA_SCALE_H * x_pos;\n"
+" float src_ypos = CHROMA_SCALE_V * y_pos;\n"
+" unsigned short u, v;\n"
+" ushort2 uv = read_chroma (tex1, tex2, src_xpos, src_ypos);\n"
+" u = uv.x;\n"
+" v = uv.y;\n"
+" u = u >> IN_SHIFT;\n"
+" v = v >> IN_SHIFT;\n"
+" u = do_scale_pixel (u);\n"
+" v = do_scale_pixel (v);\n"
+" u = u << OUT_SHIFT;\n"
+" v = v << OUT_SHIFT;\n"
+" if (SWAP_UV) {\n"
+" unsigned short tmp = u;\n"
+" u = v;\n"
+" v = tmp;\n"
+" }\n"
+" write_chroma (dst1,\n"
+" dst2, u, v, x_pos, y_pos, CHROMA_PSTRIDE, stride, MASK);\n"
+" }\n"
+"}\n"
+"\n"
+"}";
+
+/* CUDA kernel source for from YUV to RGB conversion and scale */
+static const gchar templ_YUV_TO_RGB[] =
+"extern \"C\"{\n"
+"__constant__ float offset[3] = {%f, %f, %f};\n"
+"__constant__ float rcoeff[3] = {%f, %f, %f};\n"
+"__constant__ float gcoeff[3] = {%f, %f, %f};\n"
+"__constant__ float bcoeff[3] = {%f, %f, %f};\n"
+"\n"
+"__constant__ float SCALE_H = %f;\n"
+"__constant__ float SCALE_V = %f;\n"
+"__constant__ float CHROMA_SCALE_H = %f;\n"
+"__constant__ float CHROMA_SCALE_V = %f;\n"
+"__constant__ int WIDTH = %d;\n"
+"__constant__ int HEIGHT = %d;\n"
+"__constant__ int CHROMA_WIDTH = %d;\n"
+"__constant__ int CHROMA_HEIGHT = %d;\n"
+"__constant__ int IN_DEPTH = %d;\n"
+"__constant__ int OUT_DEPTH = %d;\n"
+"__constant__ int PSTRIDE = %d;\n"
+"__constant__ int CHROMA_PSTRIDE = %d;\n"
+"__constant__ int IN_SHIFT = %d;\n"
+"__constant__ int OUT_SHIFT = %d;\n"
+"__constant__ int MASK = %d;\n"
+"__constant__ int SWAP_UV = %d;\n"
+"__constant__ int MAX_IN_VAL = %d;\n"
+"__constant__ int R_IDX = %d;\n"
+"__constant__ int G_IDX = %d;\n"
+"__constant__ int B_IDX = %d;\n"
+"__constant__ int A_IDX = %d;\n"
+"__constant__ int X_IDX = %d;\n"
+"\n"
+"__device__ unsigned short\n"
+"do_scale_pixel (unsigned short val) \n"
+"{\n"
+" unsigned int diff;\n"
+" if (OUT_DEPTH > IN_DEPTH) {\n"
+" diff = OUT_DEPTH - IN_DEPTH;\n"
+" return (val << diff) | (val >> (IN_DEPTH - diff));\n"
+" } else if (IN_DEPTH > OUT_DEPTH) {\n"
+" return val >> (IN_DEPTH - OUT_DEPTH);\n"
+" }\n"
+" return val;\n"
+"}\n"
+"\n"
+"__device__ float\n"
+"dot(float3 val, float *coeff)\n"
+"{\n"
+" return val.x * coeff[0] + val.y * coeff[1] + val.z * coeff[2];\n"
+"}\n"
+"\n"
+"__device__ uint3\n"
+"yuv_to_rgb (unsigned short y, unsigned short u, unsigned short v, unsigned int max_val)\n"
+"{\n"
+" float3 yuv = make_float3 (y, u, v);\n"
+" uint3 rgb;\n"
+" rgb.x = max ((unsigned int)(dot (yuv, rcoeff) + offset[0]), 0);\n"
+" rgb.y = max ((unsigned int)(dot (yuv, gcoeff) + offset[1]), 0);\n"
+" rgb.z = max ((unsigned int)(dot (yuv, bcoeff) + offset[2]), 0);\n"
+" rgb.x = min (rgb.x, max_val);\n"
+" rgb.y = min (rgb.y, max_val);\n"
+" rgb.z = min (rgb.z, max_val);\n"
+" return rgb;\n"
+"}\n"
+"\n"
+/* __device__ ushort2
+ * read_chroma (cudaTextureObject_t tex1, cudaTextureObject_t tex2, float x, float y);
+ */
+"%s\n"
+"\n"
+"__global__ void\n"
+GST_CUDA_KERNEL_FUNC
+"(cudaTextureObject_t tex0, cudaTextureObject_t tex1, cudaTextureObject_t tex2,\n"
+" unsigned char *dstRGB, int stride)\n"
+"{\n"
+" int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
+" int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+" if (x_pos < WIDTH && y_pos < HEIGHT) {\n"
+" float src_xpos = SCALE_H * x_pos;\n"
+" float src_ypos = SCALE_V * y_pos;\n"
+" unsigned short y = tex2D<unsigned short>(tex0, src_xpos, src_ypos);\n"
+" ushort2 uv;\n"
+" unsigned short u, v;\n"
+" uint3 rgb;\n"
+" unsigned int clip_max = MAX_IN_VAL;\n"
+" src_xpos = CHROMA_SCALE_H * x_pos;\n"
+" src_ypos = CHROMA_SCALE_V * y_pos;\n"
+" uv = read_chroma (tex1, tex2, src_xpos, src_ypos);\n"
+" u = uv.x;\n"
+" v = uv.y;\n"
+" y = y >> IN_SHIFT;\n"
+" u = u >> IN_SHIFT;\n"
+" v = v >> IN_SHIFT;\n"
+" if (SWAP_UV) {\n"
+" unsigned short tmp = u;\n"
+" u = v;\n"
+" v = tmp;\n"
+" }\n"
+ /* conversion matrix is scaled to higher bitdepth between in/out formats */
+" if (OUT_DEPTH > IN_DEPTH) {\n"
+" y = do_scale_pixel (y);\n"
+" u = do_scale_pixel (u);\n"
+" v = do_scale_pixel (v);\n"
+" clip_max = MASK;\n"
+" }"
+" rgb = yuv_to_rgb (y, u, v, clip_max);\n"
+" if (OUT_DEPTH < IN_DEPTH) {\n"
+" rgb.x = do_scale_pixel (rgb.x);\n"
+" rgb.y = do_scale_pixel (rgb.y);\n"
+" rgb.z = do_scale_pixel (rgb.z);\n"
+" }"
+" if (OUT_DEPTH > 8) {\n"
+" unsigned int packed_rgb = 0;\n"
+ /* A is always MSB, we support only little endian system */
+" packed_rgb = 0xc000 << 16;\n"
+" packed_rgb |= (rgb.x << (30 - (R_IDX * 10)));\n"
+" packed_rgb |= (rgb.y << (30 - (G_IDX * 10)));\n"
+" packed_rgb |= (rgb.z << (30 - (B_IDX * 10)));\n"
+" *(unsigned int *)&dstRGB[x_pos * PSTRIDE + y_pos * stride] = packed_rgb;\n"
+" } else {\n"
+" dstRGB[x_pos * PSTRIDE + R_IDX + y_pos * stride] = (unsigned char) rgb.x;\n"
+" dstRGB[x_pos * PSTRIDE + G_IDX + y_pos * stride] = (unsigned char) rgb.y;\n"
+" dstRGB[x_pos * PSTRIDE + B_IDX + y_pos * stride] = (unsigned char) rgb.z;\n"
+" if (A_IDX >= 0 || X_IDX >= 0)\n"
+" dstRGB[x_pos * PSTRIDE + A_IDX + y_pos * stride] = 0xff;\n"
+" }\n"
+" }\n"
+"}\n"
+"\n"
+"}";
+
+/**
+ * GST_CUDA_KERNEL_FUNC_TO_ARGB:
+ * @srcRGB: a CUDA global memory containing a RGB image
+ * @dstRGB: a CUDA global memory to store unpacked ARGB image
+ * @width: the width of @srcRGB and @dstRGB
+ * @height: the height of @srcRGB and @dstRGB
+ * @src_stride: the stride of @srcRGB
+ * @src_pstride: the pixel stride of @srcRGB
+ * @dst_stride: the stride of @dstRGB
+ * @r_idx: the index of red component of @srcRGB
+ * @g_idx: the index of green component of @srcRGB
+ * @b_idx: the index of blue component of @srcRGB
+ * @a_idx: the index of alpha component of @srcRGB
+ *
+ * Unpack a RGB image from @srcRGB and write the unpacked data into @dstRGB
+ */
+static const gchar unpack_to_ARGB[] =
+"__global__ void\n"
+GST_CUDA_KERNEL_FUNC_TO_ARGB
+"(unsigned char *srcRGB, unsigned char *dstRGB, int width, int height,\n"
+" int src_stride, int src_pstride, int dst_stride,\n"
+" int r_idx, int g_idx, int b_idx, int a_idx)\n"
+"{\n"
+" int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
+" int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+" if (x_pos < width && y_pos < height) {\n"
+" if (a_idx >= 0) {\n"
+" dstRGB[x_pos * 4 + y_pos * dst_stride] =\n"
+" srcRGB[x_pos * src_pstride + a_idx + y_pos * src_stride];\n"
+" } else {\n"
+" dstRGB[x_pos * 4 + y_pos * dst_stride] = 0xff;\n"
+" }\n"
+" dstRGB[x_pos * 4 + 1 + y_pos * dst_stride] =\n"
+" srcRGB[x_pos * src_pstride + r_idx + y_pos * src_stride];\n"
+" dstRGB[x_pos * 4 + 2 + y_pos * dst_stride] =\n"
+" srcRGB[x_pos * src_pstride + g_idx + y_pos * src_stride];\n"
+" dstRGB[x_pos * 4 + 3 + y_pos * dst_stride] =\n"
+" srcRGB[x_pos * src_pstride + b_idx + y_pos * src_stride];\n"
+" }\n"
+"}\n";
+
+/**
+ * GST_CUDA_KERNEL_FUNC_TO_ARGB:
+ * @srcRGB: a CUDA global memory containing a RGB image
+ * @dstRGB: a CUDA global memory to store unpacked ARGB64 image
+ * @width: the width of @srcRGB and @dstRGB
+ * @height: the height of @srcRGB and @dstRGB
+ * @src_stride: the stride of @srcRGB
+ * @src_pstride: the pixel stride of @srcRGB
+ * @dst_stride: the stride of @dstRGB
+ * @r_idx: the index of red component of @srcRGB
+ * @g_idx: the index of green component of @srcRGB
+ * @b_idx: the index of blue component of @srcRGB
+ * @a_idx: the index of alpha component of @srcRGB
+ *
+ * Unpack a RGB image from @srcRGB and write the unpacked data into @dstRGB
+ */
+static const gchar unpack_to_ARGB64[] =
+"__global__ void\n"
+GST_CUDA_KERNEL_FUNC_TO_ARGB
+"(unsigned char *srcRGB, unsigned char *dstRGB, int width, int height,\n"
+" int src_stride, int src_pstride, int dst_stride,\n"
+" int r_idx, int g_idx, int b_idx, int a_idx)\n"
+"{\n"
+" int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
+" int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+" if (x_pos < width && y_pos < height) {\n"
+" unsigned short a, r, g, b;\n"
+" unsigned int read_val;\n"
+" read_val = *(unsigned int *)&srcRGB[x_pos * src_pstride + y_pos * src_stride];\n"
+" a = (read_val >> 30) & 0x03;\n"
+" a = (a << 14) | (a << 12) | (a << 10) | (a << 8) | (a << 6) | (a << 4) | (a << 2) | (a << 0);\n"
+" r = ((read_val >> (30 - (r_idx * 10))) & 0x3ff);\n"
+" r = (r << 6) | (r >> 4);\n"
+" g = ((read_val >> (30 - (g_idx * 10))) & 0x3ff);\n"
+" g = (g << 6) | (g >> 4);\n"
+" b = ((read_val >> (30 - (b_idx * 10))) & 0x3ff);\n"
+" b = (b << 6) | (b >> 4);\n"
+" *(unsigned short *)&dstRGB[x_pos * 8 + y_pos * dst_stride] = 0xffff;\n"
+" *(unsigned short *)&dstRGB[x_pos * 8 + 2 + y_pos * dst_stride] = r;\n"
+" *(unsigned short *)&dstRGB[x_pos * 8 + 4 + y_pos * dst_stride] = g;\n"
+" *(unsigned short *)&dstRGB[x_pos * 8 + 6 + y_pos * dst_stride] = b;\n"
+" }\n"
+"}\n";
+
+/* CUDA kernel source for from RGB to YUV conversion and scale */
+static const gchar templ_RGB_TO_YUV[] =
+"extern \"C\"{\n"
+"__constant__ float offset[3] = {%f, %f, %f};\n"
+"__constant__ float ycoeff[3] = {%f, %f, %f};\n"
+"__constant__ float ucoeff[3] = {%f, %f, %f};\n"
+"__constant__ float vcoeff[3] = {%f, %f, %f};\n"
+"\n"
+"__constant__ float SCALE_H = %f;\n"
+"__constant__ float SCALE_V = %f;\n"
+"__constant__ float CHROMA_SCALE_H = %f;\n"
+"__constant__ float CHROMA_SCALE_V = %f;\n"
+"__constant__ int WIDTH = %d;\n"
+"__constant__ int HEIGHT = %d;\n"
+"__constant__ int CHROMA_WIDTH = %d;\n"
+"__constant__ int CHROMA_HEIGHT = %d;\n"
+"__constant__ int IN_DEPTH = %d;\n"
+"__constant__ int OUT_DEPTH = %d;\n"
+"__constant__ int PSTRIDE = %d;\n"
+"__constant__ int CHROMA_PSTRIDE = %d;\n"
+"__constant__ int IN_SHIFT = %d;\n"
+"__constant__ int OUT_SHIFT = %d;\n"
+"__constant__ int MASK = %d;\n"
+"__constant__ int SWAP_UV = %d;\n"
+"\n"
+"__device__ unsigned short\n"
+"do_scale_pixel (unsigned short val) \n"
+"{\n"
+" unsigned int diff;\n"
+" if (OUT_DEPTH > IN_DEPTH) {\n"
+" diff = OUT_DEPTH - IN_DEPTH;\n"
+" return (val << diff) | (val >> (IN_DEPTH - diff));\n"
+" } else if (IN_DEPTH > OUT_DEPTH) {\n"
+" return val >> (IN_DEPTH - OUT_DEPTH);\n"
+" }\n"
+" return val;\n"
+"}\n"
+"\n"
+"__device__ float\n"
+"dot(float3 val, float *coeff)\n"
+"{\n"
+" return val.x * coeff[0] + val.y * coeff[1] + val.z * coeff[2];\n"
+"}\n"
+"\n"
+"__device__ uint3\n"
+"rgb_to_yuv (unsigned short r, unsigned short g, unsigned short b,\n"
+" unsigned int max_val)\n"
+"{\n"
+" float3 rgb = make_float3 (r, g, b);\n"
+" uint3 yuv;\n"
+" yuv.x = max ((unsigned int)(dot (rgb, ycoeff) + offset[0]), 0);\n"
+" yuv.y = max ((unsigned int)(dot (rgb, ucoeff) + offset[1]), 0);\n"
+" yuv.z = max ((unsigned int)(dot (rgb, vcoeff) + offset[2]), 0);\n"
+" yuv.x = min (yuv.x, max_val);\n"
+" yuv.y = min (yuv.y, max_val);\n"
+" yuv.z = min (yuv.z, max_val);\n"
+" return yuv;\n"
+"}\n"
+"\n"
+/* __global__ void
+ * GST_CUDA_KERNEL_FUNC_TO_ARGB
+ */
+"%s\n"
+"\n"
+/* __device__ ushort2
+ * read_chroma (cudaTextureObject_t tex1, cudaTextureObject_t tex2, float x, float y);
+ */
+"%s\n"
+"\n"
+/* __device__ void
+ * write_chroma (unsigned char *dst1, unsigned char *dst2, unsigned short u,
+ * unsigned short v, int x, int y, int pstride, int stride, int mask);
+ */
+"%s\n"
+"\n"
+"__global__ void\n"
+GST_CUDA_KERNEL_FUNC_TO_Y444
+"(cudaTextureObject_t srcRGB, unsigned char *dstY, int y_stride,\n"
+" unsigned char *dstU, int u_stride, unsigned char *dstV, int v_stride,\n"
+" int width, int height, int dst_pstride, int in_depth)\n"
+"{\n"
+" int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
+" int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+" if (x_pos < width && y_pos < height) {\n"
+" ushort4 argb = tex2D<ushort4>(srcRGB, x_pos, y_pos);\n"
+" uint3 yuv;\n"
+" yuv = rgb_to_yuv (argb.y, argb.z, argb.w, (1 << in_depth) - 1);\n"
+" if (in_depth > 8) {\n"
+" *(unsigned short *)&dstY[x_pos * dst_pstride + y_pos * y_stride] = yuv.x;\n"
+" *(unsigned short *)&dstU[x_pos * dst_pstride + y_pos * u_stride] = yuv.y;\n"
+" *(unsigned short *)&dstV[x_pos * dst_pstride + y_pos * v_stride] = yuv.z;\n"
+" } else {\n"
+" dstY[x_pos * dst_pstride + y_pos * y_stride] = yuv.x;\n"
+" dstU[x_pos * dst_pstride + y_pos * u_stride] = yuv.y;\n"
+" dstV[x_pos * dst_pstride + y_pos * v_stride] = yuv.z;\n"
+" }\n"
+" }\n"
+"}\n"
+"\n"
+"__global__ void\n"
+GST_CUDA_KERNEL_FUNC_Y444_TO_YUV
+"(cudaTextureObject_t tex0, cudaTextureObject_t tex1, cudaTextureObject_t tex2,\n"
+" unsigned char *dst0, unsigned char *dst1, unsigned char *dst2,\n"
+" int stride)\n"
+"{\n"
+" int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
+" int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+" if (x_pos < WIDTH && y_pos < HEIGHT) {\n"
+" float src_xpos = SCALE_H * x_pos;\n"
+" float src_ypos = SCALE_V * y_pos;\n"
+" unsigned short y = tex2D<unsigned short>(tex0, src_xpos, src_ypos);\n"
+" y = y >> IN_SHIFT;\n"
+" y = do_scale_pixel (y);\n"
+" y = y << OUT_SHIFT;\n"
+" if (OUT_DEPTH > 8) {\n"
+" *(unsigned short *)&dst0[x_pos * PSTRIDE + y_pos * stride] = (y & MASK);\n"
+" } else {\n"
+" dst0[x_pos * PSTRIDE + y_pos * stride] = y;\n"
+" }\n"
+" }\n"
+" if (x_pos < CHROMA_WIDTH && y_pos < CHROMA_HEIGHT) {\n"
+" float src_xpos = CHROMA_SCALE_H * x_pos;\n"
+" float src_ypos = CHROMA_SCALE_V * y_pos;\n"
+" unsigned short u, v;\n"
+" ushort2 uv;\n"
+" uv = read_chroma (tex1, tex2, src_xpos, src_ypos);\n"
+" u = uv.x;\n"
+" v = uv.y;\n"
+" u = u >> IN_SHIFT;\n"
+" v = v >> IN_SHIFT;\n"
+" u = do_scale_pixel (u);\n"
+" v = do_scale_pixel (v);\n"
+" u = u << OUT_SHIFT;\n"
+" v = v << OUT_SHIFT;\n"
+" if (SWAP_UV) {\n"
+" unsigned short tmp = u;\n"
+" u = v;\n"
+" v = tmp;\n"
+" }\n"
+" write_chroma (dst1,\n"
+" dst2, u, v, x_pos, y_pos, CHROMA_PSTRIDE, stride, MASK);\n"
+" }\n"
+"}\n"
+"\n"
+"}";
+
+/* CUDA kernel source for from RGB to RGB conversion and scale */
+static const gchar templ_RGB_to_RGB[] =
+"extern \"C\"{\n"
+"__constant__ float SCALE_H = %f;\n"
+"__constant__ float SCALE_V = %f;\n"
+"__constant__ int WIDTH = %d;\n"
+"__constant__ int HEIGHT = %d;\n"
+"__constant__ int IN_DEPTH = %d;\n"
+"__constant__ int OUT_DEPTH = %d;\n"
+"__constant__ int PSTRIDE = %d;\n"
+"__constant__ int R_IDX = %d;\n"
+"__constant__ int G_IDX = %d;\n"
+"__constant__ int B_IDX = %d;\n"
+"__constant__ int A_IDX = %d;\n"
+"__constant__ int X_IDX = %d;\n"
+"\n"
+"__device__ unsigned short\n"
+"do_scale_pixel (unsigned short val) \n"
+"{\n"
+" unsigned int diff;\n"
+" if (OUT_DEPTH > IN_DEPTH) {\n"
+" diff = OUT_DEPTH - IN_DEPTH;\n"
+" return (val << diff) | (val >> (IN_DEPTH - diff));\n"
+" } else if (IN_DEPTH > OUT_DEPTH) {\n"
+" return val >> (IN_DEPTH - OUT_DEPTH);\n"
+" }\n"
+" return val;\n"
+"}\n"
+"\n"
+/* __global__ void
+ * GST_CUDA_KERNEL_FUNC_TO_ARGB
+ */
+"%s\n"
+"\n"
+/* convert ARGB or ARGB64 to other RGB formats with scale */
+"__global__ void\n"
+GST_CUDA_KERNEL_FUNC_SCALE_RGB
+"(cudaTextureObject_t srcRGB, unsigned char *dstRGB, int dst_stride)\n"
+"{\n"
+" int x_pos = blockIdx.x * blockDim.x + threadIdx.x;\n"
+" int y_pos = blockIdx.y * blockDim.y + threadIdx.y;\n"
+" if (x_pos < WIDTH && y_pos < HEIGHT) {\n"
+" float src_xpos = SCALE_H * x_pos;\n"
+" float src_ypos = SCALE_V * y_pos;\n"
+" ushort4 argb = tex2D<ushort4>(srcRGB, src_xpos, src_ypos);\n"
+" argb.x = do_scale_pixel(argb.x);\n"
+" argb.y = do_scale_pixel(argb.y);\n"
+" argb.z = do_scale_pixel(argb.z);\n"
+" argb.w = do_scale_pixel(argb.w);\n"
+ /* FIXME: RGB10A2_LE or BGR10A2_LE only */
+" if (OUT_DEPTH > 8) {\n"
+" unsigned int packed_rgb = 0;\n"
+" unsigned int a, r, g, b;"
+" a = (argb.x >> 8) & 0x3;\n"
+" r = argb.y & 0x3ff;\n"
+" g = argb.z & 0x3ff;\n"
+" b = argb.w & 0x3ff;\n"
+ /* A is always MSB, we support only little endian system */
+" packed_rgb = a << 30;\n"
+" packed_rgb |= (r << (30 - (R_IDX * 10)));\n"
+" packed_rgb |= (g << (30 - (G_IDX * 10)));\n"
+" packed_rgb |= (b << (30 - (B_IDX * 10)));\n"
+" *(unsigned int *)&dstRGB[x_pos * 4 + y_pos * dst_stride] = packed_rgb;\n"
+" } else {\n"
+" if (A_IDX >= 0) {\n"
+" argb.x = do_scale_pixel(argb.x);\n"
+" dstRGB[x_pos * PSTRIDE + A_IDX + y_pos * dst_stride] = argb.x;\n"
+" } else if (X_IDX >= 0) {\n"
+" dstRGB[x_pos * PSTRIDE + X_IDX + y_pos * dst_stride] = 0xff;\n"
+" }\n"
+" dstRGB[x_pos * PSTRIDE + R_IDX + y_pos * dst_stride] = argb.y;\n"
+" dstRGB[x_pos * PSTRIDE + G_IDX + y_pos * dst_stride] = argb.z;\n"
+" dstRGB[x_pos * PSTRIDE + B_IDX + y_pos * dst_stride] = argb.w;\n"
+" }\n"
+" }\n"
+"}\n"
+"\n"
+"}";
+/* *INDENT-ON* */
+
+typedef struct
+{
+ gint R;
+ gint G;
+ gint B;
+ gint A;
+ gint X;
+} GstCudaRGBOrder;
+
+typedef struct
+{
+ CUdeviceptr device_ptr;
+ gsize cuda_stride;
+} GstCudaStageBuffer;
+
+#define CONVERTER_MAX_NUM_FUNC 4
+
+struct _GstCudaConverter
+{
+ GstVideoInfo in_info;
+ GstVideoInfo out_info;
+ gboolean keep_size;
+
+ gint texture_alignment;
+
+ GstCudaContext *cuda_ctx;
+ CUmodule cuda_module;
+ CUfunction kernel_func[CONVERTER_MAX_NUM_FUNC];
+ const gchar *func_names[CONVERTER_MAX_NUM_FUNC];
+ gchar *kernel_source;
+ gchar *ptx;
+ GstCudaStageBuffer fallback_buffer[GST_VIDEO_MAX_PLANES];
+
+ gboolean (*convert) (GstCudaConverter * convert, const GstCudaMemory * src,
+ GstVideoInfo * in_info, GstCudaMemory * dst, GstVideoInfo * out_info,
+ CUstream cuda_stream);
+
+ const CUdeviceptr src;
+ GstVideoInfo *cur_in_info;
+
+ CUdeviceptr dest;
+ GstVideoInfo *cur_out_info;
+
+ /* rgb to {rgb, yuv} only */
+ GstCudaRGBOrder in_rgb_order;
+ GstCudaStageBuffer unpack_surface;
+ GstCudaStageBuffer y444_surface[GST_VIDEO_MAX_PLANES];
+};
+
+#define LOAD_CUDA_FUNC(module,func,name) G_STMT_START { \
+ if (!gst_cuda_result (CuModuleGetFunction (&(func), (module), name))) { \
+ GST_ERROR ("failed to get %s function", (name)); \
+ goto error; \
+ } \
+} G_STMT_END
+
+/**
+ * gst_cuda_converter_new:
+ * @in_info: a #GstVideoInfo
+ * @out_info: a #GstVideoInfo
+ * @cuda_ctx: (transfer none): a #GstCudaContext
+ *
+ * Create a new converter object to convert between @in_info and @out_info
+ * with @config.
+ *
+ * Returns: a #GstCudaConverter or %NULL if conversion is not possible.
+ */
+GstCudaConverter *
+gst_cuda_converter_new (GstVideoInfo * in_info, GstVideoInfo * out_info,
+ GstCudaContext * cuda_ctx)
+{
+ GstCudaConverter *convert;
+ gint i;
+
+ g_return_val_if_fail (in_info != NULL, NULL);
+ g_return_val_if_fail (out_info != NULL, NULL);
+ g_return_val_if_fail (cuda_ctx != NULL, NULL);
+ /* we won't ever do framerate conversion */
+ g_return_val_if_fail (in_info->fps_n == out_info->fps_n, NULL);
+ g_return_val_if_fail (in_info->fps_d == out_info->fps_d, NULL);
+ /* we won't ever do deinterlace */
+ g_return_val_if_fail (in_info->interlace_mode == out_info->interlace_mode,
+ NULL);
+
+ convert = g_new0 (GstCudaConverter, 1);
+
+ convert->in_info = *in_info;
+ convert->out_info = *out_info;
+
+ /* FIXME: should return kernel source */
+ if (!gst_cuda_context_push (cuda_ctx)) {
+ GST_ERROR ("cannot push context");
+ goto error;
+ }
+
+ if (!cuda_converter_lookup_path (convert))
+ goto error;
+
+ convert->ptx = gst_cuda_nvrtc_compile (convert->kernel_source);
+ if (!convert->ptx) {
+ GST_ERROR ("no PTX data to load");
+ goto error;
+ }
+
+ GST_TRACE ("compiled convert ptx \n%s", convert->ptx);
+
+ if (!gst_cuda_result (CuModuleLoadData (&convert->cuda_module, convert->ptx))) {
+ gst_cuda_context_pop (NULL);
+ GST_ERROR ("failed to load cuda module data");
+
+ goto error;
+ }
+
+ for (i = 0; i < CONVERTER_MAX_NUM_FUNC; i++) {
+ if (!convert->func_names[i])
+ break;
+
+ LOAD_CUDA_FUNC (convert->cuda_module, convert->kernel_func[i],
+ convert->func_names[i]);
+ GST_DEBUG ("kernel function \"%s\" loaded", convert->func_names[i]);
+ }
+
+ gst_cuda_context_pop (NULL);
+ convert->cuda_ctx = gst_object_ref (cuda_ctx);
+ convert->texture_alignment =
+ gst_cuda_context_get_texture_alignment (cuda_ctx);
+
+ g_free (convert->kernel_source);
+ g_free (convert->ptx);
+ convert->kernel_source = NULL;
+ convert->ptx = NULL;
+
+ return convert;
+
+error:
+ gst_cuda_context_pop (NULL);
+ gst_cuda_converter_free (convert);
+
+ return NULL;
+}
+
+/**
+ * gst_video_converter_free:
+ * @convert: a #GstCudaConverter
+ *
+ * Free @convert
+ */
+void
+gst_cuda_converter_free (GstCudaConverter * convert)
+{
+ g_return_if_fail (convert != NULL);
+
+ if (convert->cuda_ctx) {
+ if (gst_cuda_context_push (convert->cuda_ctx)) {
+ gint i;
+
+ if (convert->cuda_module) {
+ gst_cuda_result (CuModuleUnload (convert->cuda_module));
+ }
+
+ for (i = 0; i < GST_VIDEO_MAX_PLANES; i++) {
+ if (convert->fallback_buffer[i].device_ptr)
+ gst_cuda_result (CuMemFree (convert->fallback_buffer[i].device_ptr));
+ if (convert->y444_surface[i].device_ptr)
+ gst_cuda_result (CuMemFree (convert->y444_surface[i].device_ptr));
+ }
+
+ if (convert->unpack_surface.device_ptr)
+ gst_cuda_result (CuMemFree (convert->unpack_surface.device_ptr));
+
+ gst_cuda_context_pop (NULL);
+ }
+
+ gst_object_unref (convert->cuda_ctx);
+ }
+
+ g_free (convert->kernel_source);
+ g_free (convert->ptx);
+ g_free (convert);
+}
+
+/**
+ * gst_cuda_converter_frame:
+ * @convert: a #GstCudaConverter
+ * @src: a #GstCudaMemory
+ * @in_info: a #GstVideoInfo representing @src
+ * @dst: a #GstCudaMemory
+ * @out_info: a #GstVideoInfo representing @dst
+ * @cuda_stream: a #CUstream
+ *
+ * Convert the pixels of @src into @dest using @convert.
+ * Called without gst_cuda_context_push() and gst_cuda_context_pop() by caller
+ */
+gboolean
+gst_cuda_converter_frame (GstCudaConverter * convert, const GstCudaMemory * src,
+ GstVideoInfo * in_info, GstCudaMemory * dst, GstVideoInfo * out_info,
+ CUstream cuda_stream)
+{
+ gboolean ret;
+
+ g_return_val_if_fail (convert, FALSE);
+ g_return_val_if_fail (src, FALSE);
+ g_return_val_if_fail (in_info, FALSE);
+ g_return_val_if_fail (dst, FALSE);
+ g_return_val_if_fail (out_info, FALSE);
+
+ gst_cuda_context_push (convert->cuda_ctx);
+
+ ret = gst_cuda_converter_frame_unlocked (convert,
+ src, in_info, dst, out_info, cuda_stream);
+
+ gst_cuda_context_pop (NULL);
+
+ return ret;
+}
+
+/**
+ * gst_cuda_converter_frame_unlocked:
+ * @convert: a #GstCudaConverter
+ * @src: a #GstCudaMemory
+ * @in_info: a #GstVideoInfo representing @src
+ * @dst: a #GstCudaMemory
+ * @out_info: a #GstVideoInfo representing @dest
+ * @cuda_stream: a #CUstream
+ *
+ * Convert the pixels of @src into @dest using @convert.
+ * Caller should call this method after gst_cuda_context_push()
+ */
+gboolean
+gst_cuda_converter_frame_unlocked (GstCudaConverter * convert,
+ const GstCudaMemory * src, GstVideoInfo * in_info, GstCudaMemory * dst,
+ GstVideoInfo * out_info, CUstream cuda_stream)
+{
+ g_return_val_if_fail (convert, FALSE);
+ g_return_val_if_fail (src, FALSE);
+ g_return_val_if_fail (in_info, FALSE);
+ g_return_val_if_fail (dst, FALSE);
+ g_return_val_if_fail (out_info, FALSE);
+
+ return convert->convert (convert, src, in_info, dst, out_info, cuda_stream);
+}
+
+/* allocate fallback memory for texture alignment requirement */
+static gboolean
+convert_ensure_fallback_memory (GstCudaConverter * convert,
+ GstVideoInfo * info, guint plane)
+{
+ CUresult ret;
+ guint element_size = 8;
+
+ if (convert->fallback_buffer[plane].device_ptr)
+ return TRUE;
+
+ if (GST_VIDEO_INFO_COMP_DEPTH (info, 0) > 8)
+ element_size = 16;
+
+ ret = CuMemAllocPitch (&convert->fallback_buffer[plane].device_ptr,
+ &convert->fallback_buffer[plane].cuda_stride,
+ GST_VIDEO_INFO_COMP_WIDTH (info, plane) *
+ GST_VIDEO_INFO_COMP_PSTRIDE (info, plane),
+ GST_VIDEO_INFO_COMP_HEIGHT (info, plane), element_size);
+
+ if (!gst_cuda_result (ret)) {
+ GST_ERROR ("failed to allocated fallback memory");
+ return FALSE;
+ }
+
+ return TRUE;
+}
+
+/* create a 2D CUDA texture without alignment check */
+static CUtexObject
+convert_create_texture_unchecked (const CUdeviceptr src, gint width,
+ gint height, gint channels, gint stride, CUarray_format format,
+ CUfilter_mode mode, CUstream cuda_stream)
+{
+ CUDA_TEXTURE_DESC texture_desc;
+ CUDA_RESOURCE_DESC resource_desc;
+ CUtexObject texture = 0;
+ CUresult cuda_ret;
+
+ memset (&texture_desc, 0, sizeof (CUDA_TEXTURE_DESC));
+ memset (&resource_desc, 0, sizeof (CUDA_RESOURCE_DESC));
+
+ resource_desc.resType = CU_RESOURCE_TYPE_PITCH2D;
+ resource_desc.res.pitch2D.format = format;
+ resource_desc.res.pitch2D.numChannels = channels;
+ resource_desc.res.pitch2D.width = width;
+ resource_desc.res.pitch2D.height = height;
+ resource_desc.res.pitch2D.pitchInBytes = stride;
+ resource_desc.res.pitch2D.devPtr = src;
+
+ texture_desc.filterMode = mode;
+ texture_desc.flags = CU_TRSF_READ_AS_INTEGER;
+
+ gst_cuda_result (CuStreamSynchronize (cuda_stream));
+ cuda_ret = CuTexObjectCreate (&texture, &resource_desc, &texture_desc, NULL);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("couldn't create texture");
+
+ return 0;
+ }
+
+ return texture;
+}
+
+static CUtexObject
+convert_create_texture (GstCudaConverter * convert, const GstCudaMemory * src,
+ GstVideoInfo * info, guint plane, CUstream cuda_stream)
+{
+ CUarray_format format = CU_AD_FORMAT_UNSIGNED_INT8;
+ guint channels = 1;
+ CUdeviceptr src_ptr;
+ gsize stride;
+ CUresult cuda_ret;
+ CUfilter_mode mode;
+
+ if (GST_VIDEO_INFO_COMP_DEPTH (info, plane) > 8)
+ format = CU_AD_FORMAT_UNSIGNED_INT16;
+
+ /* FIXME: more graceful method ? */
+ if (plane != 0 &&
+ GST_VIDEO_INFO_N_PLANES (info) != GST_VIDEO_INFO_N_COMPONENTS (info)) {
+ channels = 2;
+ }
+
+ src_ptr = src->data + src->offset[plane];
+ stride = src->stride;
+
+ if (convert->texture_alignment && (src_ptr % convert->texture_alignment)) {
+ CUDA_MEMCPY2D copy_params = { 0, };
+
+ if (!convert_ensure_fallback_memory (convert, info, plane))
+ return 0;
+
+ GST_LOG ("device memory was not aligned, copy to fallback memory");
+
+ copy_params.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+ copy_params.srcPitch = stride;
+ copy_params.srcDevice = (CUdeviceptr) src_ptr;
+
+ copy_params.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+ copy_params.dstPitch = convert->fallback_buffer[plane].cuda_stride;
+ copy_params.dstDevice = convert->fallback_buffer[plane].device_ptr;
+ copy_params.WidthInBytes = GST_VIDEO_INFO_COMP_WIDTH (info, plane)
+ * GST_VIDEO_INFO_COMP_PSTRIDE (info, plane);
+ copy_params.Height = GST_VIDEO_INFO_COMP_HEIGHT (info, plane);
+
+ cuda_ret = CuMemcpy2DAsync (©_params, cuda_stream);
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("failed to copy to fallback buffer");
+ return 0;
+ }
+
+ src_ptr = convert->fallback_buffer[plane].device_ptr;
+ stride = convert->fallback_buffer[plane].cuda_stride;
+ }
+
+ /* Use h/w linear interpolation only when resize is required.
+ * Otherwise the image might be blurred */
+ if (convert->keep_size)
+ mode = CU_TR_FILTER_MODE_POINT;
+ else
+ mode = CU_TR_FILTER_MODE_LINEAR;
+
+ return convert_create_texture_unchecked (src_ptr,
+ GST_VIDEO_INFO_COMP_WIDTH (info, plane),
+ GST_VIDEO_INFO_COMP_HEIGHT (info, plane), channels, stride, format, mode,
+ cuda_stream);
+}
+
+/* main conversion function for YUV to YUV conversion */
+static gboolean
+convert_YUV_TO_YUV (GstCudaConverter * convert,
+ const GstCudaMemory * src, GstVideoInfo * in_info, GstCudaMemory * dst,
+ GstVideoInfo * out_info, CUstream cuda_stream)
+{
+ CUtexObject texture[GST_VIDEO_MAX_PLANES] = { 0, };
+ CUresult cuda_ret;
+ gboolean ret = FALSE;
+ CUdeviceptr dst_ptr[GST_VIDEO_MAX_PLANES] = { 0, };
+ gint dst_stride;
+ gint width, height;
+ gint i;
+
+ gpointer kernel_args[] = { &texture[0], &texture[1], &texture[2],
+ &dst_ptr[0], &dst_ptr[1], &dst_ptr[2], &dst_stride
+ };
+
+ /* conversion step
+ * STEP 1: create CUtexObject per plane
+ * STEP 2: call YUV to YUV conversion kernel function.
+ * resize, uv reordering and bitdepth conversion will be performed in
+ * the CUDA kernel function
+ */
+
+ /* map CUDA device memory to CUDA texture object */
+ for (i = 0; i < GST_VIDEO_INFO_N_PLANES (in_info); i++) {
+ texture[i] = convert_create_texture (convert, src, in_info, i, cuda_stream);
+ if (!texture[i]) {
+ GST_ERROR ("couldn't create texture for %d th plane", i);
+ goto done;
+ }
+ }
+
+ for (i = 0; i < GST_VIDEO_INFO_N_PLANES (out_info); i++)
+ dst_ptr[i] = dst->data + dst->offset[i];
+
+ dst_stride = dst->stride;
+
+ width = GST_VIDEO_INFO_WIDTH (out_info);
+ height = GST_VIDEO_INFO_HEIGHT (out_info);
+
+ cuda_ret =
+ CuLaunchKernel (convert->kernel_func[0], DIV_UP (width, CUDA_BLOCK_X),
+ DIV_UP (height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0,
+ cuda_stream, kernel_args, NULL);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("could not rescale plane");
+ goto done;
+ }
+
+ ret = TRUE;
+ gst_cuda_result (CuStreamSynchronize (cuda_stream));
+
+done:
+ for (i = 0; i < GST_VIDEO_INFO_N_PLANES (in_info); i++) {
+ if (texture[i])
+ gst_cuda_result (CuTexObjectDestroy (texture[i]));
+ }
+
+ return ret;
+}
+
+/* main conversion function for YUV to RGB conversion */
+static gboolean
+convert_YUV_TO_RGB (GstCudaConverter * convert,
+ const GstCudaMemory * src, GstVideoInfo * in_info, GstCudaMemory * dst,
+ GstVideoInfo * out_info, CUstream cuda_stream)
+{
+ CUtexObject texture[GST_VIDEO_MAX_PLANES] = { 0, };
+ CUresult cuda_ret;
+ gboolean ret = FALSE;
+ CUdeviceptr dstRGB = 0;
+ gint dst_stride;
+ gint width, height;
+ gint i;
+
+ gpointer kernel_args[] = { &texture[0], &texture[1], &texture[2],
+ &dstRGB, &dst_stride
+ };
+
+ /* conversion step
+ * STEP 1: create CUtexObject per plane
+ * STEP 2: call YUV to RGB conversion kernel function.
+ * resizing, argb ordering and bitdepth conversion will be performed in
+ * the CUDA kernel function
+ */
+
+ /* map CUDA device memory to CUDA texture object */
+ for (i = 0; i < GST_VIDEO_INFO_N_PLANES (in_info); i++) {
+ texture[i] = convert_create_texture (convert, src, in_info, i, cuda_stream);
+ if (!texture[i]) {
+ GST_ERROR ("couldn't create texture for %d th plane", i);
+ goto done;
+ }
+ }
+
+ dstRGB = dst->data;
+ dst_stride = dst->stride;
+
+ width = GST_VIDEO_INFO_WIDTH (out_info);
+ height = GST_VIDEO_INFO_HEIGHT (out_info);
+
+ cuda_ret =
+ CuLaunchKernel (convert->kernel_func[0], DIV_UP (width, CUDA_BLOCK_X),
+ DIV_UP (height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0,
+ cuda_stream, kernel_args, NULL);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("could not rescale plane");
+ goto done;
+ }
+
+ ret = TRUE;
+ gst_cuda_result (CuStreamSynchronize (cuda_stream));
+
+done:
+ for (i = 0; i < GST_VIDEO_INFO_N_PLANES (in_info); i++) {
+ if (texture[i])
+ gst_cuda_result (CuTexObjectDestroy (texture[i]));
+ }
+
+ return ret;
+}
+
+static gboolean
+convert_UNPACK_RGB (GstCudaConverter * convert, CUfunction kernel_func,
+ CUstream cuda_stream, const GstCudaMemory * src, GstVideoInfo * in_info,
+ CUdeviceptr dst, gint dst_stride, GstCudaRGBOrder * rgb_order)
+{
+ CUdeviceptr srcRGB = 0;
+ gint width, height;
+ gint src_stride, src_pstride;
+ CUresult cuda_ret;
+
+ gpointer unpack_kernel_args[] = { &srcRGB, &dst,
+ &width, &height,
+ &src_stride, &src_pstride, &dst_stride,
+ &convert->in_rgb_order.R, &convert->in_rgb_order.G,
+ &convert->in_rgb_order.B, &convert->in_rgb_order.A,
+ };
+
+ srcRGB = src->data;
+ src_stride = src->stride;
+
+ width = GST_VIDEO_INFO_WIDTH (in_info);
+ height = GST_VIDEO_INFO_HEIGHT (in_info);
+ src_pstride = GST_VIDEO_INFO_COMP_PSTRIDE (in_info, 0);
+
+ cuda_ret =
+ CuLaunchKernel (kernel_func, DIV_UP (width, CUDA_BLOCK_X),
+ DIV_UP (height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0,
+ cuda_stream, unpack_kernel_args, NULL);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("could not unpack rgb");
+ return FALSE;
+ }
+
+ return TRUE;
+}
+
+static gboolean
+convert_TO_Y444 (GstCudaConverter * convert, CUfunction kernel_func,
+ CUstream cuda_stream, CUtexObject srcRGB, CUdeviceptr dstY, gint y_stride,
+ CUdeviceptr dstU, gint u_stride, CUdeviceptr dstV, gint v_stride,
+ gint width, gint height, gint pstride, gint bitdepth)
+{
+ CUresult cuda_ret;
+
+ gpointer kernel_args[] = { &srcRGB, &dstY, &y_stride, &dstU, &u_stride, &dstV,
+ &v_stride, &width, &height, &pstride, &bitdepth,
+ };
+
+ cuda_ret =
+ CuLaunchKernel (kernel_func, DIV_UP (width, CUDA_BLOCK_X),
+ DIV_UP (height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0,
+ cuda_stream, kernel_args, NULL);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("could not unpack rgb");
+ return FALSE;
+ }
+
+ return TRUE;
+}
+
+/* main conversion function for RGB to YUV conversion */
+static gboolean
+convert_RGB_TO_YUV (GstCudaConverter * convert,
+ const GstCudaMemory * src, GstVideoInfo * in_info, GstCudaMemory * dst,
+ GstVideoInfo * out_info, CUstream cuda_stream)
+{
+ CUtexObject texture = 0;
+ CUtexObject yuv_texture[3] = { 0, };
+ CUdeviceptr dst_ptr[GST_VIDEO_MAX_PLANES] = { 0, };
+ CUresult cuda_ret;
+ gboolean ret = FALSE;
+ gint in_width, in_height;
+ gint out_width, out_height;
+ gint dst_stride;
+ CUarray_format format = CU_AD_FORMAT_UNSIGNED_INT8;
+ CUfilter_mode mode = CU_TR_FILTER_MODE_POINT;
+ gint pstride = 1;
+ gint bitdepth = 8;
+ gint i;
+
+ gpointer kernel_args[] = { &yuv_texture[0], &yuv_texture[1], &yuv_texture[2],
+ &dst_ptr[0], &dst_ptr[1], &dst_ptr[2], &dst_stride
+ };
+
+ /* conversion step
+ * STEP 1: unpack src RGB into ARGB or ARGB64 format
+ * STEP 2: convert unpacked ARGB (or ARGB64) to Y444 (or Y444_16LE)
+ * STEP 3: convert Y444 (or Y444_16LE) to final YUV format.
+ * resizing, bitdepth conversion, uv reordering will be performed in
+ * the CUDA kernel function
+ */
+ if (!convert_UNPACK_RGB (convert, convert->kernel_func[0], cuda_stream,
+ src, in_info, convert->unpack_surface.device_ptr,
+ convert->unpack_surface.cuda_stride, &convert->in_rgb_order)) {
+ GST_ERROR ("could not unpack input rgb");
+
+ goto done;
+ }
+
+ in_width = GST_VIDEO_INFO_WIDTH (in_info);
+ in_height = GST_VIDEO_INFO_HEIGHT (in_info);
+
+ out_width = GST_VIDEO_INFO_WIDTH (out_info);
+ out_height = GST_VIDEO_INFO_HEIGHT (out_info);
+ dst_stride = dst->stride;
+
+ if (GST_VIDEO_INFO_COMP_DEPTH (in_info, 0) > 8) {
+ pstride = 2;
+ bitdepth = 16;
+ format = CU_AD_FORMAT_UNSIGNED_INT16;
+ }
+
+ texture =
+ convert_create_texture_unchecked (convert->unpack_surface.device_ptr,
+ in_width, in_height, 4, convert->unpack_surface.cuda_stride, format,
+ mode, cuda_stream);
+
+ if (!texture) {
+ GST_ERROR ("could not create texture");
+ goto done;
+ }
+
+ if (!convert_TO_Y444 (convert, convert->kernel_func[1], cuda_stream, texture,
+ convert->y444_surface[0].device_ptr,
+ convert->y444_surface[0].cuda_stride,
+ convert->y444_surface[1].device_ptr,
+ convert->y444_surface[1].cuda_stride,
+ convert->y444_surface[2].device_ptr,
+ convert->y444_surface[2].cuda_stride, in_width, in_height, pstride,
+ bitdepth)) {
+ GST_ERROR ("could not convert to Y444 or Y444_16LE");
+ goto done;
+ }
+
+ /* Use h/w linear interpolation only when resize is required.
+ * Otherwise the image might be blurred */
+ if (convert->keep_size)
+ mode = CU_TR_FILTER_MODE_POINT;
+ else
+ mode = CU_TR_FILTER_MODE_LINEAR;
+
+ for (i = 0; i < 3; i++) {
+ yuv_texture[i] =
+ convert_create_texture_unchecked (convert->y444_surface[i].device_ptr,
+ in_width, in_height, 1, convert->y444_surface[i].cuda_stride, format,
+ mode, cuda_stream);
+
+ if (!yuv_texture[i]) {
+ GST_ERROR ("could not create %dth yuv texture", i);
+ goto done;
+ }
+ }
+
+ for (i = 0; i < GST_VIDEO_INFO_N_PLANES (out_info); i++)
+ dst_ptr[i] = dst->data + dst->offset[i];
+
+ cuda_ret =
+ CuLaunchKernel (convert->kernel_func[2], DIV_UP (out_width, CUDA_BLOCK_X),
+ DIV_UP (out_height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0,
+ cuda_stream, kernel_args, NULL);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("could not rescale plane");
+ goto done;
+ }
+
+ ret = TRUE;
+ gst_cuda_result (CuStreamSynchronize (cuda_stream));
+
+done:
+ if (texture)
+ gst_cuda_result (CuTexObjectDestroy (texture));
+ for (i = 0; i < 3; i++) {
+ if (yuv_texture[i])
+ gst_cuda_result (CuTexObjectDestroy (yuv_texture[i]));
+ }
+
+ return ret;
+}
+
+/* main conversion function for RGB to RGB conversion */
+static gboolean
+convert_RGB_TO_RGB (GstCudaConverter * convert,
+ const GstCudaMemory * src, GstVideoInfo * in_info, GstCudaMemory * dst,
+ GstVideoInfo * out_info, CUstream cuda_stream)
+{
+ CUtexObject texture = 0;
+ CUresult cuda_ret;
+ gboolean ret = FALSE;
+ CUdeviceptr dstRGB = 0;
+ gint in_width, in_height;
+ gint out_width, out_height;
+ gint dst_stride;
+ CUfilter_mode mode;
+ CUarray_format format = CU_AD_FORMAT_UNSIGNED_INT8;
+
+ gpointer rescale_kernel_args[] = { &texture, &dstRGB, &dst_stride };
+
+ /* conversion step
+ * STEP 1: unpack src RGB into ARGB or ARGB64 format
+ * STEP 2: convert ARGB (or ARGB64) to final RGB format.
+ * resizing, bitdepth conversion, argb reordering will be performed in
+ * the CUDA kernel function
+ */
+
+ if (!convert_UNPACK_RGB (convert, convert->kernel_func[0], cuda_stream,
+ src, in_info, convert->unpack_surface.device_ptr,
+ convert->unpack_surface.cuda_stride, &convert->in_rgb_order)) {
+ GST_ERROR ("could not unpack input rgb");
+
+ goto done;
+ }
+
+ in_width = GST_VIDEO_INFO_WIDTH (in_info);
+ in_height = GST_VIDEO_INFO_HEIGHT (in_info);
+
+ out_width = GST_VIDEO_INFO_WIDTH (out_info);
+ out_height = GST_VIDEO_INFO_HEIGHT (out_info);
+
+ dstRGB = dst->data;
+ dst_stride = dst->stride;
+
+ if (GST_VIDEO_INFO_COMP_DEPTH (in_info, 0) > 8)
+ format = CU_AD_FORMAT_UNSIGNED_INT16;
+
+ /* Use h/w linear interpolation only when resize is required.
+ * Otherwise the image might be blurred */
+ if (convert->keep_size)
+ mode = CU_TR_FILTER_MODE_POINT;
+ else
+ mode = CU_TR_FILTER_MODE_LINEAR;
+
+ texture =
+ convert_create_texture_unchecked (convert->unpack_surface.device_ptr,
+ in_width, in_height, 4, convert->unpack_surface.cuda_stride, format,
+ mode, cuda_stream);
+
+ if (!texture) {
+ GST_ERROR ("could not create texture");
+ goto done;
+ }
+
+ cuda_ret =
+ CuLaunchKernel (convert->kernel_func[1], DIV_UP (out_width, CUDA_BLOCK_X),
+ DIV_UP (out_height, CUDA_BLOCK_Y), 1, CUDA_BLOCK_X, CUDA_BLOCK_Y, 1, 0,
+ cuda_stream, rescale_kernel_args, NULL);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("could not rescale plane");
+ goto done;
+ }
+
+ ret = TRUE;
+ gst_cuda_result (CuStreamSynchronize (cuda_stream));
+
+done:
+ if (texture)
+ gst_cuda_result (CuTexObjectDestroy (texture));
+
+ return ret;
+}
+
+/* from video-converter.c */
+typedef struct
+{
+ gdouble dm[4][4];
+} MatrixData;
+
+static void
+color_matrix_set_identity (MatrixData * m)
+{
+ gint i, j;
+
+ for (i = 0; i < 4; i++) {
+ for (j = 0; j < 4; j++) {
+ m->dm[i][j] = (i == j);
+ }
+ }
+}
+
+static void
+color_matrix_copy (MatrixData * d, const MatrixData * s)
+{
+ gint i, j;
+
+ for (i = 0; i < 4; i++)
+ for (j = 0; j < 4; j++)
+ d->dm[i][j] = s->dm[i][j];
+}
+
+/* Perform 4x4 matrix multiplication:
+ * - @dst@ = @a@ * @b@
+ * - @dst@ may be a pointer to @a@ andor @b@
+ */
+static void
+color_matrix_multiply (MatrixData * dst, MatrixData * a, MatrixData * b)
+{
+ MatrixData tmp;
+ gint i, j, k;
+
+ for (i = 0; i < 4; i++) {
+ for (j = 0; j < 4; j++) {
+ gdouble x = 0;
+ for (k = 0; k < 4; k++) {
+ x += a->dm[i][k] * b->dm[k][j];
+ }
+ tmp.dm[i][j] = x;
+ }
+ }
+ color_matrix_copy (dst, &tmp);
+}
+
+static void
+color_matrix_offset_components (MatrixData * m, gdouble a1, gdouble a2,
+ gdouble a3)
+{
+ MatrixData a;
+
+ color_matrix_set_identity (&a);
+ a.dm[0][3] = a1;
+ a.dm[1][3] = a2;
+ a.dm[2][3] = a3;
+ color_matrix_multiply (m, &a, m);
+}
+
+static void
+color_matrix_scale_components (MatrixData * m, gdouble a1, gdouble a2,
+ gdouble a3)
+{
+ MatrixData a;
+
+ color_matrix_set_identity (&a);
+ a.dm[0][0] = a1;
+ a.dm[1][1] = a2;
+ a.dm[2][2] = a3;
+ color_matrix_multiply (m, &a, m);
+}
+
+static void
+color_matrix_debug (const MatrixData * s)
+{
+ GST_DEBUG ("[%f %f %f %f]", s->dm[0][0], s->dm[0][1], s->dm[0][2],
+ s->dm[0][3]);
+ GST_DEBUG ("[%f %f %f %f]", s->dm[1][0], s->dm[1][1], s->dm[1][2],
+ s->dm[1][3]);
+ GST_DEBUG ("[%f %f %f %f]", s->dm[2][0], s->dm[2][1], s->dm[2][2],
+ s->dm[2][3]);
+ GST_DEBUG ("[%f %f %f %f]", s->dm[3][0], s->dm[3][1], s->dm[3][2],
+ s->dm[3][3]);
+}
+
+static void
+color_matrix_YCbCr_to_RGB (MatrixData * m, gdouble Kr, gdouble Kb)
+{
+ gdouble Kg = 1.0 - Kr - Kb;
+ MatrixData k = {
+ {
+ {1., 0., 2 * (1 - Kr), 0.},
+ {1., -2 * Kb * (1 - Kb) / Kg, -2 * Kr * (1 - Kr) / Kg, 0.},
+ {1., 2 * (1 - Kb), 0., 0.},
+ {0., 0., 0., 1.},
+ }
+ };
+
+ color_matrix_multiply (m, &k, m);
+}
+
+static void
+color_matrix_RGB_to_YCbCr (MatrixData * m, gdouble Kr, gdouble Kb)
+{
+ gdouble Kg = 1.0 - Kr - Kb;
+ MatrixData k;
+ gdouble x;
+
+ k.dm[0][0] = Kr;
+ k.dm[0][1] = Kg;
+ k.dm[0][2] = Kb;
+ k.dm[0][3] = 0;
+
+ x = 1 / (2 * (1 - Kb));
+ k.dm[1][0] = -x * Kr;
+ k.dm[1][1] = -x * Kg;
+ k.dm[1][2] = x * (1 - Kb);
+ k.dm[1][3] = 0;
+
+ x = 1 / (2 * (1 - Kr));
+ k.dm[2][0] = x * (1 - Kr);
+ k.dm[2][1] = -x * Kg;
+ k.dm[2][2] = -x * Kb;
+ k.dm[2][3] = 0;
+
+ k.dm[3][0] = 0;
+ k.dm[3][1] = 0;
+ k.dm[3][2] = 0;
+ k.dm[3][3] = 1;
+
+ color_matrix_multiply (m, &k, m);
+}
+
+static void
+compute_matrix_to_RGB (GstCudaConverter * convert, MatrixData * data,
+ GstVideoInfo * info)
+{
+ gdouble Kr = 0, Kb = 0;
+ gint offset[4], scale[4];
+
+ /* bring color components to [0..1.0] range */
+ gst_video_color_range_offsets (info->colorimetry.range, info->finfo, offset,
+ scale);
+
+ color_matrix_offset_components (data, -offset[0], -offset[1], -offset[2]);
+ color_matrix_scale_components (data, 1 / ((float) scale[0]),
+ 1 / ((float) scale[1]), 1 / ((float) scale[2]));
+
+ if (!GST_VIDEO_INFO_IS_RGB (info)) {
+ /* bring components to R'G'B' space */
+ if (gst_video_color_matrix_get_Kr_Kb (info->colorimetry.matrix, &Kr, &Kb))
+ color_matrix_YCbCr_to_RGB (data, Kr, Kb);
+ }
+ color_matrix_debug (data);
+}
+
+static void
+compute_matrix_to_YUV (GstCudaConverter * convert, MatrixData * data,
+ GstVideoInfo * info)
+{
+ gdouble Kr = 0, Kb = 0;
+ gint offset[4], scale[4];
+
+ if (!GST_VIDEO_INFO_IS_RGB (info)) {
+ /* bring components to YCbCr space */
+ if (gst_video_color_matrix_get_Kr_Kb (info->colorimetry.matrix, &Kr, &Kb))
+ color_matrix_RGB_to_YCbCr (data, Kr, Kb);
+ }
+
+ /* bring color components to nominal range */
+ gst_video_color_range_offsets (info->colorimetry.range, info->finfo, offset,
+ scale);
+
+ color_matrix_scale_components (data, (float) scale[0], (float) scale[1],
+ (float) scale[2]);
+ color_matrix_offset_components (data, offset[0], offset[1], offset[2]);
+
+ color_matrix_debug (data);
+}
+
+static gboolean
+cuda_converter_get_matrix (GstCudaConverter * convert, MatrixData * matrix,
+ GstVideoInfo * in_info, GstVideoInfo * out_info)
+{
+ gboolean same_matrix, same_bits;
+ guint in_bits, out_bits;
+
+ in_bits = GST_VIDEO_INFO_COMP_DEPTH (in_info, 0);
+ out_bits = GST_VIDEO_INFO_COMP_DEPTH (out_info, 0);
+
+ same_bits = in_bits == out_bits;
+ same_matrix = in_info->colorimetry.matrix == out_info->colorimetry.matrix;
+
+ GST_DEBUG ("matrix %d -> %d (%d)", in_info->colorimetry.matrix,
+ out_info->colorimetry.matrix, same_matrix);
+ GST_DEBUG ("bits %d -> %d (%d)", in_bits, out_bits, same_bits);
+
+ color_matrix_set_identity (matrix);
+
+ if (same_bits && same_matrix) {
+ GST_DEBUG ("conversion matrix is not required");
+
+ return FALSE;
+ }
+
+ if (in_bits < out_bits) {
+ gint scale = 1 << (out_bits - in_bits);
+ color_matrix_scale_components (matrix,
+ 1 / (float) scale, 1 / (float) scale, 1 / (float) scale);
+ }
+
+ GST_DEBUG ("to RGB matrix");
+ compute_matrix_to_RGB (convert, matrix, in_info);
+ GST_DEBUG ("current matrix");
+ color_matrix_debug (matrix);
+
+ GST_DEBUG ("to YUV matrix");
+ compute_matrix_to_YUV (convert, matrix, out_info);
+ GST_DEBUG ("current matrix");
+ color_matrix_debug (matrix);
+
+ if (in_bits > out_bits) {
+ gint scale = 1 << (in_bits - out_bits);
+ color_matrix_scale_components (matrix,
+ (float) scale, (float) scale, (float) scale);
+ }
+
+ GST_DEBUG ("final matrix");
+ color_matrix_debug (matrix);
+
+ return TRUE;
+}
+
+static gboolean
+is_uv_swapped (GstVideoFormat format)
+{
+ static GstVideoFormat swapped_formats[] = {
+ GST_VIDEO_FORMAT_YV12,
+ GST_VIDEO_FORMAT_NV21,
+ };
+ gint i;
+
+ for (i = 0; i < G_N_ELEMENTS (swapped_formats); i++) {
+ if (format == swapped_formats[i])
+ return TRUE;
+ }
+
+ return FALSE;
+}
+
+typedef struct
+{
+ const gchar *read_chroma;
+ const gchar *write_chroma;
+ const gchar *unpack_function;
+ gfloat scale_h, scale_v;
+ gfloat chroma_scale_h, chroma_scale_v;
+ gint width, height;
+ gint chroma_width, chroma_height;
+ gint in_depth;
+ gint out_depth;
+ gint pstride, chroma_pstride;
+ gint in_shift, out_shift;
+ gint mask;
+ gint swap_uv;
+ /* RGBA specific variables */
+ gint max_in_val;
+ GstCudaRGBOrder rgb_order;
+} GstCudaKernelTempl;
+
+static gchar *
+cuda_converter_generate_yuv_to_yuv_kernel_code (GstCudaConverter * convert,
+ GstCudaKernelTempl * templ)
+{
+ return g_strdup_printf (templ_YUV_TO_YUV,
+ templ->scale_h, templ->scale_v, templ->chroma_scale_h,
+ templ->chroma_scale_v, templ->width, templ->height, templ->chroma_width,
+ templ->chroma_height, templ->in_depth, templ->out_depth, templ->pstride,
+ templ->chroma_pstride, templ->in_shift, templ->out_shift, templ->mask,
+ templ->swap_uv, templ->read_chroma, templ->write_chroma);
+}
+
+static gchar *
+cuda_converter_generate_yuv_to_rgb_kernel_code (GstCudaConverter * convert,
+ GstCudaKernelTempl * templ, MatrixData * matrix)
+{
+ return g_strdup_printf (templ_YUV_TO_RGB,
+ matrix->dm[0][3], matrix->dm[1][3], matrix->dm[2][3],
+ matrix->dm[0][0], matrix->dm[0][1], matrix->dm[0][2],
+ matrix->dm[1][0], matrix->dm[1][1], matrix->dm[1][2],
+ matrix->dm[2][0], matrix->dm[2][1], matrix->dm[2][2],
+ templ->scale_h, templ->scale_v, templ->chroma_scale_h,
+ templ->chroma_scale_v, templ->width, templ->height, templ->chroma_width,
+ templ->chroma_height, templ->in_depth, templ->out_depth, templ->pstride,
+ templ->chroma_pstride, templ->in_shift, templ->out_shift, templ->mask,
+ templ->swap_uv, templ->max_in_val, templ->rgb_order.R,
+ templ->rgb_order.G, templ->rgb_order.B, templ->rgb_order.A,
+ templ->rgb_order.X, templ->read_chroma);
+}
+
+static gchar *
+cuda_converter_generate_rgb_to_yuv_kernel_code (GstCudaConverter * convert,
+ GstCudaKernelTempl * templ, MatrixData * matrix)
+{
+ return g_strdup_printf (templ_RGB_TO_YUV,
+ matrix->dm[0][3], matrix->dm[1][3], matrix->dm[2][3],
+ matrix->dm[0][0], matrix->dm[0][1], matrix->dm[0][2],
+ matrix->dm[1][0], matrix->dm[1][1], matrix->dm[1][2],
+ matrix->dm[2][0], matrix->dm[2][1], matrix->dm[2][2],
+ templ->scale_h, templ->scale_v, templ->chroma_scale_h,
+ templ->chroma_scale_v, templ->width, templ->height, templ->chroma_width,
+ templ->chroma_height, templ->in_depth, templ->out_depth, templ->pstride,
+ templ->chroma_pstride, templ->in_shift, templ->out_shift, templ->mask,
+ templ->swap_uv, templ->unpack_function, templ->read_chroma,
+ templ->write_chroma);
+}
+
+static gchar *
+cuda_converter_generate_rgb_to_rgb_kernel_code (GstCudaConverter * convert,
+ GstCudaKernelTempl * templ)
+{
+ return g_strdup_printf (templ_RGB_to_RGB,
+ templ->scale_h, templ->scale_v,
+ templ->width, templ->height,
+ templ->in_depth, templ->out_depth, templ->pstride,
+ templ->rgb_order.R, templ->rgb_order.G,
+ templ->rgb_order.B, templ->rgb_order.A, templ->rgb_order.X,
+ templ->unpack_function);
+}
+
+#define SET_ORDER(o,r,g,b,a,x) G_STMT_START { \
+ (o)->R = (r); \
+ (o)->G = (g); \
+ (o)->B = (b); \
+ (o)->A = (a); \
+ (o)->X = (x); \
+} G_STMT_END
+
+static void
+cuda_converter_get_rgb_order (GstVideoFormat format, GstCudaRGBOrder * order)
+{
+ switch (format) {
+ case GST_VIDEO_FORMAT_RGBA:
+ SET_ORDER (order, 0, 1, 2, 3, -1);
+ break;
+ case GST_VIDEO_FORMAT_RGBx:
+ SET_ORDER (order, 0, 1, 2, -1, 3);
+ break;
+ case GST_VIDEO_FORMAT_BGRA:
+ SET_ORDER (order, 2, 1, 0, 3, -1);
+ break;
+ case GST_VIDEO_FORMAT_BGRx:
+ SET_ORDER (order, 2, 1, 0, -1, 3);
+ break;
+ case GST_VIDEO_FORMAT_ARGB:
+ SET_ORDER (order, 1, 2, 3, 0, -1);
+ break;
+ case GST_VIDEO_FORMAT_ABGR:
+ SET_ORDER (order, 3, 2, 1, 0, -1);
+ break;
+ case GST_VIDEO_FORMAT_RGB:
+ SET_ORDER (order, 0, 1, 2, -1, -1);
+ break;
+ case GST_VIDEO_FORMAT_BGR:
+ SET_ORDER (order, 2, 1, 0, -1, -1);
+ break;
+ case GST_VIDEO_FORMAT_BGR10A2_LE:
+ SET_ORDER (order, 1, 2, 3, 0, -1);
+ break;
+ case GST_VIDEO_FORMAT_RGB10A2_LE:
+ SET_ORDER (order, 3, 2, 1, 0, -1);
+ break;
+ default:
+ g_assert_not_reached ();
+ break;
+ }
+}
+
+static gboolean
+cuda_converter_lookup_path (GstCudaConverter * convert)
+{
+ GstVideoFormat in_format, out_format;
+ gboolean src_yuv, dst_yuv;
+ gboolean src_planar, dst_planar;
+ GstCudaKernelTempl templ = { 0, };
+ GstVideoInfo *in_info, *out_info;
+ gboolean ret = FALSE;
+ CUresult cuda_ret;
+
+ in_info = &convert->in_info;
+ out_info = &convert->out_info;
+
+ in_format = GST_VIDEO_INFO_FORMAT (in_info);
+ out_format = GST_VIDEO_INFO_FORMAT (out_info);
+
+ src_yuv = GST_VIDEO_INFO_IS_YUV (in_info);
+ dst_yuv = GST_VIDEO_INFO_IS_YUV (out_info);
+
+ src_planar = GST_VIDEO_INFO_N_PLANES (in_info) ==
+ GST_VIDEO_INFO_N_COMPONENTS (in_info);
+ dst_planar = GST_VIDEO_INFO_N_PLANES (out_info) ==
+ GST_VIDEO_INFO_N_COMPONENTS (out_info);
+
+ convert->keep_size = (GST_VIDEO_INFO_WIDTH (&convert->in_info) ==
+ GST_VIDEO_INFO_WIDTH (&convert->out_info) &&
+ GST_VIDEO_INFO_HEIGHT (&convert->in_info) ==
+ GST_VIDEO_INFO_HEIGHT (&convert->out_info));
+
+ templ.scale_h = (gfloat) GST_VIDEO_INFO_COMP_WIDTH (in_info, 0) /
+ (gfloat) GST_VIDEO_INFO_COMP_WIDTH (out_info, 0);
+ templ.scale_v = (gfloat) GST_VIDEO_INFO_COMP_HEIGHT (in_info, 0) /
+ (gfloat) GST_VIDEO_INFO_COMP_HEIGHT (out_info, 0);
+ templ.chroma_scale_h = (gfloat) GST_VIDEO_INFO_COMP_WIDTH (in_info, 1) /
+ (gfloat) GST_VIDEO_INFO_COMP_WIDTH (out_info, 1);
+ templ.chroma_scale_v = (gfloat) GST_VIDEO_INFO_COMP_HEIGHT (in_info, 1) /
+ (gfloat) GST_VIDEO_INFO_COMP_HEIGHT (out_info, 1);
+ templ.width = GST_VIDEO_INFO_COMP_WIDTH (out_info, 0);
+ templ.height = GST_VIDEO_INFO_COMP_HEIGHT (out_info, 0);
+ templ.chroma_width = GST_VIDEO_INFO_COMP_WIDTH (out_info, 1);
+ templ.chroma_height = GST_VIDEO_INFO_COMP_HEIGHT (out_info, 1);
+
+ templ.in_depth = GST_VIDEO_INFO_COMP_DEPTH (in_info, 0);
+ templ.out_depth = GST_VIDEO_INFO_COMP_DEPTH (out_info, 0);
+ templ.pstride = GST_VIDEO_INFO_COMP_PSTRIDE (out_info, 0);
+ templ.chroma_pstride = GST_VIDEO_INFO_COMP_PSTRIDE (out_info, 1);
+ templ.in_shift = in_info->finfo->shift[0];
+ templ.out_shift = out_info->finfo->shift[0];
+ templ.mask = ((1 << templ.out_depth) - 1) << templ.out_shift;
+ templ.swap_uv = (is_uv_swapped (in_format) != is_uv_swapped (out_format));
+
+ if (src_yuv && dst_yuv) {
+ convert->convert = convert_YUV_TO_YUV;
+
+ if (src_planar && dst_planar) {
+ templ.read_chroma = READ_CHROMA_FROM_PLANAR;
+ templ.write_chroma = WRITE_CHROMA_TO_PLANAR;
+ } else if (!src_planar && dst_planar) {
+ templ.read_chroma = READ_CHROMA_FROM_SEMI_PLANAR;
+ templ.write_chroma = WRITE_CHROMA_TO_PLANAR;
+ } else if (src_planar && !dst_planar) {
+ templ.read_chroma = READ_CHROMA_FROM_PLANAR;
+ templ.write_chroma = WRITE_CHROMA_TO_SEMI_PLANAR;
+ } else {
+ templ.read_chroma = READ_CHROMA_FROM_SEMI_PLANAR;
+ templ.write_chroma = WRITE_CHROMA_TO_SEMI_PLANAR;
+ }
+
+ convert->kernel_source =
+ cuda_converter_generate_yuv_to_yuv_kernel_code (convert, &templ);
+ convert->func_names[0] = GST_CUDA_KERNEL_FUNC;
+
+ ret = TRUE;
+ } else if (src_yuv && !dst_yuv) {
+ MatrixData matrix;
+
+ if (src_planar) {
+ templ.read_chroma = READ_CHROMA_FROM_PLANAR;
+ } else {
+ templ.read_chroma = READ_CHROMA_FROM_SEMI_PLANAR;
+ }
+
+ templ.max_in_val = (1 << templ.in_depth) - 1;
+ cuda_converter_get_rgb_order (out_format, &templ.rgb_order);
+
+ cuda_converter_get_matrix (convert, &matrix, in_info, out_info);
+ convert->kernel_source =
+ cuda_converter_generate_yuv_to_rgb_kernel_code (convert,
+ &templ, &matrix);
+ convert->func_names[0] = GST_CUDA_KERNEL_FUNC;
+
+ convert->convert = convert_YUV_TO_RGB;
+
+ ret = TRUE;
+ } else if (!src_yuv && dst_yuv) {
+ MatrixData matrix;
+ gsize element_size = 8;
+ GstVideoFormat unpack_format;
+ GstVideoFormat y444_format;
+ GstVideoInfo unpack_info;
+ GstVideoInfo y444_info;
+ gint i;
+
+ if (dst_planar) {
+ templ.write_chroma = WRITE_CHROMA_TO_PLANAR;
+ } else {
+ templ.write_chroma = WRITE_CHROMA_TO_SEMI_PLANAR;
+ }
+ templ.read_chroma = READ_CHROMA_FROM_PLANAR;
+
+ cuda_converter_get_rgb_order (in_format, &convert->in_rgb_order);
+
+ if (templ.in_depth > 8) {
+ /* FIXME: RGB10A2_LE and BGR10A2_LE only */
+ element_size = 16;
+ unpack_format = GST_VIDEO_FORMAT_ARGB64;
+ y444_format = GST_VIDEO_FORMAT_Y444_16LE;
+ templ.unpack_function = unpack_to_ARGB64;
+ } else {
+ unpack_format = GST_VIDEO_FORMAT_ARGB;
+ y444_format = GST_VIDEO_FORMAT_Y444;
+ templ.unpack_function = unpack_to_ARGB;
+ }
+
+ gst_video_info_set_format (&unpack_info,
+ unpack_format, GST_VIDEO_INFO_WIDTH (in_info),
+ GST_VIDEO_INFO_HEIGHT (in_info));
+ gst_video_info_set_format (&y444_info,
+ y444_format, GST_VIDEO_INFO_WIDTH (in_info),
+ GST_VIDEO_INFO_HEIGHT (in_info));
+
+ templ.in_depth = GST_VIDEO_INFO_COMP_DEPTH (&unpack_info, 0);
+
+ cuda_ret = CuMemAllocPitch (&convert->unpack_surface.device_ptr,
+ &convert->unpack_surface.cuda_stride,
+ GST_VIDEO_INFO_COMP_WIDTH (&unpack_info, 0) *
+ GST_VIDEO_INFO_COMP_PSTRIDE (&unpack_info, 0),
+ GST_VIDEO_INFO_HEIGHT (&unpack_info), element_size);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("couldn't alloc unpack surface");
+ return FALSE;
+ }
+
+ for (i = 0; i < 3; i++) {
+ cuda_ret = CuMemAllocPitch (&convert->y444_surface[i].device_ptr,
+ &convert->y444_surface[i].cuda_stride,
+ GST_VIDEO_INFO_COMP_WIDTH (&y444_info, i) *
+ GST_VIDEO_INFO_COMP_PSTRIDE (&y444_info, i),
+ GST_VIDEO_INFO_COMP_HEIGHT (&y444_info, i), element_size);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("couldn't alloc %dth y444 surface", i);
+ return FALSE;
+ }
+ }
+
+ cuda_converter_get_matrix (convert, &matrix, &unpack_info, &y444_info);
+
+ convert->kernel_source =
+ cuda_converter_generate_rgb_to_yuv_kernel_code (convert,
+ &templ, &matrix);
+
+ convert->func_names[0] = GST_CUDA_KERNEL_FUNC_TO_ARGB;
+ convert->func_names[1] = GST_CUDA_KERNEL_FUNC_TO_Y444;
+ convert->func_names[2] = GST_CUDA_KERNEL_FUNC_Y444_TO_YUV;
+
+ convert->convert = convert_RGB_TO_YUV;
+
+ ret = TRUE;
+ } else {
+ gsize element_size = 8;
+ GstVideoFormat unpack_format;
+ GstVideoInfo unpack_info;
+
+ cuda_converter_get_rgb_order (in_format, &convert->in_rgb_order);
+ cuda_converter_get_rgb_order (out_format, &templ.rgb_order);
+
+ if (templ.in_depth > 8) {
+ /* FIXME: RGB10A2_LE and BGR10A2_LE only */
+ element_size = 16;
+ unpack_format = GST_VIDEO_FORMAT_ARGB64;
+ templ.unpack_function = unpack_to_ARGB64;
+ } else {
+ unpack_format = GST_VIDEO_FORMAT_ARGB;
+ templ.unpack_function = unpack_to_ARGB;
+ }
+
+ gst_video_info_set_format (&unpack_info,
+ unpack_format, GST_VIDEO_INFO_WIDTH (in_info),
+ GST_VIDEO_INFO_HEIGHT (in_info));
+
+ templ.in_depth = GST_VIDEO_INFO_COMP_DEPTH (&unpack_info, 0);
+
+ cuda_ret = CuMemAllocPitch (&convert->unpack_surface.device_ptr,
+ &convert->unpack_surface.cuda_stride,
+ GST_VIDEO_INFO_COMP_WIDTH (&unpack_info, 0) *
+ GST_VIDEO_INFO_COMP_PSTRIDE (&unpack_info, 0),
+ GST_VIDEO_INFO_HEIGHT (&unpack_info), element_size);
+
+ if (!gst_cuda_result (cuda_ret)) {
+ GST_ERROR ("couldn't alloc unpack surface");
+ return FALSE;
+ }
+
+ convert->kernel_source =
+ cuda_converter_generate_rgb_to_rgb_kernel_code (convert, &templ);
+
+ convert->func_names[0] = GST_CUDA_KERNEL_FUNC_TO_ARGB;
+ convert->func_names[1] = GST_CUDA_KERNEL_FUNC_SCALE_RGB;
+
+ convert->convert = convert_RGB_TO_RGB;
+
+ ret = TRUE;
+ }
+
+ if (!ret) {
+ GST_DEBUG ("no path found");
+
+ return FALSE;
+ }
+
+ GST_TRACE ("configured CUDA kernel source\n%s", convert->kernel_source);
+
+ return TRUE;
+}
projects / platform / upstream / gstreamer.git / commitdiff