Add beginning of edge detection

Clear out the approximation portion of the wavelet, and inverse transform.
This likely needs a threshold -> saturate to emphasize the edges.
Create a buffer of logo ROI data only
Make the edges dilate and accumulate the logo

mythtv/programs/mythgpucommflag/videoConvert.cl

@@ -40,6 +40,26 @@ void videoYUVToSNORM(__read_only image2d_t in, __write_only image2d_t out)
     write_imagef(out, (int2)(x, y), val);
 }
 
+// Convert the YUV from use with SNORM buffers (for dumping)
+__kernel
+void videoYUVFromSNORM(__read_only image2d_t in, __write_only image2d_t out)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP_TO_EDGE |
+                              CLK_FILTER_NEAREST;
+
+    float4 val;
+
+    val = read_imagef(in,  sampler, (int2)(x, y));
+    val.x += 1.0;
+    val.x /= 2.0;
+    val.w  = 1.0;
+
+    write_imagef(out, (int2)(x, y), val);
+}
+
 
 // Convert Packed YUV into RGB
 __kernel
@@ -64,6 +84,147 @@ void videoYUVToRGB(__read_only image2d_t YUV, __write_only image2d_t RGB)
     write_imagef(RGB, (int2)(x, y), rgbval);
 }
 
+// Zero a region of a buffer
+__kernel
+void videoZeroRegion(__read_only image2d_t in, __write_only image2d_t out,
+                     int2 regionStart, int2 regionEnd)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP_TO_EDGE |
+                              CLK_FILTER_NEAREST;
+    float4 output;
+
+    if (x >= regionStart.x && x < regionEnd.x &&
+        y >= regionStart.y && y < regionEnd.y)
+    {
+        output = 0.0;
+        output.x = -1.0;    // since we are in SNORM mode.
+    }
+    else
+    {
+        output = read_imagef(in,  sampler, (int2)(x, y));
+    }
+
+    output.w = 1.0;
+
+    write_imagef(out, (int2)(x, y), output);
+}
+
+// Threshold => saturate
+__kernel
+void videoThreshSaturate(__read_only image2d_t in, __write_only image2d_t out,
+                         int threshold)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP_TO_EDGE |
+                              CLK_FILTER_NEAREST;
+    float4 output;
+    float thresh = ((convert_float(threshold) / 255.0) - 0.5) * 2.0;
+    float thresh2 = convert_float(threshold) / 255.0;
+
+    output = read_imagef(in,  sampler, (int2)(x, y));
+
+    output.x = (ceil(fmax(output.x - thresh, 0.0)) - 0.5) * 2.0;
+    output.y = ceil(fmax(fabs(output.y) - thresh2, 0.0));
+    output.z = ceil(fmax(fabs(output.z) - thresh2, 0.0));
+    output.w = 1.0;
+
+    write_imagef(out, (int2)(x, y), output);
+}
+
+__kernel
+void videoCopyLogoROI(__read_only image2d_t in, __write_only image2d_t out,
+                      int2 regionStart, int2 regionEnd)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP_TO_EDGE |
+                              CLK_FILTER_NEAREST;
+    float4 output;
+
+    int xover = min(x / regionStart.x, 1);
+    int yover = min(y / regionStart.y, 1);
+    int xin = x + (xover * (regionEnd.x - regionStart.x));
+    int yin = y + (yover * (regionEnd.y - regionStart.y));
+
+    output = read_imagef(in,  sampler, (int2)(xin, yin));
+    write_imagef(out, (int2)(x, y), output);
+}
+
+__kernel
+void videoLogoMSE(__read_only image2d_t ref, __read_only image2d_t in,
+                  __write_only image2d_t out)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP_TO_EDGE |
+                              CLK_FILTER_NEAREST;
+    float4 in1;
+    float4 in2;
+    float4 output;
+
+    in1 = read_imagef(ref, sampler, (int2)(x, y));
+    in2 = read_imagef(in, sampler, (int2)(x, y));
+    output = 1.0;
+    output.x = in2.x - in1.x;
+    output.x *= output.x * 2.0;
+    output.x -= 1.0;
+    write_imagef(out, (int2)(x, y), output);
+}
+
+__kernel
+void videoMultiply(__read_only image2d_t ref, __read_only image2d_t in,
+                   __write_only image2d_t out)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP_TO_EDGE |
+                              CLK_FILTER_NEAREST;
+    float4 in1;
+    float4 in2;
+    float4 output;
+
+    in1 = (read_imagef(ref, sampler, (int2)(x, y)) + 1.0) / 2.0;
+    in2 = (read_imagef(in, sampler, (int2)(x, y)) + 1.0) / 2.0;
+    output = 0.0;
+    output.x = ((in2.x * in1.x) * 2.0) - 1.0;
+    write_imagef(out, (int2)(x, y), output);
+}
+
+__kernel
+void videoDilate3x3(__read_only image2d_t in, __write_only image2d_t out)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP |
+                              CLK_FILTER_NEAREST;
+    float4 output = -1.0;
+
+    for (int i = -1; i <= 1; i++)
+    {
+        for (int j = -1; j <= 1; j++)
+        {
+            int locx = x + i;
+            int locy = y + j;
+
+            output = fmax(output, read_imagef(in, sampler, (int2)(locx, locy)));
+        }
+    }
+    output.yz = 0.0;
+    output.w  = 1.0;
+
+    write_imagef(out, (int2)(x, y), output);
+}
+
+
 
 /*
  * vim:ts=4:sw=4:ai:et:si:sts=4:filetype=c

mythtv/programs/mythgpucommflag/videoInverse.cl

@@ -0,0 +1,19 @@
+__kernel
+void videoInvert(__read_only image2d_t frameIn,
+                 __write_only image2d_t frameOut)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                              CLK_ADDRESS_CLAMP_TO_EDGE |
+                              CLK_FILTER_NEAREST;
+
+    float4 val;
+    val = read_imagef(frameIn, sampler, (int2)(x, y));
+    val.x *= -1.0;
+    write_imagef(frameOut, (int2)(x, y), val);
+}
+
+/*
+ * vim:ts=4:sw=4:ai:et:si:sts=4:filetype=c
+ */

mythtv/programs/mythgpucommflag/videoconsumer.cpp

@@ -187,13 +187,13 @@ void VideoConsumer::ProcessPacket(Packet *packet)
     if (m_dev)
     {
         // Push frame to the GPU via OpenGL/OpenCL
-        videoFrame = new VideoPacket(m_decoder, &mpa_pic);
         count++;
+        videoFrame = new VideoPacket(m_decoder, &mpa_pic, count);
         if ((count <= 100) && videoFrame)
         {
-            videoFrame->m_frameRaw->Dump("frame", count);
-            videoFrame->m_frameYUVSNORM->Dump("yuv", count);
-            videoFrame->m_wavelet->Dump("wavelet", count);
+            // videoFrame->m_frameRaw->Dump("frame", count);
+            // videoFrame->m_frameYUVSNORM->Dump("yuv", count);
+            // videoFrame->m_wavelet->Dump("wavelet", count);
             VideoSurface rgb(m_dev, kSurfaceRGB,
                              videoFrame->m_frameRaw->m_width,
                              videoFrame->m_frameRaw->m_height);
@@ -204,7 +204,13 @@ void VideoConsumer::ProcessPacket(Packet *packet)
                              videoFrame->m_frameRaw->m_width,
                              videoFrame->m_frameRaw->m_height);
             OpenCLWaveletInverse(m_dev, videoFrame->m_wavelet, &yuv);
-            yuv.Dump("unwavelet", count);
+            // yuv.Dump("unwaveletYUV", count);
+            VideoSurface yuv2(m_dev, kSurfaceYUV,
+                              videoFrame->m_frameRaw->m_width,
+                              videoFrame->m_frameRaw->m_height);
+            OpenCLYUVFromSNORM(m_dev, &yuv, &yuv2);
+            OpenCLYUVToRGB(m_dev, &yuv2, &rgb);
+            rgb.Dump("unwaveletRGB", count);
         }
     }
 

mythtv/programs/mythgpucommflag/videopacket.cpp

@@ -13,8 +13,8 @@ extern "C" {
 
 VideoPacketMap videoPacketMap;
 
-VideoPacket::VideoPacket(VideoDecoder *decoder, AVFrame *frame) :
-    m_decoder(decoder), m_frameIn(frame)
+VideoPacket::VideoPacket(VideoDecoder *decoder, AVFrame *frame, int num) :
+    m_num(num), m_decoder(decoder), m_frameIn(frame)
 {
     m_frameRaw = m_decoder->DecodeFrame(frame);
 

mythtv/programs/mythgpucommflag/videopacket.h

@@ -14,9 +14,10 @@ extern "C" {
 class VideoPacket
 {
   public:
-    VideoPacket(VideoDecoder *decoder, AVFrame *frame);
+    VideoPacket(VideoDecoder *decoder, AVFrame *frame, int num);
     ~VideoPacket();
 
+    uint64_t m_num;
     VideoDecoder *m_decoder;
     AVFrame *m_frameIn;
     VideoSurface *m_frameRaw;

mythtv/programs/mythgpucommflag/videoprocessor.h

@@ -56,6 +56,20 @@ void OpenCLYUVToRGB(OpenCLDevice *dev, VideoSurface *yuvframe,
                     VideoSurface *rgbframe);
 void OpenCLYUVToSNORM(OpenCLDevice *dev, VideoSurface *inframe,
                       VideoSurface *outframe);
+void OpenCLYUVFromSNORM(OpenCLDevice *dev, VideoSurface *inframe,
+                        VideoSurface *outframe);
+void OpenCLZeroRegion(OpenCLDevice *dev, VideoSurface *inframe,
+                      VideoSurface *outframe, int x1, int y1, int x2, int y2);
+void OpenCLThreshSat(OpenCLDevice *dev, VideoSurface *inframe,
+                     VideoSurface *outframe, int threshold);
+void OpenCLCopyLogoROI(OpenCLDevice *dev, VideoSurface *inframe,
+                       VideoSurface *outframe, int roiWidth, int roiHeight);
+void OpenCLLogoMSE(OpenCLDevice *dev, VideoSurface *ref, VideoSurface *in,
+                   VideoSurface *out);
+void OpenCLInvert(OpenCLDevice *dev, VideoSurface *in, VideoSurface *out);
+void OpenCLMultiply(OpenCLDevice *dev, VideoSurface *ref, VideoSurface *in,
+                    VideoSurface *out);
+void OpenCLDilate3x3(OpenCLDevice *dev, VideoSurface *in, VideoSurface *out);
 
 #endif
 

mythtv/programs/mythgpucommflag/videosurface.cpp

@@ -26,7 +26,9 @@ VideoSurface::VideoSurface(OpenCLDevice *dev, uint32_t width, uint32_t height,
     memset(m_clBuffer, 0, m_bufCount * sizeof(cl_mem));
     m_render.surface = m_vdpSurface;
 
-    char *zeros = new char[m_width*m_height/2];
+    m_realWidth = m_width;
+    m_realHeight = m_height / 2;
+    char *zeros = new char[m_realWidth*m_realHeight];
 
     glGenTextures(4, m_glVDPAUTex);
     glGenTextures(4, m_glOpenCLTex);
@@ -65,8 +67,8 @@ VideoSurface::VideoSurface(OpenCLDevice *dev, uint32_t width, uint32_t height,
 
         glTexImage2D(GL_TEXTURE_RECTANGLE_ARB, 0,
                      (i < 2) ? GL_RED : GL_RG,
-                     (i < 2) ? m_width : m_width / 2,
-                     (i < 2) ? m_height / 2 : m_height / 4,
+                     (i < 2) ? m_realWidth : m_realWidth / 2,
+                     (i < 2) ? m_realHeight : m_realHeight / 2,
                      0, 
                      (i < 2) ? GL_RED : GL_RG,
                      GL_UNSIGNED_BYTE, 
@@ -108,30 +110,44 @@ VideoSurface::VideoSurface(OpenCLDevice *dev, VideoSurfaceType type,
     cl_image_format format;
     uint32_t bufWidth;
     uint32_t bufHeight;
+    uint32_t temp = 0;
 
     m_bufCount = 2;
+    bufWidth = m_width;
+    bufHeight = m_height / 2;
 
     switch (m_type)
     {
         case kSurfaceWavelet:
         case kSurfaceYUV2:
-        case kSurfaceYUV:
             format.image_channel_data_type = CL_SNORM_INT8;
             m_bufCount = 4;
             break;
+        case kSurfaceLogoROI:
+            bufHeight /= 2;
+            temp = (m_height / 3) * 4;
+            bufWidth = m_width - (temp / 2);
+        case kSurfaceYUV:
+            format.image_channel_data_type = CL_SNORM_INT8;
+            break;
+        case kSurfaceLogoRGB:
+            bufHeight /= 2;
+            temp = (m_height / 3) * 4;
+            bufWidth = m_width - (temp / 2);
         case kSurfaceRGB:
             format.image_channel_data_type = CL_UNORM_INT8;
             break;
         default:
             return;
     }
 
+    m_realWidth = bufWidth;
+    m_realHeight = bufHeight;
+
     m_clBuffer = new cl_mem[m_bufCount];
     memset(m_clBuffer, 0, m_bufCount * sizeof(cl_mem));
 
     format.image_channel_order = CL_RGBA;
-    bufWidth = m_width;
-    bufHeight = m_height / 2;
 
     for (int i = 0; i < m_bufCount; i++)
     {
@@ -178,8 +194,8 @@ void VideoSurface::Bind(void)
         for (int j = 0; j < 2; j++)
         {
             GLenum bufnum = GL_COLOR_ATTACHMENT0_EXT + j;
-            uint32_t width  = (i < 1) ? m_width : m_width / 2;
-            uint32_t height = (i < 1) ? m_height / 2 : m_height / 4; 
+            uint32_t width  = (i < 1) ? m_realWidth : m_realWidth / 2;
+            uint32_t height = (i < 1) ? m_realHeight : m_realHeight / 2; 
 
             glReadBuffer(bufnum);
             glDrawBuffer(bufnum);
@@ -208,6 +224,8 @@ void VideoSurface::Dump(QString basename, int framenum)
             break;
         case kSurfaceYUV:
         case kSurfaceRGB:
+        case kSurfaceLogoROI:
+        case kSurfaceLogoRGB:
             extension = "ppm";
             break;
         default:

mythtv/programs/mythgpucommflag/videosurface.h

@@ -25,7 +25,9 @@ typedef enum {
     kSurfaceWavelet,
     kSurfaceYUV,
     kSurfaceYUV2,
-    kSurfaceRGB
+    kSurfaceRGB,
+    kSurfaceLogoROI,
+    kSurfaceLogoRGB
 } VideoSurfaceType;
 
 class VideoSurface
@@ -44,6 +46,8 @@ class VideoSurface
     OpenCLDevice       *m_dev;
     uint32_t            m_width;
     uint32_t            m_height;
+    uint32_t            m_realWidth;
+    uint32_t            m_realHeight;
     PixelFormat         m_pixFmt;
     VdpVideoSurface     m_vdpSurface;
     vdpau_render_state  m_render;