Ticket #7214: 0001-Xilinx-NGcodec-hevc-ffmpeg-plugin.patch

new file .gitignore

@@ +1 @@
+*.a
+*.o
+*.o.*
+*.d
+*.def
+*.dll
+*.dylib
+*.exe
+*.exp
+*.gcda
+*.gcno
+*.h.c
+*.ilk
+*.lib
+*.pc
+*.pdb
+*.so
+*.so.*
+*.swp
+*.ver
+*_g
+\#*
+.\#*
+/.config
+/.version
+/ffmpeg
+/ffplay
+/ffprobe
+/ffserver
+/config.*
+/coverage.info
+/avversion.h
+/lcov/
+/src
+/mapfile
+doc/*
+bin/*
+include/*
+share/*
+build/*
+bavcodec/bsf_list.c
+libavformat/protocol_list.c
+libavutil/avconfig.h
+libavutil/ffversion.h
+libavcodec/bsf_list.c
+ffmpeg_build/*

ffmpeg.c

@@ -1175 +1175 @@
     ost->last_dropped = nb_frames == nb0_frames && next_picture;
 
   /* duplicates frame if needed */
+  if (ost->enc_ctx->codec->name) {
+      if (0 == strcmp(ost->enc_ctx->codec->name, "xlnx_hevc_enc")) {
+          nb_frames = 1;
+      }
+  }
   for (i = 0; i < nb_frames; i++) {
     AVFrame *in_picture;
     av_init_packet(&pkt);
@@ -3800 +3805 @@
     int i;
     int64_t opts_min = INT64_MAX;
     OutputStream *ost_min = NULL;
-
+    uint64_t frames_min = UINT64_MAX;
+    static int frame_based = 1;
+    int all_valid_ts = 1;
     for (i = 0; i < nb_output_streams; i++) {
         OutputStream *ost = output_streams[i];
         int64_t opts = ost->st->cur_dts == AV_NOPTS_VALUE ? INT64_MIN :
@@ -3809 +3816 @@
         if (ost->st->cur_dts == AV_NOPTS_VALUE)
             av_log(NULL, AV_LOG_DEBUG, "cur_dts is invalid (this is harmless if it occurs once at the start per stream)\n");
 
-        if (!ost->initialized && !ost->inputs_done)
+        if (!ost->initialized && !ost->inputs_done) {
             return ost;
-
-        if (!ost->finished && opts < opts_min) {
-            opts_min = opts;
-            ost_min  = ost->unavailable ? NULL : ost;
         }
+        if (frame_based && (!strcmp(ost->enc_ctx->codec->name, "xlnx_hevc_enc"))) {
+            if (!ost->finished && ost->frames_encoded < frames_min) {
+                frames_min = ost->frames_encoded;
+                ost_min  = ost->unavailable ? NULL : ost;
+            }
+            if (opts == INT64_MIN) {
+                all_valid_ts = 0;
+            }
+        } else {
+            if (!ost->finished && opts < opts_min) {
+                opts_min = opts;
+                ost_min  = ost->unavailable ? NULL : ost;
+            }
+        }
+    }
+    if (frame_based && all_valid_ts) {
+        frame_based = 0;
     }
     return ost_min;
 }

libavcodec/Makefile

@@ -49 +49 @@
        utils.o                                                          \
        vorbis_parser.o                                                  \
        xiph.o                                                           \
+           xlnx_h264/hw_h264.o                                                                                          \
+           xlnx_h264/h264e_csregs.o                                                                             \
+           xlnx_h264/h264e_sdaccel.o                                                                            \
+           xlnx_h264/h264e_sdaccel_utils.o
 
 # subsystems
 OBJS-$(CONFIG_AANDCTTABLES)            += aandcttab.o
@@ -454 +458 @@
 OBJS-$(CONFIG_PBM_ENCODER)             += pnmenc.o
 OBJS-$(CONFIG_PCX_DECODER)             += pcx.o
 OBJS-$(CONFIG_PCX_ENCODER)             += pcxenc.o
+OBJS-$(CONFIG_XLNX_H264_ENCODER)       += xlnx_h264_enc.o  
 OBJS-$(CONFIG_PGM_DECODER)             += pnmdec.o pnm.o
 OBJS-$(CONFIG_PGM_ENCODER)             += pnmenc.o
 OBJS-$(CONFIG_PGMYUV_DECODER)          += pnmdec.o pnm.o
@@ -916 +921 @@
 OBJS-$(CONFIG_LIBXAVS_ENCODER)            += libxavs.o
 OBJS-$(CONFIG_LIBXVID_ENCODER)            += libxvid.o
 OBJS-$(CONFIG_LIBZVBI_TELETEXT_DECODER)   += libzvbi-teletextdec.o ass.o
+OBJS-$(CONFIG_XLNX_HEVC_ENCODER)          += xlnx_hevc_enc.o xlnx_hw/xlnx_hw.o \
+                                             xlnx_hw/common/src/xlnx_hrm.o \
+                                             xlnx_hw/common/src/xcl.o \
+                                             xlnx_hw/common/src/xlnx_queue.o \
+                                             xlnx_hw/common/src/xlnx_ts_queue.o \
+                                             xlnx_hw/hevc/enc_vu9p.o \
+                                             xlnx_hw/hevc/enc_vu9p_2c.o
 
 # parsers
 OBJS-$(CONFIG_AAC_LATM_PARSER)         += latm_parser.o

libavcodec/allcodecs.c

@@ -673 +673 @@
     REGISTER_ENCODER(VP8_VAAPI,         vp8_vaapi);
     REGISTER_DECODER(VP9_CUVID,         vp9_cuvid);
     REGISTER_DECODER(VP9_MEDIACODEC,    vp9_mediacodec);
+    REGISTER_ENCODER(XLNX_HEVC,         xlnx_hevc);
+    REGISTER_ENCODER(XLNX_H264,         xlnx_h264);
 
     /* parsers */
     REGISTER_PARSER(AAC,                aac);

new file libavcodec/xlnx_h264/Makefile

@@ +1 @@
+OBJS                            += xlnx_h264/hw_h264.o \
+                                                                   xlnx_h264/h264e_csregs.o \
+                                                                   xlnx_h264/h264e_sdaccel.o \
+                                                                   xlnx_h264/h264e_sdaccel_utils.o

new file libavcodec/xlnx_h264/h264e_csregs.c

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#include <stdio.h>
+#include "h264e_csregs.h"
+#include "h264e_csregs_private.h"
+
+void h264eSetControlRegsDefaults(H264E_CFG_PVT* ph264eCFG);
+void h264eSetReg(uint *puiRegFile, int regOffs, int regFieldHI, int regFieldLO, int regFieldVal);
+uint h264eGetReg(uint *puiRegFile, int regOffs, int regFieldHI, int regFieldLO);
+
+void h264ePrintStatusRegs(int* pRegFile)
+{
+    for(int i = 0; i < 64; i++)
+    {
+        H264EMSG(printf("H264-E CORE STATUS[0x%.2X] = 0x%.8X\n", i, pRegFile[i]));
+    }
+    return;
+}
+
+void h264ePrintControlRegs(int* pRegFile)
+{
+    for(int i = 0; i < 64; i++)
+    {
+        H264EMSG(printf("CREG[0x%.2X] = 0x%.8X\n", i, pRegFile[i]));
+    }
+    return;
+}
+
+int h264eGetNalSize(int* pRegFile)
+{
+    return pRegFile[CSREG_PICLEN];
+}
+
+int h264eGetErrorCode(int* pRegFile)
+{
+    return pRegFile[CSREG_ERRID];
+}
+
+int h264eGetDbgStatus(int* pRegFile)
+{
+    return pRegFile[CSREG_DEBUG];
+}
+
+void h264eSetControlRegs(H264E_USR* ph264eUsr, char* CRegsFile)
+{
+    H264E_CFG_PVT  h264eCFGPvt;
+    H264E_CFG_PVT* ph264eCFG;
+    if(!ph264eUsr->bUseUserContextAsCfg)
+    {
+        ph264eCFG = &h264eCFGPvt;
+        h264eSetControlRegsDefaults(ph264eCFG);
+        ph264eCFG->X = ph264eUsr->FrameX;
+        ph264eCFG->Y = ph264eUsr->FrameY;
+        if(ph264eUsr->HWOptimizeScanFormat)
+        {
+             ph264eCFG->VInputFormat = ph264eUsr->ScanFormat == H264E_MACROBLOCK ? 0 : 2;
+        }
+        else
+        {
+            ph264eCFG->VInputFormat = ph264eUsr->ScanFormat == H264E_MACROBLOCK       ? 0 :
+                                      ph264eUsr->ScanFormat == H264E_PLANAR_PACKED    ? 1 :
+                                      ph264eUsr->ScanFormat == H264E_PLANAR_SEPARATED ? 1 : 2;
+        }
+        ph264eCFG->Timescale            = ph264eUsr->Timescale;
+        ph264eCFG->NUIT                 = ph264eUsr->Nuit;
+        ph264eCFG->Level                = ph264eUsr->Level;
+        ph264eCFG->IPeriod              = ph264eUsr->GOPLength         == -1 ? ph264eCFG->IPeriod              : ph264eUsr->GOPLength;
+        ph264eCFG->IRefColumns          = ph264eUsr->IntraRefreshCol   == -1 ? ph264eCFG->IRefColumns          : ph264eUsr->IntraRefreshCol;
+        ph264eCFG->nMBRowsInSliceReq    = ph264eUsr->MBRowsPerSlice    == -1 ? ph264eCFG->nMBRowsInSliceReq    : ph264eUsr->MBRowsPerSlice;
+        ph264eCFG->UseDeBlockingFilter  = ph264eUsr->DeblockingFilter  == -1 ? ph264eCFG->UseDeBlockingFilter  :
+                                          ph264eUsr->DeblockingFilter  ==  0 ? 0                               :
+                                          ph264eUsr->DBRestrict        ==  1 ? 2                               : 1;
+        ph264eCFG->InterMVSliceRestrict = ph264eUsr->MVRestrict        == -1 ? ph264eCFG->InterMVSliceRestrict : (bool)ph264eUsr->MVRestrict;
+        ph264eCFG->QPI                  = ph264eUsr->QPI               == -1 ? ph264eCFG->QPI                  : ph264eUsr->QPI;
+        ph264eCFG->QPP                  = ph264eUsr->QPP               == -1 ? ph264eCFG->QPP                  : ph264eUsr->QPP;
+        ph264eCFG->QPChrOffs            = ph264eUsr->QPChrOffset       == -1 ? ph264eCFG->QPChrOffs            : ph264eUsr->QPChrOffset;
+        ph264eCFG->Rate                 = ph264eUsr->Bitrate;
+        ph264eCFG->nMBsInBasicUnit      = ph264eUsr->BUSize            == -1 ? ph264eCFG->nMBsInBasicUnit      : ph264eUsr->BUSize;
+        ph264eCFG->FIFOSize             = ph264eUsr->RCFIFOSize        == -1 ? ph264eCFG->FIFOSize             : ph264eUsr->RCFIFOSize;
+        ph264eCFG->bCABAC               = ph264eUsr->UseCABAC          == -1 ? ph264eCFG->bCABAC               : (bool)ph264eUsr->UseCABAC;
+    }
+    else
+    {
+        ph264eCFG = (H264E_CFG_PVT*)ph264eUsr->UserContext;
+        if(ph264eUsr->HWOptimizeScanFormat)
+        {
+            ph264eCFG->VInputFormat = ph264eUsr->ScanFormat == H264E_MACROBLOCK ? 0 : 2;
+        }
+        if(ph264eCFG->VInputFormat == 2 && ph264eCFG->UseFrameBuffer)
+        {
+            ph264eCFG->VInputFormat = 3;
+        }
+    }
+    uint vidinres_reg;
+    uint vidinscan_reg;
+    uint vidints_reg;
+    uint vidinuit_reg;
+    uint vidinlevel_reg;
+    uint vidinvuist_reg;
+    uint gopattr_reg;
+    uint sliceattr_reg;
+    uint qpsliceattr_reg;
+    uint rcattr_reg;
+    uint rcqplims_reg;
+    uint rcxopts_reg;
+    uint predattr_reg;
+    uint encopts_reg;
+    uint xpredattr_reg;
+    uint xrcxopts_reg;
+    uint xrcxopts_pid_reg[5];
+    uint instruction_reg;
+    uint xtmemptr_reg;
+    static const int LevelLUT[52] = {
+       0, 0, 0, 0, 0, 0, 0, 0,  /*  0-7 */
+       0, 1, 0, 2,              /*  8-11 */
+       3, 4, 0, 0, 0, 0, 0, 0,  /* 12-19 */
+       5, 6, 7, 0, 0, 0, 0, 0,  /* 20-27 */
+       0, 0, 8, 9,10, 0, 0, 0,  /* 28-35 */
+       0, 0, 0, 0,11,12,13, 0,  /* 36-43 */
+       0, 0, 0, 0, 0, 0,14,15   /* 44-51 */
+    };
+    vidinres_reg = 0;
+    h264eSetReg(&vidinres_reg, 0, FRAME_X_HI, FRAME_X_LO, ph264eCFG->X);
+    h264eSetReg(&vidinres_reg, 0, FRAME_Y_HI, FRAME_Y_LO, ph264eCFG->Y * (ph264eCFG->bInterlaced ? 2 : 1));
+    vidinscan_reg = 0;
+#ifdef USE_UHT_CORE
+    h264eSetReg(&vidinscan_reg, 0, SCAN_FORMAT_HI, SCAN_FORMAT_LO, ph264eCFG->csp);
+#else
+    h264eSetReg(&vidinscan_reg, 0, SCAN_FORMAT_HI, SCAN_FORMAT_LO, (ph264eCFG->csp == YUV420 ? ph264eCFG->VInputFormat : 4 + ph264eCFG->VInputFormat));
+#endif
+    h264eSetReg(&vidinscan_reg, 0, PSP_HI,         PSP_LO,         ph264eCFG->bpp);
+    h264eSetReg(&vidinscan_reg, 0, FRAMEDROP_FLAG, FRAMEDROP_FLAG, (int) ph264eCFG->bFrameDrop);
+    vidints_reg = 0;
+    h264eSetReg(&vidints_reg, 0, TIMESCALE_HI, TIMESCALE_LO, ph264eCFG->Timescale);
+    vidinuit_reg = 0;
+    h264eSetReg(&vidinuit_reg, 0, NUIT_HI, NUIT_LO, ph264eCFG->NUIT);
+    vidinlevel_reg = 0;
+    h264eSetReg(&vidinlevel_reg, 0, LEVEL_HI, LEVEL_LO, LevelLUT[ph264eCFG->Level]);
+    vidinvuist_reg = 0;
+    h264eSetReg(&vidinvuist_reg, 0, V_VIDEO_SIGNAL_TYPE_PRESENT_FLAG, V_VIDEO_SIGNAL_TYPE_PRESENT_FLAG, (ph264eCFG->VUI_VideoFormat > 0 || ph264eCFG->VUI_VideoFullRangeFlag ? 1 : 0));
+    h264eSetReg(&vidinvuist_reg, 0, V_VIDEO_FORMAT_HI,                V_VIDEO_FORMAT_LO,                (ph264eCFG->VUI_VideoFormat > 0 ? ph264eCFG->VUI_VideoFormat : 0));
+    h264eSetReg(&vidinvuist_reg, 0, V_VIDEO_FULL_RANGE_FLAG,          V_VIDEO_FULL_RANGE_FLAG,           ph264eCFG->VUI_VideoFullRangeFlag);
+    gopattr_reg = 0;
+    if (ph264eCFG->IPeriod == 0xFFFFFFFF)
+    {
+       h264eSetReg(&gopattr_reg, 0, IPERIOD_HI, IPERIOD_LO, 0);
+    }
+    else
+    {
+       h264eSetReg(&gopattr_reg, 0, IPERIOD_HI, IPERIOD_LO, ph264eCFG->IPeriod);
+    }
+    h264eSetReg(&gopattr_reg, 0, IREFCOLS_HI, IREFCOLS_LO, ph264eCFG->IRefColumns);
+    sliceattr_reg = 0;
+    h264eSetReg(&sliceattr_reg, 0, SLICE_MBROWS_HI,      SLICE_MBROWS_L0,      ph264eCFG->nMBRowsInSliceReq);
+    h264eSetReg(&sliceattr_reg, 0, DB_EN_FLAG,           DB_EN_FLAG,           ph264eCFG->UseDeBlockingFilter  > 0 ? 1 : 0);
+#ifndef USE_UHT_CORE
+    h264eSetReg(&sliceattr_reg, 0, DB_SLCBNDRS_FLAG,     DB_SLCBNDRS_FLAG,     ph264eCFG->UseDeBlockingFilter == 2 ? 1 : 0);
+#endif
+    h264eSetReg(&sliceattr_reg, 0, DB_ALPHA_HI,          DB_ALPHA_LO,          ph264eCFG->sliceAlphaOffset);
+    h264eSetReg(&sliceattr_reg, 0, DB_BETA_HI,           DB_BETA_LO,           ph264eCFG->sliceBetaOffset);
+#ifndef USE_UHT_CORE
+    h264eSetReg(&sliceattr_reg, 0, INTERSLICECROSS_FLAG, INTERSLICECROSS_FLAG, (int) ph264eCFG->InterMVSliceRestrict);
+#endif
+    qpsliceattr_reg = 0;
+    h264eSetReg(&qpsliceattr_reg, 0, QPINIT_I_HI,  QPINIT_I_LO,  ph264eCFG->QPI + ph264eCFG->QpBdOffset);
+    h264eSetReg(&qpsliceattr_reg, 0, QPINIT_P_HI,  QPINIT_P_LO,  ph264eCFG->QPP + ph264eCFG->QpBdOffset);
+    h264eSetReg(&qpsliceattr_reg, 0, QPCHROFFS_HI, QPCHROFFS_LO, ph264eCFG->QPChrOffs);
+    h264eSetReg(&qpsliceattr_reg, 0, QPTHR_HI,     QPTHR_LO,     ph264eCFG->bTQThresholding ? 1 : 0);
+    rcattr_reg = 0;
+    h264eSetReg(&rcattr_reg, 0, RCBITRATE_HI,    RCBITRATE_LO,    ph264eCFG->Rate    >> 3);
+    h264eSetReg(&rcattr_reg, 0, RCMAXBITRATE_HI, RCMAXBITRATE_LO, ph264eCFG->MaxRate >> 3);
+    rcqplims_reg = 0;
+    h264eSetReg(&rcqplims_reg, 0, QPI_MIN_HI, QPI_MIN_LO, ph264eCFG->QPI_min + ph264eCFG->QpBdOffset);
+    h264eSetReg(&rcqplims_reg, 0, QPI_MAX_HI, QPI_MAX_LO, ph264eCFG->QPI_max + ph264eCFG->QpBdOffset);
+    h264eSetReg(&rcqplims_reg, 0, QPP_MIN_HI, QPP_MIN_LO, ph264eCFG->QPP_min + ph264eCFG->QpBdOffset);
+    h264eSetReg(&rcqplims_reg, 0, QPP_MAX_HI, QPP_MAX_LO, ph264eCFG->QPP_max + ph264eCFG->QpBdOffset);
+    rcxopts_reg = 0;
+    h264eSetReg(&rcxopts_reg, 0, BUSIZE_HI, BUSIZE_LO, (ph264eCFG->nMBsInBasicUnit < 0 ? 0 : ph264eCFG->nMBsInBasicUnit));
+    if (ph264eCFG->RCMethod == RCM_PID)
+    {
+        h264eSetReg(&rcxopts_reg, 0, RCPERIOD_HI, RCPERIOD_LO, ph264eCFG->FIFOSize >> 3);
+    }
+    else
+    {
+        h264eSetReg(&rcxopts_reg, 0, RCPERIOD_HI, RCPERIOD_LO, ph264eCFG->RCPeriod);
+    }
+    encopts_reg = 0;
+    h264eSetReg(&encopts_reg, 0, CABAC_FLAG,      CABAC_FLAG,      (int) ph264eCFG->bCABAC);
+    h264eSetReg(&encopts_reg, 0, FLUSHFRAME_FLAG, FLUSHFRAME_FLAG, 1);
+    h264eSetReg(&encopts_reg, 0, NALTRZWORD_FLAG, NALTRZWORD_FLAG, (ph264eCFG->NALTrZeros == 2 ? 1 : 0));
+#ifdef USE_UHT_CORE
+    h264eSetReg(&encopts_reg, 0, ACTCENGS_HI,     ACTCENGS_LO,     (int) ph264eCFG->nCEngines);
+#endif
+    predattr_reg = 0;
+    h264eSetReg(&predattr_reg, 0, MEVBS_P16X16_FLAG, MEVBS_P16X16_FLAG, ph264eCFG->bVBSEnableMask[0]);
+    h264eSetReg(&predattr_reg, 0, MEVBS_P16X8_FLAG,  MEVBS_P16X8_FLAG,  ph264eCFG->bInterVBSenable & ph264eCFG->bVBSEnableMask[1]);
+    h264eSetReg(&predattr_reg, 0, MEVBS_P8X16_FLAG,  MEVBS_P8X16_FLAG,  ph264eCFG->bInterVBSenable & ph264eCFG->bVBSEnableMask[2]);
+    h264eSetReg(&predattr_reg, 0, MEVBS_P8X8_FLAG,   MEVBS_P8X8_FLAG,   ph264eCFG->bInterVBSenable & ph264eCFG->bVBSEnableMask[3]);
+    h264eSetReg(&predattr_reg, 0, MEPXLAC_HI,        MEPXLAC_LO,        ph264eCFG->InterAccuracy/2);
+    h264eSetReg(&predattr_reg, 0, XTRACOST_FLAG,     XTRACOST_FLAG,    !ph264eCFG->bVBSOptions[VBS_USE_XTRA_PART_COST]);
+    h264eSetReg(&predattr_reg, 0, XTRACOSTP8x8_FLAG, XTRACOSTP8x8_FLAG, ph264eCFG->bVBSOptions[VBS_USE_XTRA_PART_COST_P8x8]);
+    h264eSetReg(&predattr_reg, 0, INTRA_IN_P_FLAG,   INTRA_IN_P_FLAG,   ph264eCFG->bUseIntraInInter);
+    h264eSetReg(&predattr_reg, 0, I16x16_FLAG,       I16x16_FLAG,      !ph264eCFG->bI16x16Dis);
+    h264eSetReg(&predattr_reg, 0, I4x4_FLAG,         I4x4_FLAG,        !ph264eCFG->bI4x4Dis);
+    int i, mask;
+    for (mask=0, i=0; i<4; i++) mask |= ((int) ph264eCFG->bIntraLuma16x16Mask[i]) << i;
+    h264eSetReg(&predattr_reg, 0, I16x16_MASK_HI,    I16x16_MASK_LO,    mask);
+    for (mask=0, i=0; i<4; i++) mask |= ((int) ph264eCFG->bIntraChromaMask[i]) << i;
+    h264eSetReg(&predattr_reg, 0, ICHR_MASK_HI,      ICHR_MASK_LO,      mask);
+    for (mask=0, i=0; i<9; i++) mask |= ((int) ph264eCFG->bIntraLuma4x4Mask[i]) << i;
+    h264eSetReg(&predattr_reg, 0, I4x4_MASK_HI,      I4x4_MASK_LO,      mask);
+    xpredattr_reg = 0;
+    h264eSetReg(&xpredattr_reg, 0, SRX_HI,   SRX_LO,   ph264eCFG->InterSearchRangeWidth/8);
+    h264eSetReg(&xpredattr_reg, 0, SRY_HI,   SRY_LO,   ph264eCFG->InterSearchRangeHeight/2);
+    h264eSetReg(&xpredattr_reg, 0, VIMPW_HI, VIMPW_LO, ph264eCFG->MPPriorities[0]);
+    h264eSetReg(&xpredattr_reg, 0, VIMPR_HI, VIMPR_LO, ph264eCFG->MPPriorities[1]);
+    h264eSetReg(&xpredattr_reg, 0, VOMPW_HI, VOMPW_LO, ph264eCFG->MPPriorities[2]);
+    h264eSetReg(&xpredattr_reg, 0, REFWR_HI, REFWR_LO, ph264eCFG->MPPriorities[3]);
+    xrcxopts_reg = 0;
+    h264eSetReg(&xrcxopts_reg, 0, QPADJSTRENGTH_HI,  QPADJSTRENGTH_LO,  ph264eCFG->AQStrength);
+    h264eSetReg(&xrcxopts_reg, 0, USEINITQP_FLAG,    USEINITQP_FLAG,    ph264eCFG->bRcExtraOpt[0]);
+    h264eSetReg(&xrcxopts_reg, 0, USESCENECUT_FLAG,  USESCENECUT_FLAG,  ph264eCFG->bRcExtraOpt[1]);
+    h264eSetReg(&xrcxopts_reg, 0, USERCOVERDUE_FLAG, USERCOVERDUE_FLAG, ph264eCFG->bRcExtraOpt[2]);
+    h264eSetReg(&xrcxopts_reg, 0, STRICTCBR_FLAG,    STRICTCBR_FLAG,    ph264eCFG->bRcExtraOpt[3]);
+    h264eSetReg(&xrcxopts_reg, 0, DISMADTH_FLAG,     DISMADTH_FLAG,     ph264eCFG->bRcExtraOpt[4]);
+    h264eSetReg(&xrcxopts_reg, 0, VBRF_HI,           VBRF_LO,           ph264eCFG->rcVbrFactor);
+    h264eSetReg(&xrcxopts_reg, 0, ICPBRDFR_HI,       ICPBRDFR_LO,       ph264eCFG->initialCbpRdInfr);
+    for (int i=0; i<5; i++)
+    {
+        xrcxopts_pid_reg[i] = ph264eCFG->xrcxopts_pid[i];
+    }
+    instruction_reg = 0;
+    h264eSetReg(&instruction_reg, 0, SWCLR_FLAG,     SWCLR_FLAG,     0);
+    h264eSetReg(&instruction_reg, 0, SWSUSPEND_FLAG, SWSUSPEND_FLAG, 0);
+    xtmemptr_reg = 0;
+    h264eSetReg(&xtmemptr_reg, 0, XTMEMPTR_HI, XTMEMPTR_LO, ph264eCFG->XtmemPtr);
+    h264eSetReg((uint*)CRegsFile, CSREG_VIDINRES,      31, 0, vidinres_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_VIDINSCAN,     31, 0, vidinscan_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_VIDINTS,       31, 0, vidints_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_VIDINUIT,      31, 0, vidinuit_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_VIDINLEVEL,    31, 0, vidinlevel_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_VIDINVUIST,    31, 0, vidinvuist_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_GOPATTR,       31, 0, gopattr_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_SLICEATTR,     31, 0, sliceattr_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_QPSLICEATTR,   31, 0, qpsliceattr_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_RCATTR,        31, 0, rcattr_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_RCQPLIMS,      31, 0, rcqplims_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_RCXOPTS,       31, 0, rcxopts_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_PREDATTR,      31, 0, predattr_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_ENCOPTS,       31, 0, encopts_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_XPREDATTR,     31, 0, xpredattr_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_XRCXOPTS,      31, 0, xrcxopts_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_INSTRUCTION,   31, 0, instruction_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_XTMEMPTR,      31, 0, xtmemptr_reg);
+    h264eSetReg((uint*)CRegsFile, CSREG_XRCXOPTS_PID0, 31, 0, xrcxopts_pid_reg[0]);
+    h264eSetReg((uint*)CRegsFile, CSREG_XRCXOPTS_PID1, 31, 0, xrcxopts_pid_reg[1]);
+    h264eSetReg((uint*)CRegsFile, CSREG_XRCXOPTS_PID2, 31, 0, xrcxopts_pid_reg[2]);
+    h264eSetReg((uint*)CRegsFile, CSREG_XRCXOPTS_PID3, 31, 0, xrcxopts_pid_reg[3]);
+    h264eSetReg((uint*)CRegsFile, CSREG_XRCXOPTS_PID4, 31, 0, xrcxopts_pid_reg[4]);
+    if(ph264eCFG->PrintControlRegisters)
+    {
+        h264ePrintControlRegs((int*)CRegsFile);
+    }
+    return;
+}
+
+void h264eSetControlRegsDefaults(H264E_CFG_PVT* ph264eCFG)
+{
+    ph264eCFG->X                      = 1280;
+    ph264eCFG->Y                      = 720;
+    ph264eCFG->bInterlaced            = false;
+    ph264eCFG->VInputFormat           = VIPLANAR;
+    ph264eCFG->bpp                    = 8;
+    ph264eCFG->csp                    = YUV420;
+    ph264eCFG->Timescale              = 50000;
+    ph264eCFG->NUIT                   = 1000;
+    ph264eCFG->Level                  = 31;
+    ph264eCFG->QPI                    = 25;
+    ph264eCFG->QPP                    = 25;
+    ph264eCFG->QPChrOffs              = 0;
+    ph264eCFG->QpBdOffset             = 0;
+    ph264eCFG->bTQThresholding        = true;
+    ph264eCFG->UseDeBlockingFilter    = 1;
+    ph264eCFG->sliceAlphaOffset       = 0;
+    ph264eCFG->sliceBetaOffset        = 0;
+    ph264eCFG->nMBRowsInSliceReq      = 0;
+    ph264eCFG->InterSearchRangeWidth  = 32;
+    ph264eCFG->InterSearchRangeHeight = 20;
+    ph264eCFG->InterAccuracy          = 4;
+    ph264eCFG->InterMVSliceRestrict   = false;
+    ph264eCFG->bInterVBSenable        = true;
+    ph264eCFG->bUseIntraInInter       = true;
+    ph264eCFG->NALTrZeros             = 0;
+    ph264eCFG->NALFlushOnFrame        = 1;
+    ph264eCFG->bFrameDrop             = false;
+    ph264eCFG->VUI_VideoFormat        = -1;
+    ph264eCFG->VUI_VideoFullRangeFlag = 0;
+#ifdef USE_CABAC
+    ph264eCFG->bCABAC                 = true;
+#else
+    ph264eCFG->bCABAC                 = false;
+#endif
+    ph264eCFG->AQStrength             = 4;
+    ph264eCFG->XtmemPtr               = 0;
+    ph264eCFG->nCEngines              = 1;
+    ph264eCFG->MPPriorities[0]        = 3;
+    ph264eCFG->MPPriorities[1]        = 2;
+    ph264eCFG->MPPriorities[2]        = 1;
+    ph264eCFG->MPPriorities[3]        = 1;
+    ph264eCFG->Rate                   = 0;
+    ph264eCFG->FIFOSize               = 0;
+    ph264eCFG->MaxRate                = 0;
+    ph264eCFG->nMBsInBasicUnit        = 0;
+    ph264eCFG->RCPeriod               = 0;
+    ph264eCFG->QPI_min                = 0;
+    ph264eCFG->QPI_max                = 51;
+    ph264eCFG->QPP_min                = 0;
+    ph264eCFG->QPP_max                = 51;
+    ph264eCFG->IPeriod                = 30;
+    ph264eCFG->IRefColumns            = 0;
+    ph264eCFG->RCMethod               = RCM_PID;
+    ph264eCFG->initialCbpRdInfr       = 0;
+    ph264eCFG->rcVbrFactor            = 16;
+    for (int i = 1; i < RC_EXTRA_OPT_MASK_WIDTH; i++)
+    {
+        ph264eCFG->bRcExtraOpt[i] = false;
+    }
+    for(int i = 0; i < 7; i++)
+    {
+        if(i < 4)
+        {
+            ph264eCFG->bVBSEnableMask[i] = true;
+        }
+        else
+        {
+            ph264eCFG->bVBSEnableMask[i] = false;
+        }
+        if(i < 5)
+        {
+            ph264eCFG->bVBSOptions[i] = true;
+        }
+        else
+        {
+            ph264eCFG->bVBSOptions[i] = false;
+        }
+    }
+    ph264eCFG->bVBSOptions[3] = false;
+    ph264eCFG->bI16x16Dis     = 0;
+    ph264eCFG->bI4x4Dis       = false;
+    for(int i = 0; i < 9; i++)
+    {
+        if(i < 4)
+        {
+            ph264eCFG->bIntraLuma16x16Mask[i] = true;
+            ph264eCFG->bIntraChromaMask[i]    = true;
+        }
+        ph264eCFG->bIntraLuma4x4Mask[i] = true;
+    }
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[0], 0, PID_KU_IONPY_HI,             PID_KU_IONPY_LO,              PID_KU_IONPY);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[0], 0, PID_KU_IP_HI,                PID_KU_IP_LO,                 PID_KU_IP);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[0], 0, PID_KI_UPD_IONPY_HI,         PID_KI_UPD_IONPY_LO,          PID_KI_UPD_IONPY);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[0], 0, PID_KI_UPD_IP_HI,            PID_KI_UPD_IP_LO,             PID_KI_UPD_IP);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[0], 0, PID_USE_IBU_THR_FLAG,        PID_USE_IBU_THR_FLAG,         PID_USE_IBU_THR);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[0], 0, PID_FIFO_STATUS_TARGET_HI,   PID_FIFO_STATUS_TARGET_LO,    PID_FIFO_STATUS_TARGET);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_USE_COLOCATED_FLAG,      PID_USE_COLOCATED_FLAG,       PID_USE_COLOCATED);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_IFRAME_BUFFER_FACTOR_HI, PID_IFRAME_BUFFER_FACTOR_LO,  PID_IFRAME_BUFFER_FACTOR);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_QP_STEP_FR_HI,           PID_QP_STEP_FR_LO,            PID_QP_STEP_FR);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_QP_STEP_IBU_HI,          PID_QP_STEP_IBU_LO,           PID_QP_STEP_IBU);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_FSTATUS_0_HI,            PID_FSTATUS_0_LO,             PID_FSTATUS_0);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_THRFRMLIM_0_FLAG,        PID_THRFRMLIM_0_FLAG,         PID_THRFRMLIM_0);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_FSTATUS_1_HI,            PID_FSTATUS_1_LO,             PID_FSTATUS_1);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[1], 0, PID_THRFRMLIM_1_FLAG,        PID_THRFRMLIM_1_FLAG,         PID_THRFRMLIM_1);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_FSTATUS_2_HI,            PID_FSTATUS_2_LO,             PID_FSTATUS_2);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_THRFRMLIM_2_FLAG,        PID_THRFRMLIM_2_FLAG,         PID_THRFRMLIM_2);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_FSTATUS_3_HI,            PID_FSTATUS_3_LO,             PID_FSTATUS_3);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_THRFRMLIM_3_FLAG,        PID_THRFRMLIM_3_FLAG,         PID_THRFRMLIM_3);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_FSTATUS_4_HI,            PID_FSTATUS_4_LO,             PID_FSTATUS_4);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_THRFRMLIM_4_FLAG,        PID_THRFRMLIM_4_FLAG,         PID_THRFRMLIM_4);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_QP_STEP_0_HI,            PID_QP_STEP_0_LO,             PID_QP_STEP_0);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[2], 0, PID_QP_THR_0_HI,             PID_QP_THR_0_LO,              PID_QP_THR_0);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_STEP_1_HI,            PID_QP_STEP_1_LO,             PID_QP_STEP_1);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_THR_1_HI,             PID_QP_THR_1_LO,              PID_QP_THR_1);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_STEP_2_HI,            PID_QP_STEP_2_LO,             PID_QP_STEP_2);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_THR_2_HI,             PID_QP_THR_2_LO,              PID_QP_THR_2);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_STEP_3_HI,            PID_QP_STEP_3_LO,             PID_QP_STEP_3);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_THR_3_HI,             PID_QP_THR_3_LO,              PID_QP_THR_3);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_STEP_4_HI,            PID_QP_STEP_4_LO,             PID_QP_STEP_4);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[3], 0, PID_QP_THR_4_HI,             PID_QP_THR_4_LO,              PID_QP_THR_4);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_RAMP_METHOD_HI,          PID_RAMP_METHOD_LO,           PID_RAMP_METHOD);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_USE_KD_FLAG,             PID_USE_KD_FLAG,              PID_USE_KD);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_USE_KI_UPD_FLAG,         PID_USE_KI_UPD_FLAG,          PID_USE_KI_UPD);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_USE_KI_UPDM_FLAG,        PID_USE_KI_UPDM_FLAG,         PID_USE_KI_UPDM);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_USE_SCENE_THR_FLAG,      PID_USE_SCENE_THR_FLAG,       PID_USE_SCENE_THR);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_USE_AVG_QP_THR_FLAG,     PID_USE_AVG_QP_THR_FLAG,      PID_USE_AVG_QP_THR);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_AVG_QP_THR_HI,           PID_AVG_QP_THR_LO,            PID_AVG_QP_THR);
+    h264eSetReg(&ph264eCFG->xrcxopts_pid[4], 0, PID_SCENE_THR_HI,            PID_SCENE_THR_LO,             PID_SCENE_THR);
+    return;
+}
+
+void h264eSetReg(uint      *puiRegFile,
+                 int        regOffs,
+                 int        regFieldHI,
+                 int        regFieldLO,
+                 int        regFieldVal)
+{
+    uint FieldMask;
+    uint regVal;
+    uint regFieldValP;
+    FieldMask      = 0xFFFFFFFF;
+    FieldMask    >>= ((8*sizeof(uint)) - (regFieldHI - regFieldLO + 1));
+    regFieldValP   = regFieldVal & FieldMask;
+    FieldMask    <<= regFieldLO;
+    regFieldValP <<= regFieldLO;
+    FieldMask      = ~FieldMask;
+    regVal         = puiRegFile[regOffs];
+    regVal        &= FieldMask;
+    regVal        |= regFieldValP;
+    puiRegFile[regOffs] = regVal;
+    return;
+}
+
+uint h264eGetReg(uint *puiRegFile,
+                 int   regOffs,
+                 int   regFieldHI,
+                 int   regFieldLO)
+{
+    uint retval;
+    uint FieldMask;
+    FieldMask     = 0xFFFFFFFF;
+    FieldMask   >>= ((8*sizeof(uint)) - (regFieldHI - regFieldLO + 1));
+    retval        = puiRegFile[regOffs];
+    retval      >>= regFieldLO;
+    retval       &= FieldMask;
+    return retval;
+}

new file libavcodec/xlnx_h264/h264e_csregs.h

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#ifndef _H264E_SDACCEL_CSREGS_H_
+#define _H264E_SDACCEL_CSREGS_H_
+
+#include "h264e_sdaccel.h"
+
+void h264eSetControlRegs   (H264E_USR* ph264eUsr, char* CRegsFile);
+void h264ePrintStatusRegs  (int* pRegFile);
+void h264ePrintControlRegs (int* pRegFile);
+int  h264eGetNalSize       (int* pRegFile);
+int  h264eGetErrorCode     (int* pRegFile);
+int  h264eGetDbgStatus     (int* pRegFile);
+
+#endif // _H264E_SDACCEL_CSREGS_H_
+
+

new file libavcodec/xlnx_h264/h264e_csregs_private.h

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#ifndef _H264E_CSREGS_PRIVATE_H_
+#define _H264E_CSREGS_PRIVATE_H_
+
+#include <stdbool.h>
+
+typedef unsigned int uint;
+
+typedef struct _QPs {
+
+   int QPIPic;
+   int QPPPic;
+   int dQPISlice;
+   int dQPPSlice;
+   int dQPIMB;
+   int dQPPMB;
+   int QPEffectiveI;
+   int QPEffectiveP;
+   int QpBdOffset;
+
+} QPs;
+
+#define RC_EXTRA_OPT_MASK_WIDTH 5
+
+typedef struct _H264E_CFG_PVT {
+
+   int      X;
+   int      Y;
+   int      nMBRowsInSlice;
+   int      nMBRowsInSliceReq;
+   int      QPI;
+   int      QPP;
+   int      QPChrOffs;
+   int      QpBdOffset;
+   bool     bTQThresholding;
+   unsigned IPeriod;
+   unsigned IRefColumns;
+   int      UseDeBlockingFilter;
+   int      sliceAlphaOffset;
+   int      sliceBetaOffset;
+   float    FrameRate;
+   unsigned Timescale;
+   unsigned NUIT;
+   int      nMBsInBasicUnit;
+   unsigned RCMethod;
+   int      RCPeriod;
+   int      Rate;
+   int      MaxRate;
+   int      FIFOSize;
+   int      QPI_min;
+   int      QPI_max;
+   int      QPP_min;
+   int      QPP_max;
+   int      FramesToBeEncoded;
+   bool     bInputFromPipe;
+   int      pipeType;
+   int      InterSearchAlgorithm;
+   int      InterSearchRangeWidth;
+   int      InterSearchRangeHeight;
+   int      InterSRWOriginPolicy;
+   int      InterAccuracy;
+   int      InterQPELWindow;
+   int      InterPSKIPPolicy;
+   int      InterSADBits;
+   bool     InterMVSliceRestrict;
+   int      InterIMENeighborA;
+   int      InterFMENeighborA;
+   int      MPPriorities[4];
+   bool     bInterUseP0;
+   bool     bInterQP2S;
+   bool     bInterVBSenable;
+   bool     bVBSEnableMask[7];
+   bool     bVBSOptions[7];
+   int      bI16x16Dis;
+   bool     bIntraLuma16x16Mask[4];
+   bool     bI4x4Dis;
+   bool     bIntraLuma4x4Mask[9];
+   bool     bIntraChromaMask[4];
+   bool     bUseIntraInInter;
+   bool     bForceIPCMInIntra;
+   int      NALTrZeros;
+   int      NALFlushOnFrame;
+   bool     bFrameDrop;
+   QPs      h264eQPs;
+   int      VInputFormat;
+   int      bpp;
+   int      csp;
+   int      subWidthC;
+   int      subHeightC;
+   int      endianess;
+   int      Level;
+   unsigned XtmemPtr;
+   int      TBA;
+   bool     TBASMask[9];
+   void*    pPFIFile;
+   char     cClipInfo[2048];
+   int      Profile;
+   bool     bCABAC;
+   bool     bInterlaced;
+   bool     bIsTopFieldYUV;
+   bool     bBatch;
+   bool     bRcExtraOpt[RC_EXTRA_OPT_MASK_WIDTH];
+   char     rcExtraOpt[512];
+   unsigned rcVbrFactor;
+   unsigned estatsLevel;
+   char     rcRegsFileName[512];
+   FILE*    rcRegsFile;
+   bool     bUseSMAPData;
+   bool     bHrdEnable;
+   bool     bCbr;
+   uint     cpbSizeInKbps;
+   uint     initialCbpRdIn90Khz;
+   uint     initialCbpRdInfr;
+   bool     bRcHrd;
+   int      AQStrength;
+   uint     xrcxopts_pid[5];
+   bool     bDumpIOFrames;
+   uint     nNALOutWidth;
+   uint     nCEngines;
+   uint     YSecNomPad16;
+   int      TBDBG;
+   int      VUI_VideoFormat;
+   int      VUI_VideoFullRangeFlag;
+   int      UseHW;
+   char     KernelFile[2048];
+   char     TargetDevice[2048];
+   char     KernelName[2048];
+   FILE*    pYUVFile;
+   int      UsePlanarSeparated;
+   int      PackPlanarSeparated;
+   int      PlanarToInterleavedInHW;
+   int      HWBench;
+   int      UseFrameBuffer;
+   int      FrameOffset;
+   int      PrintControlRegisters;
+
+} H264E_CFG_PVT;
+
+#define CSREG_VIDINRES      0x00
+#define CSREG_VIDINSCAN     0x01
+#define CSREG_VIDINTS       0x02
+#define CSREG_VIDINUIT      0x03
+#define CSREG_VIDINLEVEL    0x04
+#define CSREG_VIDINVUIST    0x05
+#define CSREG_GOPATTR       0x10
+#define CSREG_SLICEATTR     0x11
+#define CSREG_QPSLICEATTR   0x12
+#define CSREG_RCATTR        0x13
+#define CSREG_RCQPLIMS      0x14
+#define CSREG_RCXOPTS       0x15
+#define CSREG_PREDATTR      0x16
+#define CSREG_ENCOPTS       0x17
+#define CSREG_XPREDATTR     0x3F
+#define CSREG_XRCXOPTS      0x3E
+#define CSREG_COREVER       0x3D
+#define CSREG_CORECFG       0x3C
+#define CSREG_XRCXOPTS_PID0 0x3B
+#define CSREG_XRCXOPTS_PID1 0x3A
+#define CSREG_XRCXOPTS_PID2 0x39
+#define CSREG_XRCXOPTS_PID3 0x38
+#define CSREG_XRCXOPTS_PID4 0x37
+#define CSREG_RCSTATUS      0x36
+#define CSREG_INSTRUCTION   0x20
+#define CSREG_STATUS        0x21
+#define CSREG_ERRID         0x22
+#define CSREG_IDRCNT        0x23
+#define CSREG_FRAMECNT      0x24
+#define CSREG_SEQLEN        0x25
+#define CSREG_PICLEN        0x26
+#define CSREG_SLICELEN      0x27
+#define CSREG_FRAMEDROPCNT  0x28
+#define CSREG_XTMEMPTR      0x30
+#define CSREG_DEBUG         0x3F
+#define FRAME_X_LO      0
+#define FRAME_X_HI     15
+#define FRAME_Y_LO     16
+#define FRAME_Y_HI     31
+#define SCAN_FORMAT_LO  0
+#define SCAN_FORMAT_HI  3
+#define PSP_LO          4
+#define PSP_HI          7
+#define FRAMEDROP_FLAG 16
+#define TIMESCALE_LO    0
+#define TIMESCALE_HI   31
+#define NUIT_LO         0
+#define NUIT_HI        31
+#define LEVEL_LO        0
+#define LEVEL_HI        4
+#define V_VIDEO_SIGNAL_TYPE_PRESENT_FLAG 0
+#define V_VIDEO_FORMAT_LO                4
+#define V_VIDEO_FORMAT_HI                7
+#define V_VIDEO_FULL_RANGE_FLAG          8
+#define IPERIOD_LO   0
+#define IPERIOD_HI  15
+#define IREFCOLS_LO 16
+#define IREFCOLS_HI 31
+#define SLICE_MBROWS_L0       0
+#define SLICE_MBROWS_HI      15
+#define DB_EN_FLAG           16
+#define DB_SLCBNDRS_FLAG     17
+#define DB_ALPHA_LO          20
+#define DB_ALPHA_HI          23
+#define DB_BETA_LO           24
+#define DB_BETA_HI           27
+#define INTERSLICECROSS_FLAG 28
+#define QPINIT_I_LO      0
+#define QPINIT_I_HI      7
+#define QPINIT_P_LO      8
+#define QPINIT_P_HI     15
+#define QPCHROFFS_LO    16
+#define QPCHROFFS_HI    23
+#define QPTHR_LO        24
+#define QPTHR_HI        25
+#define RCBITRATE_LO     0
+#define RCBITRATE_HI    15
+#define RCMAXBITRATE_LO 16
+#define RCMAXBITRATE_HI 31
+#define QPI_MIN_LO       0
+#define QPI_MIN_HI       7
+#define QPI_MAX_LO       8
+#define QPI_MAX_HI      15
+#define QPP_MIN_LO      16
+#define QPP_MIN_HI      23
+#define QPP_MAX_LO      24
+#define QPP_MAX_HI      31
+#define BUSIZE_LO        0
+#define BUSIZE_HI       15
+#define RCPERIOD_LO     16
+#define RCPERIOD_HI     31
+#define CABAC_FLAG       0
+#define FLUSHFRAME_FLAG  4
+#define NALTRZWORD_FLAG  8
+#define ACTCENGS_LO     12
+#define ACTCENGS_HI     15
+#define MEVBS_P16X16_FLAG  0
+#define MEVBS_P16X8_FLAG   1
+#define MEVBS_P8X16_FLAG   2
+#define MEVBS_P8X8_FLAG    3
+#define MEPXLAC_LO         4
+#define MEPXLAC_HI         5
+#define XTRACOST_FLAG      6
+#define XTRACOSTP8x8_FLAG  7
+#define INTRA_IN_P_FLAG    8
+#define I16x16_FLAG        9
+#define I4x4_FLAG         10
+#define I16x16_MASK_LO    12
+#define I16x16_MASK_HI    15
+#define ICHR_MASK_LO      16
+#define ICHR_MASK_HI      19
+#define I4x4_MASK_LO      20
+#define I4x4_MASK_HI      28
+#define SRX_LO             0
+#define SRX_HI             7
+#define SRY_LO             8
+#define SRY_HI            15
+#define VIMPW_LO          16
+#define VIMPW_HI          19
+#define VIMPR_LO          20
+#define VIMPR_HI          23
+#define VOMPW_LO          24
+#define VOMPW_HI          27
+#define REFWR_LO          28
+#define REFWR_HI          31
+#define QPADJSTRENGTH_LO   0
+#define QPADJSTRENGTH_HI   3
+#define USEINITQP_FLAG     4
+#define USESCENECUT_FLAG   5
+#define USERCOVERDUE_FLAG  6
+#define STRICTCBR_FLAG     7
+#define DISMADTH_FLAG      8
+#define VBRF_LO            9
+#define VBRF_HI            15
+#define ICPBRDFR_LO        16
+#define ICPBRDFR_HI        23
+#define SWCLR_FLAG     0
+#define SWSUSPEND_FLAG 1
+#define XTMEMPTR_LO    0
+#define XTMEMPTR_HI   31
+#define PID_KU_IONPY_LO              0
+#define PID_KU_IONPY_HI              3
+#define PID_KU_IP_LO                 4
+#define PID_KU_IP_HI                 7
+#define PID_KI_UPD_IONPY_LO          8
+#define PID_KI_UPD_IONPY_HI         15
+#define PID_KI_UPD_IP_LO            16
+#define PID_KI_UPD_IP_HI            23
+#define PID_FIFO_STATUS_TARGET_LO   24
+#define PID_FIFO_STATUS_TARGET_HI   29
+#define PID_USE_COLOCATED_FLAG      30
+#define PID_USE_IBU_THR_FLAG        31
+#define PID_IFRAME_BUFFER_FACTOR_LO  0
+#define PID_IFRAME_BUFFER_FACTOR_HI  7
+#define PID_QP_STEP_FR_LO            8
+#define PID_QP_STEP_FR_HI           11
+#define PID_QP_STEP_IBU_LO          12
+#define PID_QP_STEP_IBU_HI          15
+#define PID_FSTATUS_0_LO            16
+#define PID_FSTATUS_0_HI            22
+#define PID_THRFRMLIM_0_FLAG        23
+#define PID_FSTATUS_1_LO            24
+#define PID_FSTATUS_1_HI            30
+#define PID_THRFRMLIM_1_FLAG        31
+#define PID_FSTATUS_2_LO             0
+#define PID_FSTATUS_2_HI             6
+#define PID_THRFRMLIM_2_FLAG         7
+#define PID_FSTATUS_3_LO             8
+#define PID_FSTATUS_3_HI            14
+#define PID_THRFRMLIM_3_FLAG        15
+#define PID_FSTATUS_4_LO            16
+#define PID_FSTATUS_4_HI            22
+#define PID_THRFRMLIM_4_FLAG        23
+#define PID_QP_STEP_0_LO            24
+#define PID_QP_STEP_0_HI            27
+#define PID_QP_THR_0_LO             28
+#define PID_QP_THR_0_HI             31
+#define PID_QP_STEP_1_LO             0
+#define PID_QP_STEP_1_HI             3
+#define PID_QP_THR_1_LO              4
+#define PID_QP_THR_1_HI              7
+#define PID_QP_STEP_2_LO             8
+#define PID_QP_STEP_2_HI            11
+#define PID_QP_THR_2_LO             12
+#define PID_QP_THR_2_HI             15
+#define PID_QP_STEP_3_LO            16
+#define PID_QP_STEP_3_HI            19
+#define PID_QP_THR_3_LO             20
+#define PID_QP_THR_3_HI             23
+#define PID_QP_STEP_4_LO            24
+#define PID_QP_STEP_4_HI            27
+#define PID_QP_THR_4_LO             28
+#define PID_QP_THR_4_HI             31
+#define PID_RAMP_METHOD_LO           0
+#define PID_RAMP_METHOD_HI           1
+#define PID_USE_KD_FLAG              2
+#define PID_USE_KI_UPD_FLAG          3
+#define PID_USE_KI_UPDM_FLAG         4
+#define PID_USE_SCENE_THR_FLAG       5
+#define PID_USE_AVG_QP_THR_FLAG      6
+#define PID_AVG_QP_THR_LO           16
+#define PID_AVG_QP_THR_HI           23
+#define PID_SCENE_THR_LO            24
+#define PID_SCENE_THR_HI            31
+#define PID_USE_COLOCATED          0
+#define PID_RAMP_METHOD            1
+#define PID_KU_IONPY               8
+#define PID_KU_IP                  4
+#define PID_KI_UPD_IONPY          16
+#define PID_KI_UPD_IP              8
+#define PID_USE_KP                 1
+#define PID_USE_KI                 1
+#define PID_USE_KD                 1
+#define PID_USE_KI_UPD             1
+#define PID_USE_KI_UPDM            1
+#define PID_USE_IBU_THR            1
+#define PID_USE_AVG_QP_THR         1
+#define PID_USE_SCENE_THR          1
+#define PID_AVG_QP_THR            16
+#define PID_SCENE_THR             64
+#define PID_FIFO_STATUS_TARGET    25
+#define PID_IFRAME_BUFFER_FACTOR  65
+#define PID_QP_STEP_FR             3
+#define PID_QP_STEP_IBU            2
+#define PID_FSTATUS_4              0
+#define PID_QP_STEP_4              4
+#define PID_QP_THR_4              12
+#define PID_THRFRMLIM_4            0
+#define PID_FSTATUS_3             50
+#define PID_QP_STEP_3              3
+#define PID_QP_THR_3               9
+#define PID_THRFRMLIM_3            0
+#define PID_FSTATUS_2             37
+#define PID_QP_STEP_2              2
+#define PID_QP_THR_2               6
+#define PID_THRFRMLIM_2            1
+#define PID_FSTATUS_1             12
+#define PID_QP_STEP_1              1
+#define PID_QP_THR_1               3
+#define PID_THRFRMLIM_1            1
+#define PID_FSTATUS_0              5
+#define PID_QP_STEP_0              0
+#define PID_QP_THR_0               0
+#define PID_THRFRMLIM_0            1
+#define VBS_DIFF_8x16_16x8_PREDICTOR 0
+#define VBS_DIFF_8x8_PREDICTOR       1
+#define VBS_USE_XTRA_PART_COST       2
+#define VBS_USE_XTRA_PART_COST_P8x8  3
+#define PSKIP_USE_HIGH_COST          4
+#define VIMBPLANAR     0
+#define VIPLANAR       1
+#define VIINTERLEAVED  2
+#define VIINTERLEAVEDF 3
+#define TBA_DISABLE 0
+#define TBA_INIT    1
+#define TBA_APPEND  2
+#define TBA_FINISH  3
+#define TBA_SINGLE  4
+#define BASELINE  66
+#define MAIN      77
+#define HIGH     100
+#define HIGH10P  110
+#define HIGH422  122
+#define HIGH444  244
+#define YUV400 0
+#define YUV420 1
+#define YUV422 2
+#define YUV444 3
+#define RC_MASK_SMARTINIT_QP     0
+#define RC_MASK_SCENE_CUT_DETECT 1
+#define RC_MASK_DISOVERDUE       2
+#define RC_MASK_STRICTCBR        3
+#define RC_MASK_DISMATH          4
+#define RCM_PID 0
+#define RCM_M1  1
+#define RCM_M0  2
+#define RCM_M2  3
+
+static const int SUB_WIDTH_C_PVT[4]  = {0, 2, 2, 1};
+static const int SUB_HEIGHT_C_PVT[4] = {0, 2, 1, 1};
+
+
+#endif // _H264E_CSREGS_PRIVATE_H_

new file libavcodec/xlnx_h264/h264e_sdaccel.c

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <time.h>
+#include <CL/opencl.h>
+#include <CL/cl_ext.h>
+#include "libavutil/atomic.h"
+#include "h264e_sdaccel.h"
+#include "h264e_sdaccel_private.h"
+#include "h264e_sdaccel_utils.h"
+#include "h264e_csregs.h"
+
+static int kernel_number = -1;
+static const int kernel_ddr_map[] = {
+        XCL_MEM_DDR_BANK0,
+        XCL_MEM_DDR_BANK0,
+        XCL_MEM_DDR_BANK1,
+        XCL_MEM_DDR_BANK1,
+        XCL_MEM_DDR_BANK2,
+        XCL_MEM_DDR_BANK2,
+        XCL_MEM_DDR_BANK3,
+        XCL_MEM_DDR_BANK3
+        };
+
+static const int kernel_map[] = {0,2,4,6,1,3,5,7};
+
+int h264eInitSW(H264E_KCFG* ph264eKCFG);
+int h264eEndSW (H264E_KCFG* ph264eKCFG);
+int h264eInitHW(H264E_KCFG* ph264eKCFG);
+int h264eEndHW (H264E_KCFG* ph264eKCFG);
+int h264eClear (H264E_KCFG* ph264eKCFG);
+
+int h264eWriteDeviceBuffer(H264E_CL* ph264eCL,
+                           char*     pHostBuffer,
+                           cl_mem    DeviceBuffer,
+                           int       HostOffset,
+                           int       DeviceOffset,
+                           int       ByteSize,
+                           bool      bSync
+                           );
+
+int h264eReadDeviceBuffer(H264E_CL* ph264eCL,
+                          char*     pHostBuffer,
+                          cl_mem    DeviceBuffer,
+                          int       HostOffset,
+                          int       DeviceOffset,
+                          int       ByteSize
+                          );
+
+int h264eSetKernelArguments(H264E_KCFG* ph264eKCFG,
+                            int         KernelRunType
+                            );
+
+int h264eExecuteKernel(H264E_CL* ph264eCL, cl_event *event);
+
+
+int h264eWriteYUV(H264E_KERNEL_HANDLE h264eKernelHandle,
+                  char* pHostYBuffer,
+                  char* pHostUBuffer,
+                  char* pHostVBuffer);
+
+void *h264eEncoderThread(void *h264eKernelHandle);
+
+void h264eUpdateStats(H264E_TIME* ph264eTime);
+void h264eReportDebugStatus(int DebugStatus);
+
+H264E_KERNEL_HANDLE h264eInit(const H264E_USR* pch264eUsr)
+{
+    H264E_KCFG* ph264eKCFG;
+    int         ans;
+    if(pch264eUsr == NULL) return NULL;
+    ph264eKCFG = malloc(sizeof(H264E_KCFG));
+    if(ph264eKCFG == NULL) return NULL;
+    H264E_CL*   ph264eCL   = &ph264eKCFG->h264eCL;
+    H264E_USR*  ph264eUsr  = &ph264eKCFG->h264eUsr;
+    H264E_CNTX* ph264eCntx = &ph264eKCFG->h264eCntx;
+    memcpy(ph264eUsr, pch264eUsr, sizeof(H264E_USR));
+    ans  = h264eInitSW(ph264eKCFG);
+    ans |= h264eInitHW(ph264eKCFG);
+    if(ans < 0)
+    {
+        h264eEndSW(ph264eKCFG);
+        h264eEndHW(ph264eKCFG);
+
+        free(ph264eKCFG);
+        return NULL;
+    }
+    memset(ph264eCntx->CRegsFile, 0, sizeof(ph264eCntx->CRegsFile));
+    memset(ph264eCntx->SRegsFile, 0, sizeof(ph264eCntx->SRegsFile));
+    h264eSetControlRegs(ph264eUsr, ph264eCntx->CRegsFile);
+    ans = h264eWriteDeviceBuffer(ph264eCL,
+                                 ph264eCntx->CRegsFile,
+                                 ph264eCntx->DeviceBuffer,
+                                 0,
+                                 ph264eCntx->CRegsPtr,
+                                 sizeof(ph264eCntx->CRegsFile),
+                                 true
+                                 );
+    ans |= h264eSetKernelArguments(ph264eKCFG, CONTROL_RT);
+    ans |= h264eExecuteKernel(ph264eCL, NULL);
+    ans |= h264eSetKernelArguments(ph264eKCFG, STATUS_RT);
+    ans |= h264eExecuteKernel(ph264eCL, NULL);
+    if(ans < 0)
+    {
+        h264eEndSW(ph264eKCFG);
+        h264eEndHW(ph264eKCFG);
+        free(ph264eKCFG);
+        return NULL;
+    }
+    ans = h264eReadDeviceBuffer(ph264eCL,
+                                ph264eKCFG->h264eCntx.SRegsFile,
+                                ph264eKCFG->h264eCntx.DeviceBuffer,
+                                0,
+                                ph264eKCFG->h264eCntx.SRegsPtr,
+                                sizeof(ph264eKCFG->h264eCntx.SRegsFile)
+                                );
+    if(ans < 0)
+    {
+        h264eEndSW(ph264eKCFG);
+        h264eEndHW(ph264eKCFG);
+
+        free(ph264eKCFG);
+        return NULL;
+    }
+    int ErrID = h264eGetErrorCode((int*)ph264eKCFG->h264eCntx.SRegsFile);
+    if(ErrID)
+    {
+        H264EMSG(printf("ERROR: H.264 core programming error\n"));
+        h264eClear(ph264eKCFG);
+        h264eEndSW(ph264eKCFG);
+        h264eEndHW(ph264eKCFG);
+        free(ph264eKCFG);
+        return NULL;
+    }
+    else
+    {
+        return (void*)ph264eKCFG;
+    }
+}
+
+int h264eEnd(H264E_KERNEL_HANDLE h264eKernelHandle)
+{
+    int ans;
+    if(h264eKernelHandle == NULL)
+    {
+        H264EMSG(printf("ERROR: Invalid kernel handle\n"));
+        return H264E_BAD_KERNEL_HANDLE;
+    }
+    H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+    ans  = h264eEndSW(ph264eKCFG);
+    ans |= h264eEndHW(ph264eKCFG);
+    if(ans < 0)
+    {
+        return H264E_END_ERR;
+    }
+    else
+    {
+        return H264E_OK;
+    }
+}
+
+int h264eWriteYUV(H264E_KERNEL_HANDLE h264eKernelHandle,
+                  char* pHostYBuffer,
+                  char* pHostUBuffer,
+                  char* pHostVBuffer)
+{
+    H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+    int  YSizeInBytes     = 0;
+    int  ChrSizeInBytes   = 0;
+    int  FrameSizeInBytes = 0;
+    int  VDataIdx         = 0;
+    H264E_CL* ph264eCL = &ph264eKCFG->h264eCL;
+    H264E_CNTX *ph264eCNTX = &ph264eKCFG->h264eCntx;
+    VDataIdx = ph264eCNTX->VDataIdx;
+    H264ETIME(timespec tss);
+    H264ETIME(timespec tse);
+    H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tss));
+    if(ph264eKCFG->h264eUsr.ScanFormat == H264E_PLANAR_SEPARATED)
+    {
+        if(pHostYBuffer == NULL || pHostUBuffer == NULL || pHostVBuffer == NULL)
+        {
+            H264EMSG(printf("ERROR: Invalid Y, U or V data pointer(s)\n"));
+            return H264E_BAD_COMP_PTR;
+        }
+        YSizeInBytes     =  ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY;
+        ChrSizeInBytes   = (ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY) / 4;
+        FrameSizeInBytes = YSizeInBytes + 2*ChrSizeInBytes;
+        if(ph264eKCFG->h264eUsr.PackPlanarSeparated)
+        {
+            memcpy(ph264eKCFG->h264eCntx.pPackedPlanarHostBuffer,                                 pHostYBuffer, YSizeInBytes);
+            memcpy(ph264eKCFG->h264eCntx.pPackedPlanarHostBuffer + YSizeInBytes,                  pHostUBuffer, ChrSizeInBytes);
+            memcpy(ph264eKCFG->h264eCntx.pPackedPlanarHostBuffer + YSizeInBytes + ChrSizeInBytes, pHostVBuffer, ChrSizeInBytes);
+            h264eWriteDeviceBuffer(ph264eCL,
+                                   ph264eKCFG->h264eCntx.pPackedPlanarHostBuffer,
+                                   ph264eKCFG->h264eCntx.DeviceBuffer,
+                                   0,
+                                   ph264eKCFG->h264eCntx.VDataPtr[VDataIdx],
+                                   FrameSizeInBytes,
+                                   true
+                                   );
+            if(ph264eCL->bWRDeviceBufferFailed) return H264E_DEVBUFF_WR_ERR;
+        }
+        else
+        {
+            h264eWriteDeviceBuffer(ph264eCL,
+                                   pHostYBuffer,
+                                   ph264eKCFG->h264eCntx.DeviceBuffer,
+                                   0,
+                                   ph264eKCFG->h264eCntx.VDataPtr[VDataIdx],
+                                   YSizeInBytes,
+                                   true
+                                   );
+            if(ph264eCL->bWRDeviceBufferFailed) return H264E_DEVBUFF_WR_ERR;
+            h264eWriteDeviceBuffer(ph264eCL,
+                                   pHostUBuffer,
+                                   ph264eKCFG->h264eCntx.DeviceBuffer,
+                                   0,
+                                   ph264eKCFG->h264eCntx.VDataPtr[VDataIdx] + YSizeInBytes,
+                                   ChrSizeInBytes,
+                                   true
+                                   );
+            if(ph264eCL->bWRDeviceBufferFailed) return H264E_DEVBUFF_WR_ERR;
+            h264eWriteDeviceBuffer(ph264eCL,
+                                   pHostVBuffer,
+                                   ph264eKCFG->h264eCntx.DeviceBuffer,
+                                   0,
+                                   ph264eKCFG->h264eCntx.VDataPtr[VDataIdx] + YSizeInBytes + ChrSizeInBytes,
+                                   ChrSizeInBytes,
+                                   true
+                                   );
+            if(ph264eCL->bWRDeviceBufferFailed) return H264E_DEVBUFF_WR_ERR;
+        }
+
+        if(ph264eKCFG->h264eUsr.HWOptimizeScanFormat)
+        {
+              ph264eKCFG->h264eCntx.Format = PLANAR_TO_INTERLEAVED_FRMT;
+        }
+        else
+        {
+            ph264eKCFG->h264eCntx.Format = NO_CHANGE_FRMT;
+        }
+    }
+    else
+    {
+        if(pHostYBuffer == NULL)
+        {
+            H264EMSG(printf("ERROR: Invalid Y data pointer\n"));
+            return H264E_BAD_COMP_PTR;
+        }
+        if(ph264eKCFG->h264eUsr.ScanFormat == H264E_PLANAR_PACKED)
+        {
+            YSizeInBytes     =  ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY;
+            ChrSizeInBytes   = (ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY) / 4;
+            FrameSizeInBytes = YSizeInBytes + 2*ChrSizeInBytes;
+            if(ph264eKCFG->h264eUsr.HWOptimizeScanFormat)
+            {
+                ph264eKCFG->h264eCntx.Format = PLANAR_TO_INTERLEAVED_FRMT;
+            }
+            else
+            {
+                ph264eKCFG->h264eCntx.Format = NO_CHANGE_FRMT;
+            }
+        }
+        else if(ph264eKCFG->h264eUsr.ScanFormat == H264E_INTERLEAVED_YC)
+        {
+            YSizeInBytes     =  ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY;
+            ChrSizeInBytes   = (ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY) / 4;
+            FrameSizeInBytes = YSizeInBytes + 2*ChrSizeInBytes;
+            ph264eKCFG->h264eCntx.Format = NO_CHANGE_FRMT;
+        }
+        else if(ph264eKCFG->h264eUsr.ScanFormat == H264E_INTERLEAVED_CY)
+        {
+            YSizeInBytes     =  ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY;
+            ChrSizeInBytes   = (ph264eKCFG->h264eUsr.FrameX * ph264eKCFG->h264eUsr.FrameY) / 4;
+            FrameSizeInBytes = YSizeInBytes + 2*ChrSizeInBytes;
+            ph264eKCFG->h264eCntx.Format = NO_CHANGE_FRMT;
+        }
+        else if(ph264eKCFG->h264eUsr.ScanFormat == H264E_MACROBLOCK)
+        {
+            int FrameX = ph264eKCFG->h264eUsr.FrameX;
+            int FrameY = ph264eKCFG->h264eUsr.FrameY;
+            FrameX = ((FrameX + 15) / 16) * 16;
+            FrameY = ((FrameY + 15) / 16) * 16;
+            YSizeInBytes     =  FrameX * FrameY;
+            ChrSizeInBytes   = (FrameX * FrameY) / 4;
+            FrameSizeInBytes = YSizeInBytes + 2*ChrSizeInBytes;
+            ph264eKCFG->h264eCntx.Format = NO_CHANGE_FRMT;
+        }
+        h264eWriteDeviceBuffer(ph264eCL,
+                               pHostYBuffer,
+                               ph264eKCFG->h264eCntx.DeviceBuffer,
+                               0,
+                               ph264eKCFG->h264eCntx.VDataPtr[VDataIdx],
+                               FrameSizeInBytes,
+                               true
+                               );
+        if(ph264eCL->bWRDeviceBufferFailed) return H264E_DEVBUFF_WR_ERR;
+    }
+    H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tse));
+    H264ETIME(ph264eKCFG->h264eTime.WriteDeviceBufferTime = tse.tv_nsec - tss.tv_nsec);
+    return H264E_OK;
+}
+
+int h264eEncodeFrame(H264E_KERNEL_HANDLE h264eKernelHandle,
+                     char* pHostYBuffer,
+                     char* pHostUBuffer,
+                     char* pHostVBuffer,
+                     char* pHostNALBuffer,
+                     int   frame_count
+                     )
+{
+    int  rc;
+    int  NalSizeInBytes   = 0;
+    if(h264eKernelHandle == NULL)
+    {
+        H264EMSG(printf("ERROR: Invalid kernel handle\n"));
+        return H264E_BAD_KERNEL_HANDLE;
+    }
+    if(pHostNALBuffer == NULL)
+    {
+        H264EMSG(printf("ERROR: Invalid NAL pointer\n"));
+        return H264E_BAD_NAL_PTR;
+    }
+    H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+    H264E_CL* ph264eCL = &ph264eKCFG->h264eCL;
+    H264ETIME(timespec tss);
+    H264ETIME(timespec tse);
+    H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tss));
+    rc = h264eWriteYUV(h264eKernelHandle,
+                       pHostYBuffer,
+                       pHostUBuffer,
+                       pHostVBuffer);
+    if (rc != H264E_OK) return rc;
+    if (frame_count == 0)
+    {
+        H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tss));
+        h264eSetKernelArguments(ph264eKCFG, ENCODE_RT);
+        if(ph264eCL->bSetKernelArgsFailed) return H264E_ARGSET_ERR;
+        h264eExecuteKernel(ph264eCL, &ph264eCL->kernel_event);
+        if(ph264eCL->bExecuteKernelFailed) return H264E_EXEC_ERR;
+        return H264E_DEFER_DATA;
+    }
+    else
+    {
+        clWaitForEvents(1, &ph264eCL->kernel_event);
+        H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tss));
+        h264eReadDeviceBuffer(ph264eCL,
+                            ph264eKCFG->h264eCntx.SRegsFile,
+                            ph264eKCFG->h264eCntx.DeviceBuffer,
+                            0,
+                            ph264eKCFG->h264eCntx.SRegsPtr,
+                            sizeof(ph264eKCFG->h264eCntx.SRegsFile)
+                            );
+        if(ph264eCL->bRDDeviceBufferFailed) return H264E_DEVBUFF_RD_ERR;
+        H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tse));
+        H264ETIME(ph264eKCFG->h264eTime.EncodeAndReadNALSizeTime = tse.tv_nsec - tss.tv_nsec);
+        NalSizeInBytes = h264eGetNalSize((int*)ph264eKCFG->h264eCntx.SRegsFile);
+        if(NalSizeInBytes == 0)
+        {
+            H264EMSG(printf("ERROR: H.264 core returned 0 NAL size\n"));
+            return H264E_CORE_ERR;
+        }
+        H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tss));
+        h264eReadDeviceBuffer(ph264eCL,
+                              pHostNALBuffer,
+                              ph264eKCFG->h264eCntx.DeviceBuffer,
+                              0,
+                              ph264eKCFG->h264eCntx.NALPtr,
+                              NalSizeInBytes
+                            );
+        if(ph264eCL->bRDDeviceBufferFailed) return H264E_DEVBUFF_RD_ERR;
+        H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tse));
+        H264ETIME(ph264eKCFG->h264eTime.ReadNALBufferTime = tse.tv_nsec - tss.tv_nsec);
+        H264ETIME(h264eUpdateStats(&ph264eKCFG->h264eTime));
+        H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tss));
+        h264eSetKernelArguments(ph264eKCFG, ENCODE_RT);
+        if(ph264eCL->bSetKernelArgsFailed) return H264E_ARGSET_ERR;
+        h264eExecuteKernel(ph264eCL, &ph264eCL->kernel_event);
+        if(ph264eCL->bExecuteKernelFailed) return H264E_EXEC_ERR;
+        H264ETIME(clock_gettime(CLOCK_MONOTONIC_RAW, &tse));
+        H264ETIME(ph264eKCFG->h264eTime.ExecuteKernelTime = tse.tv_nsec - tss.tv_nsec);
+        rc = NalSizeInBytes;
+    }
+    return rc;
+}
+
+int h264eInitSW(H264E_KCFG* ph264eKCFG)
+{
+    int rc = 0;
+    int IPCoreMemSizeInBytes;
+    int CRegsSizeInBytes;
+    int SRegsSizeInBytes;
+    int VDataSizeInBytes;
+    int NALSizeInBytes;
+    H264E_USR*  ph264eUsr  = &ph264eKCFG->h264eUsr;
+    H264E_CNTX* ph264eCntx = &ph264eKCFG->h264eCntx;
+    CRegsSizeInBytes = 64 * 4;
+    SRegsSizeInBytes = 64 * 4;
+    VDataSizeInBytes     = ph264eUsr->FrameX * ph264eUsr->FrameY * 1.5;
+    IPCoreMemSizeInBytes = VDataSizeInBytes * 4;
+    NALSizeInBytes       = VDataSizeInBytes;
+    IPCoreMemSizeInBytes *= 1.1;
+    VDataSizeInBytes     *= 1.1;
+    NALSizeInBytes       *= 1.1;
+    IPCoreMemSizeInBytes = ((IPCoreMemSizeInBytes / 4096) + 1) * 4096;
+    CRegsSizeInBytes     = ((CRegsSizeInBytes     / 4096) + 1) * 4096;
+    SRegsSizeInBytes     = ((SRegsSizeInBytes     / 4096) + 1) * 4096;
+    VDataSizeInBytes     = ((VDataSizeInBytes     / 4096) + 1) * 4096;
+    NALSizeInBytes       = ((NALSizeInBytes       / 4096) + 1) * 4096;
+    ph264eCntx->NALPtr   = IPCoreMemSizeInBytes;
+    ph264eCntx->SRegsPtr = ph264eCntx->NALPtr   + NALSizeInBytes;
+    ph264eCntx->CRegsPtr = ph264eCntx->SRegsPtr + SRegsSizeInBytes;
+    ph264eCntx->VDataPtr[0] = ph264eCntx->CRegsPtr + CRegsSizeInBytes;
+    ph264eCntx->VDataPtr[1] = ph264eCntx->VDataPtr[0] + VDataSizeInBytes;
+    ph264eCntx->VDataIdx                = 0;
+    ph264eCntx->Format                  = 0;
+    ph264eCntx->DeviceBuffer            = 0;
+    ph264eCntx->DeviceBufferSizeInBytes = ph264eCntx->VDataPtr[1] + VDataSizeInBytes;
+    memset(&ph264eKCFG->h264eTime, 0, sizeof(H264E_TIME));
+    ph264eCntx->pPackedPlanarHostBuffer = NULL;
+    if(ph264eUsr->ScanFormat == H264E_PLANAR_SEPARATED && ph264eUsr->PackPlanarSeparated)
+    {
+        ph264eCntx->pPackedPlanarHostBuffer = malloc(ph264eUsr->FrameX *
+                                                                                                         ph264eUsr->FrameY *
+                                                                                                         2);
+        if(ph264eCntx->pPackedPlanarHostBuffer == NULL)
+        {
+            H264EMSG(printf("ERROR: Failed to allocate host buffer for packing planar separated input\n"));
+            return H264E_INITSW_ERR;
+        }
+    }
+    if(ph264eUsr->PackPlanarSeparated  < 0 || ph264eUsr->PackPlanarSeparated  > 0) ph264eUsr->PackPlanarSeparated  = 1;
+    if(ph264eUsr->HWOptimizeScanFormat < 0 || ph264eUsr->HWOptimizeScanFormat > 0) ph264eUsr->HWOptimizeScanFormat = 1;
+        printf("Creating request message queue\n");
+        rc = av_thread_message_queue_alloc(&ph264eKCFG->h264eReqMsgQ, 1, sizeof(H264E_REQ_MSG));
+        printf("Creating response message queue\n");
+        rc |= av_thread_message_queue_alloc(&ph264eKCFG->h264eRspMsgQ,
+                                                                           5, sizeof(H264E_RSP_MSG));
+        if (rc < 0)
+        {
+                printf("Failed to allocate message queue\n");
+                exit(1);
+        }
+        rc = pthread_create(&ph264eKCFG->h264eThread, NULL,
+                                                h264eEncoderThread, ph264eKCFG);
+    return H264E_OK;
+}
+
+int h264eEndSW(H264E_KCFG* ph264eKCFG)
+{
+        H264E_REQ_MSG reqMsg;
+        reqMsg.endThread = 1;
+        av_thread_message_queue_send(ph264eKCFG->h264eReqMsgQ, &reqMsg, 0);
+        pthread_join(ph264eKCFG->h264eThread, NULL);
+        av_thread_message_queue_free(&ph264eKCFG->h264eReqMsgQ);
+        av_thread_message_queue_free(&ph264eKCFG->h264eRspMsgQ);
+    if(ph264eKCFG->h264eCntx.pPackedPlanarHostBuffer != NULL)
+    {
+        free(ph264eKCFG->h264eCntx.pPackedPlanarHostBuffer);
+        ph264eKCFG->h264eCntx.pPackedPlanarHostBuffer = NULL;
+    }
+    H264EMSG(printf("\n"));
+    H264EMSG(printf("\t\t======================================================\n"));
+    H264EMSG(printf("\t\tTotal Write Device Buffer Run-Time    : %8.2f sec\n",  ph264eKCFG->h264eTime.WriteDeviceBufferTotalTimeMS    / 1000.0));
+    H264EMSG(printf("\t\tTotal Execute Kernel Run-Time         : %8.2f sec\n",  ph264eKCFG->h264eTime.ExecuteKernelTotalTimeMS        / 1000.0));
+    H264EMSG(printf("\t\tTotal Encode & Read NAL Size Run-Time : %8.2f sec\n",  ph264eKCFG->h264eTime.EncodeAndReadNALSizeTotalTimeMS / 1000.0));
+    H264EMSG(printf("\t\tTotal Read NAL Buffer Run-Time        : %8.2f sec\n",  ph264eKCFG->h264eTime.ReadNALBufferTotalTimeMS        / 1000.0));
+    H264ETIME(
+        double TotalCLTime;
+        TotalCLTime = ph264eKCFG->h264eTime.WriteDeviceBufferTotalTimeMS    +
+                      ph264eKCFG->h264eTime.ExecuteKernelTotalTimeMS        +
+                      ph264eKCFG->h264eTime.EncodeAndReadNALSizeTotalTimeMS +
+                      ph264eKCFG->h264eTime.ReadNALBufferTotalTimeMS;
+        TotalCLTime /= 1000.0;
+    );
+    H264EMSG(printf("\t\t======================================================\n"));
+    H264EMSG(printf("\t\tTotal CL Run-Time                     : %8.2f sec\n", TotalCLTime));
+    H264EMSG(printf("\t\t======================================================\n"));
+    H264EMSG(printf("\n"));
+    return H264E_OK;
+}
+
+int h264eInitHW(H264E_KCFG* ph264eKCFG)
+{
+   int k_num;
+   if(!ph264eKCFG->h264eUsr.bCLIsOK)
+   {
+           k_num = avpriv_atomic_int_add_and_fetch(&kernel_number, 1);
+           ph264eKCFG->h264eCL.kernel_num = kernel_map[k_num];
+       ph264eKCFG->h264eCL.pDeviceName        = ph264eKCFG->h264eUsr.pDeviceName;
+       ph264eKCFG->h264eCL.pKernelBinFileName = ph264eKCFG->h264eUsr.pKernelBinFileName;
+           snprintf(ph264eKCFG->h264eCL.pKernelName, 64, "%s_%d", 
+                ph264eKCFG->h264eUsr.pKernelName, 
+                ph264eKCFG->h264eCL.kernel_num);
+       h264e_cl_init(&ph264eKCFG->h264eCL);
+       if(!ph264eKCFG->h264eCL.bCLIsOK)
+       {
+           H264EMSG(printf("ERROR: Failed to initialize CL\n"));
+           return H264E_INITHW_ERR;
+       }
+   }
+   else
+   {
+       ph264eKCFG->h264eCL.commands = ph264eKCFG->h264eUsr.commands;
+       ph264eKCFG->h264eCL.context  = ph264eKCFG->h264eUsr.context;
+       ph264eKCFG->h264eCL.kernel   = ph264eKCFG->h264eUsr.kernel;
+   }
+        cl_mem_ext_ptr_t cl_mem_ptr;
+        cl_mem_ptr.obj = 0;
+        cl_mem_ptr.param = 0;
+        cl_mem_ptr.flags = kernel_ddr_map[ph264eKCFG->h264eCL.kernel_num];
+   ph264eKCFG->h264eCntx.DeviceBuffer = clCreateBuffer(ph264eKCFG->h264eCL.context,
+                                                       CL_MEM_READ_WRITE | CL_MEM_EXT_PTR_XILINX,
+                                                       ph264eKCFG->h264eCntx.DeviceBufferSizeInBytes,
+                                                       &cl_mem_ptr,
+                                                       NULL);
+   if(!ph264eKCFG->h264eCntx.DeviceBuffer)
+   {
+       H264EMSG(printf("ERROR: Failed to allocate device memory\n"));
+        return H264E_DEVBUFF_ERR;
+   }
+    return H264E_OK;
+}
+
+int h264eEndHW(H264E_KCFG* ph264eKCFG)
+{
+    int err;
+    int ans;
+    int retval = H264E_OK;
+   if(ph264eKCFG->h264eCntx.DeviceBuffer)
+   {
+        err = clReleaseMemObject(ph264eKCFG->h264eCntx.DeviceBuffer);
+        if(err != CL_SUCCESS)
+        {
+            H264EMSG(printf("ERROR: Failed to release device buffer\n"));
+            retval = H264E_DEVBUFF_ERR;
+   }
+   }
+   if(ph264eKCFG->h264eCL.bCLIsOK && !ph264eKCFG->h264eUsr.bCLIsOK)
+   {
+        ans = h264e_cl_end(&(ph264eKCFG->h264eCL));
+        if(ans < 0) retval = ans;
+   }
+   return retval;
+}
+
+int h264eWriteDeviceBuffer(H264E_CL* ph264eCL,
+                           char*     pHostBuffer,
+                           cl_mem    DeviceBuffer,
+                           int       HostOffset,
+                           int       DeviceOffset,
+                           int       ByteSize,
+                           bool      bSync
+                           )
+{
+   int      err;
+   cl_event writeevent;
+   if(bSync)
+   {
+       err = clEnqueueWriteBuffer(ph264eCL->commands,
+                                  DeviceBuffer,
+                                  CL_TRUE,
+                                  DeviceOffset,
+                                  ByteSize,
+                                  pHostBuffer + HostOffset,
+                                  0,
+                                  NULL,
+                                  &writeevent
+                                  );
+   }
+   else
+   {
+       err = clEnqueueWriteBuffer(ph264eCL->commands,
+                                  DeviceBuffer,
+                                  CL_TRUE,
+                                  DeviceOffset,
+                                  ByteSize,
+                                  pHostBuffer + HostOffset,
+                                  0,
+                                  NULL,
+                                  NULL
+                                  );
+   }
+   if(err != CL_SUCCESS)
+   {
+      H264EMSG(printf("ERROR: Failed to write to device buffer\n"));
+        ph264eCL->bWRDeviceBufferFailed = true;
+        return H264E_DEVBUFF_WR_ERR;
+   }
+   if(bSync)
+   {
+      clWaitForEvents(1, &writeevent);
+   }
+    ph264eCL->bWRDeviceBufferFailed = false;
+    return H264E_OK;
+}
+
+int h264eReadDeviceBuffer(H264E_CL* ph264eCL,
+                          char*     pHostBuffer,
+                          cl_mem    DeviceBuffer,
+                          int       HostOffset,
+                          int       DeviceOffset,
+                          int       ByteSize)
+{
+    int err;
+    cl_event readevent;
+    err = clEnqueueReadBuffer(ph264eCL->commands,
+                              DeviceBuffer,
+                              CL_TRUE,
+                              DeviceOffset,
+                              ByteSize,
+                              pHostBuffer + HostOffset,
+                              0,
+                              NULL,
+                              &readevent
+                              );
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to read device buffer (error %d)\n", err));
+        ph264eCL->bRDDeviceBufferFailed = true;
+        return H264E_DEVBUFF_RD_ERR;
+    }
+    clWaitForEvents(1, &readevent);
+    ph264eCL->bRDDeviceBufferFailed = false;
+    return H264E_OK;
+}
+
+int h264eSetKernelArguments(H264E_KCFG* ph264eKCFG,
+                            int         KernelRunType
+                            )
+{
+    int err;
+    H264E_CL*   ph264eCL   = &ph264eKCFG->h264eCL;
+    H264E_USR*  ph264eUsr  = &ph264eKCFG->h264eUsr;
+    H264E_CNTX* ph264eCntx = &ph264eKCFG->h264eCntx;
+    int CRegsPtr = ph264eCntx->CRegsPtr;
+    int SRegsPtr = ph264eCntx->SRegsPtr;
+    int VDataPtr = ph264eCntx->VDataPtr[ph264eCntx->VDataIdx];
+    int NALPtr   = ph264eCntx->NALPtr;
+    int Format   = ph264eCntx->Format;
+    int FrameX   = ph264eUsr->FrameX;
+    int FrameY   = ph264eUsr->FrameY;
+    if(ph264eUsr->ScanFormat == H264E_MACROBLOCK)
+    {
+        FrameX = ((FrameX + 15) / 16) * 16;
+        FrameY = ((FrameY + 15) / 16) * 16;
+    }
+    cl_kernel kernel       = ph264eCL->kernel;
+    cl_mem    DeviceBuffer = ph264eCntx->DeviceBuffer;
+    err  = 0;
+    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &DeviceBuffer);
+    err |= clSetKernelArg(kernel, 1, sizeof(int), &CRegsPtr);
+    err |= clSetKernelArg(kernel, 2, sizeof(int), &SRegsPtr);
+    err |= clSetKernelArg(kernel, 3, sizeof(int), &VDataPtr);
+    err |= clSetKernelArg(kernel, 4, sizeof(int), &NALPtr);
+    err |= clSetKernelArg(kernel, 5, sizeof(int), &FrameX);
+    err |= clSetKernelArg(kernel, 6, sizeof(int), &FrameY);
+    err |= clSetKernelArg(kernel, 7, sizeof(int), &Format);
+    err |= clSetKernelArg(kernel, 8, sizeof(int), &KernelRunType);
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to set kernel arguments (error %d)\n", err));
+        ph264eKCFG->h264eCL.bSetKernelArgsFailed = true;
+        return H264E_ARGSET_ERR;
+    }
+    ph264eKCFG->h264eCL.bSetKernelArgsFailed = false;
+    return H264E_OK;
+}
+
+int h264eExecuteKernel(H264E_CL* ph264eCL, cl_event *event)
+{
+    int err;
+    cl_event execwait;
+    cl_event *pEvent;
+    if (event)
+        pEvent = event;
+    else
+        pEvent = &execwait;
+    err = clEnqueueTask(ph264eCL->commands, ph264eCL->kernel, 0, NULL, pEvent);
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to execute kernel (error %d)\n", err));
+        ph264eCL->bExecuteKernelFailed = true;
+        return H264E_EXEC_ERR;
+    }
+    if (event == NULL)
+    {
+        clWaitForEvents(1, &execwait);
+        ph264eCL->bExecuteKernelFailed = false;
+    }
+    return H264E_OK;
+}
+
+const char* h264eReadStatusRegs(H264E_KERNEL_HANDLE h264eKernelHandle)
+{
+    if(h264eKernelHandle == NULL)
+    {
+        H264EMSG(printf("ERROR: Invalid kernel handle\n"));
+        return NULL;
+    }
+   H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+   h264eSetKernelArguments(ph264eKCFG, STATUS_RT);
+    if(ph264eKCFG->h264eCL.bSetKernelArgsFailed) return NULL;
+
+   h264eExecuteKernel(&ph264eKCFG->h264eCL, NULL);
+    if(ph264eKCFG->h264eCL.bExecuteKernelFailed) return NULL;
+   h264eReadDeviceBuffer(&ph264eKCFG->h264eCL,
+                         ph264eKCFG->h264eCntx.SRegsFile,
+                         ph264eKCFG->h264eCntx.DeviceBuffer,
+                         0,
+                         ph264eKCFG->h264eCntx.SRegsPtr,
+                         sizeof(ph264eKCFG->h264eCntx.SRegsFile)
+                         );
+    if(ph264eKCFG->h264eCL.bRDDeviceBufferFailed) return NULL;
+   return ph264eKCFG->h264eCntx.SRegsFile;
+}
+
+int h264eClear(H264E_KCFG* ph264eKCFG)
+{
+   h264eSetKernelArguments(ph264eKCFG, CLEAR_RT);
+    if(ph264eKCFG->h264eCL.bSetKernelArgsFailed) return H264E_ARGSET_ERR;
+   h264eExecuteKernel(&ph264eKCFG->h264eCL, NULL);
+    if(ph264eKCFG->h264eCL.bExecuteKernelFailed) return H264E_EXEC_ERR;
+    return H264E_OK;
+}
+
+int h264eSetUserContext(H264E_KERNEL_HANDLE h264eKernelHandle, void* Context)
+{
+    if(h264eKernelHandle == NULL)
+    {
+        H264EMSG(printf("ERROR: Invalid kernel handle\n"));
+        return H264E_BAD_KERNEL_HANDLE;
+    }
+    H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+    ph264eKCFG->h264eUsr.UserContext = Context;
+    return H264E_OK;
+}
+
+void* h264eGetUserContext(H264E_KERNEL_HANDLE h264eKernelHandle)
+{
+    if(h264eKernelHandle == NULL)
+    {
+        H264EMSG(printf("ERROR: Invalid kernel handle\n"));
+        return NULL;
+    }
+    H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+    return ph264eKCFG->h264eUsr.UserContext;
+}
+
+void h264eUpdateStats(H264E_TIME* ph264eTime)
+{
+    ph264eTime->nFrames++;
+    ph264eTime->WriteDeviceBufferTimeMS    = (double)ph264eTime->WriteDeviceBufferTime    / 1E6;
+    ph264eTime->ExecuteKernelTimeMS        = (double)ph264eTime->ExecuteKernelTime        / 1E6;
+    ph264eTime->EncodeAndReadNALSizeTimeMS = (double)ph264eTime->EncodeAndReadNALSizeTime / 1E6;
+    ph264eTime->ReadNALBufferTimeMS        = (double)ph264eTime->ReadNALBufferTime        / 1E6;
+    if(ph264eTime->WriteDeviceBufferTimeMS    < 0) ph264eTime->WriteDeviceBufferTimeMS    += 1000.0;
+    if(ph264eTime->ExecuteKernelTimeMS        < 0) ph264eTime->ExecuteKernelTimeMS        += 1000.0;
+    if(ph264eTime->EncodeAndReadNALSizeTimeMS < 0) ph264eTime->EncodeAndReadNALSizeTimeMS += 1000.0;
+    if(ph264eTime->ReadNALBufferTimeMS        < 0) ph264eTime->ReadNALBufferTimeMS        += 1000.0;
+    ph264eTime->WriteDeviceBufferTotalTimeMS    += ph264eTime->WriteDeviceBufferTimeMS;
+    ph264eTime->ExecuteKernelTotalTimeMS        += ph264eTime->ExecuteKernelTimeMS;
+    ph264eTime->EncodeAndReadNALSizeTotalTimeMS += ph264eTime->EncodeAndReadNALSizeTimeMS;
+    ph264eTime->ReadNALBufferTotalTimeMS        += ph264eTime->ReadNALBufferTimeMS;
+    return;
+}
+
+const H264E_TIME* h264eGetStats(H264E_KERNEL_HANDLE h264eKernelHandle)
+{
+    if(h264eKernelHandle == NULL)
+    {
+        H264EMSG(printf("ERROR: Invalid kernel handle\n"));
+        return NULL;
+    }
+    H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+    return &ph264eKCFG->h264eTime;
+}
+
+void h264eReportDebugStatus(int DebugStatus)
+{
+    int i = 0;
+    printf("TIMEOUT: xtmem_req_val     = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmem_req_rdy     = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmem_wdata_val   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmem_wdata_rdy   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmem_rdata_val   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmem_rdata_rdy   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: sreg_wrreq_val    = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: sreg_wrreq_rdy    = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_rdreq_val   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_rdreq_rdy   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: nal_wrreq_val     = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: nal_wrreq_rdy     = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_wrreq_val = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_wrreq_rdy = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_rdreq_val = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_rdreq_rdy = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_rdata_val   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_rdata_rdy   = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: nal_wdata_val     = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: nal_wdata_rdy     = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_wdata_val = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_wdata_rdy = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_rdata_val = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: xtmemin_rdata_rdy = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_dout_val    = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_dout_rdy    = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_val         = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: vdata_rdy         = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: nal_val           = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    printf("TIMEOUT: nal_rdy           = %d\n", DebugStatus & (0x20000000 >> i++) ? 1 : 0);
+    return;
+}
+
+void *h264eEncoderThread(void *h264eKernelHandle)
+{
+        H264E_KCFG* ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+    H264E_CNTX *ph264eCNTX = &ph264eKCFG->h264eCntx;
+    H264E_CL   *ph264eCL   = &ph264eKCFG->h264eCL;
+        AVThreadMessageQueue *reqMq = ph264eKCFG->h264eReqMsgQ;
+        AVThreadMessageQueue *rspMq = ph264eKCFG->h264eRspMsgQ;
+        H264E_REQ_MSG reqMsg;
+        H264E_RSP_MSG rspMsg;
+    int frame_count = 0;
+    int NalSizeInBytes;
+        printf("Starting thread for kernel %d\n", ph264eCL->kernel_num);
+    printf("Using latest version of Alma plugin\n");
+        while(1)
+        {
+        ph264eCNTX->VDataIdx = frame_count & 1;
+                av_thread_message_queue_recv(reqMq, &reqMsg, 0);
+                if (reqMsg.endThread)
+        {
+            clWaitForEvents(1, &ph264eCL->kernel_event);
+                        break;
+        }
+                else
+                {
+                        NalSizeInBytes = h264eEncodeFrame(ph264eKCFG,
+                                                                  reqMsg.pHostYBuffer,
+                                                                  reqMsg.pHostUBuffer,
+                                                                  reqMsg.pHostVBuffer,
+                                                                  reqMsg.pHostNALBuffer,
+                                              frame_count);
+            frame_count++;
+            if (NalSizeInBytes == H264E_DEFER_DATA)
+                continue;
+            rspMsg.NalSizeInBytes = NalSizeInBytes;
+                        rspMsg.pHostNALBuffer = reqMsg.pHostNALBuffer;
+                        rspMsg.pFrame = reqMsg.pFrame;
+                        av_thread_message_queue_send(rspMq, &rspMsg, 0);
+                }
+        }
+        pthread_exit(NULL);
+}

new file libavcodec/xlnx_h264/h264e_sdaccel.h

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#ifndef _H264E_SDACCEL_H_
+#define _H264E_SDACCEL_H_
+
+#include <stdbool.h>
+#include <CL/opencl.h>
+#include <CL/cl_ext.h>
+#include "libavcodec/avcodec.h"
+
+struct _H264E_RSP_MSG;
+typedef struct _H264E_RSP_MSG H264E_RSP_MSG;
+
+#ifdef H264E_DEBUG
+    #define H264E_PROFILING
+    #define H264EMSG(s) s
+#else
+    #define H264EMSG(s) ;
+#endif
+
+#ifdef H264E_PROFILING
+    #define H264ETIME(s) s
+#else
+    #define H264ETIME(s) ;
+#endif
+
+typedef void* H264E_KERNEL_HANDLE;
+#define H264E_PLANAR_PACKED     0
+#define H264E_PLANAR_SEPARATED  1
+#define H264E_INTERLEAVED_YC    2
+#define H264E_INTERLEAVED_CY    3
+#define H264E_MACROBLOCK        4
+
+typedef struct _H264E_USR {
+    char*            pKernelBinFileName;
+    char*            pDeviceName;
+    char*            pKernelName;
+    bool             bCLIsOK;
+    cl_context        context;
+    cl_command_queue  commands;
+    cl_kernel         kernel;
+    int               FrameX;
+    int               FrameY;
+    int               ScanFormat;
+    int               PackPlanarSeparated;
+    int               HWOptimizeScanFormat;
+    int               Timescale;
+    int               Nuit;
+    int               Level;
+    int               GOPLength;
+    int               IntraRefreshCol;
+    int               MBRowsPerSlice;
+    int               DeblockingFilter;
+    int               DBRestrict;
+    int               MVRestrict;
+    int               QPI;
+    int               QPP;
+    int               QPChrOffset;
+    int               Bitrate;
+    int               BUSize;
+    int               RCFIFOSize;
+    int               UseCABAC;
+    void*             UserContext;
+    bool              bUseUserContextAsCfg;
+
+} H264E_USR;
+
+typedef struct _H264E_TIME {
+
+    long int WriteDeviceBufferTime;
+    long int ExecuteKernelTime;
+    long int EncodeAndReadNALSizeTime;
+    long int ReadNALBufferTime;
+
+    double   WriteDeviceBufferTimeMS;
+    double   ExecuteKernelTimeMS;
+    double   EncodeAndReadNALSizeTimeMS;
+    double   ReadNALBufferTimeMS;
+
+    double   WriteDeviceBufferTotalTimeMS;
+    double   ExecuteKernelTotalTimeMS;
+    double   EncodeAndReadNALSizeTotalTimeMS;
+    double   ReadNALBufferTotalTimeMS;
+
+    unsigned long int nFrames;
+
+} H264E_TIME;
+
+H264E_KERNEL_HANDLE h264eInit(const H264E_USR* ph264eUsr);
+int                 h264eEnd (H264E_KERNEL_HANDLE h264eKernelHandle);
+
+int h264eEncodeFrame(H264E_KERNEL_HANDLE h264eKernelHandle,
+                     char* pHostYBuffer,
+                     char* pHostUBuffer,
+                     char* pHostVBuffer,
+                     char* pHostNALBuffer,
+                     int   frame_count
+                     );
+
+void h264eSendEncodeFrame(H264E_KERNEL_HANDLE h264eKernelHandle,
+                          AVFrame *pic, char *pCurrentBuffer);
+
+int h264eRecvEncodeFrame(H264E_KERNEL_HANDLE h264eKernelHandle,
+                                                 H264E_RSP_MSG           *rspMsg);
+
+const char* h264eReadStatusRegs(H264E_KERNEL_HANDLE h264eKernelHandle);
+int               h264eSetUserContext(H264E_KERNEL_HANDLE h264eKernelHandle, void* Context);
+void*       h264eGetUserContext(H264E_KERNEL_HANDLE h264eKernelHandle);
+const H264E_TIME* h264eGetStats      (H264E_KERNEL_HANDLE h264eKernelHandle);
+
+#define H264E_OK                 0
+#define H264E_END_ERR           -1
+#define H264E_BAD_KERNEL_HANDLE -2
+#define H264E_BAD_COMP_PTR      -3
+#define H264E_BAD_NAL_PTR       -4
+#define H264E_CORE_ERR          -5
+#define H264E_INITSW_ERR        -6
+#define H264E_INITHW_ERR        -7
+#define H264E_DEVBUFF_ERR       -8
+#define H264E_DEVBUFF_WR_ERR    -9
+#define H264E_DEVBUFF_RD_ERR    -10
+#define H264E_ARGSET_ERR        -11
+#define H264E_EXEC_ERR          -12
+#define H264E_CL_INIT_ERR       -13
+#define H264E_CL_END_ERR        -14
+#define H264E_DEFER_DATA        -15
+
+#endif // _H264E_SDACCEL_H_

new file libavcodec/xlnx_h264/h264e_sdaccel_private.h

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#ifndef _H264E_SDACCEL_PRIVATE_H_
+#define _H264E_SDACCEL_PRIVATE_H_
+
+#include <CL/opencl.h>
+#include <pthread.h>
+#include <libavutil/threadmessage.h>
+#include "libavcodec/avcodec.h"
+#include "h264e_sdaccel.h"
+
+#define CLEAR_RT          0
+#define CONTROL_RT        1
+#define STATUS_RT         2
+#define CONTROL_STATUS_RT 3
+#define CONTROL_ENCODE_RT 4
+#define ENCODE_RT         5
+#define TEST0_RT          6
+#define TEST1_RT          7
+#define NO_CHANGE_FRMT                     0
+#define PLANAR_TO_INTERLEAVED_FRMT         1
+#define INTERLEAVED_YC_TO_INTERLEAVED_FRMT 2
+#define INTERLEAVED_CY_TO_INTERLEAVED_FRMT 3
+
+typedef struct _H264E_CL {
+
+    char*            pDeviceName;
+    char*            pKernelBinFileName;
+    char             pKernelName[64];
+    cl_platform_id   platform_id;
+    cl_device_id     device_id;
+    cl_context       context;
+    cl_command_queue commands;
+    cl_program       program;
+    cl_kernel        kernel;
+    cl_event         kernel_event;
+        int              kernel_num;
+    bool             bCLIsOK;
+    bool             bWRDeviceBufferFailed;
+    bool             bRDDeviceBufferFailed;
+    bool             bSetKernelArgsFailed;
+    bool             bExecuteKernelFailed;
+
+} H264E_CL;
+
+typedef struct _H264E_CNTX {
+
+    cl_mem DeviceBuffer;
+    int    DeviceBufferSizeInBytes;
+    int    CRegsPtr;
+    int    SRegsPtr;
+    int    VDataIdx;
+    int    VDataPtr[2];
+    int    NALPtr;
+    int    Format;
+    char   CRegsFile[256];
+    char   SRegsFile[256];
+    char*  pPackedPlanarHostBuffer;
+
+} H264E_CNTX;
+
+typedef struct _H264E_KCFG {
+
+    H264E_USR                    h264eUsr;
+    H264E_CL                     h264eCL;
+    H264E_CNTX                   h264eCntx;
+    H264E_TIME           h264eTime;
+        pthread_t                        h264eThread;
+        AVThreadMessageQueue *h264eReqMsgQ;
+        AVThreadMessageQueue *h264eRspMsgQ;
+
+} H264E_KCFG;
+
+typedef struct _H264E_REQ_MSG {
+        char    *pHostYBuffer;
+        char    *pHostUBuffer;
+        char    *pHostVBuffer;
+        char    *pHostNALBuffer;
+        AVFrame *pFrame;
+        int              endThread;
+} H264E_REQ_MSG;
+
+typedef struct _H264E_RSP_MSG {
+        int              NalSizeInBytes;
+        char    *pHostNALBuffer;
+        AVFrame *pFrame;
+} H264E_RSP_MSG;
+
+#endif // _H264E_SDACCEL_PRIVATE_H_

new file libavcodec/xlnx_h264/h264e_sdaccel_utils.c

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <pthread.h>
+#include "libavutil/log.h"
+#include "h264e_sdaccel.h"
+#include "h264e_sdaccel_utils.h"
+
+static pthread_mutex_t g_cl_mutex = PTHREAD_MUTEX_INITIALIZER;
+static cl_program      g_cl_program;
+static int             g_is_loaded = 0;
+
+#define TARGET_DEVICE "xilinx:xil-accel-rd-vu9p:4ddr-xpr:4.1"
+
+#if defined(SDX_PLATFORM) && !defined(TARGET_DEVICE)
+   #define STR_VALUE(arg)   #arg
+   #define GET_STRING(name) STR_VALUE(name)
+   #define TARGET_DEVICE    GET_STRING(SDX_PLATFORM)
+#endif
+
+int load_file_to_memory(const char *filename, unsigned char **result);
+
+int h264e_cl_init(H264E_CL* ph264eCL)
+{
+    char             cl_platform_vendor[1001];
+    char             target_device_name[2048];
+    char             kernel_name[2048];
+   cl_platform_id   platform_id;
+   cl_device_id     device_id;
+   cl_context       context;
+   cl_command_queue commands;
+   cl_kernel        kernel;
+    cl_int           err;
+    const char* default_target_device_name = TARGET_DEVICE;
+    const char* default_kernel_name        = "ALMA_H264_E_KERNEL";
+   ph264eCL->bCLIsOK = false;
+    if(ph264eCL->pDeviceName != NULL)
+   {
+        strcpy(target_device_name, ph264eCL->pDeviceName);
+   }
+    else
+    {
+        strcpy(target_device_name, default_target_device_name);
+    }
+    if(ph264eCL->pKernelName != NULL)
+    {
+        strcpy(kernel_name, ph264eCL->pKernelName);
+    }
+    else
+    {
+        strcpy(kernel_name, default_kernel_name);
+    }
+   cl_platform_id platforms[16];
+   cl_uint        platform_count;
+   int            platform_found = 0;
+   err = clGetPlatformIDs(16, platforms, &platform_count);
+   if(err != CL_SUCCESS)
+   {
+      H264EMSG(printf("ERROR: Failed to find an OpenCL platform\n"));
+        return H264E_CL_INIT_ERR;
+   }
+   H264EMSG(printf("INFO: Found %d platforms\n", platform_count));
+   for(unsigned int iplat=0; iplat<platform_count; iplat++)
+   {
+      err = clGetPlatformInfo(platforms[iplat], CL_PLATFORM_VENDOR, 1000, (void *)cl_platform_vendor,NULL);
+
+      if(err != CL_SUCCESS)
+      {
+            H264EMSG(printf("ERROR: clGetPlatformInfo(CL_PLATFORM_VENDOR) failed\n"));
+            return H264E_CL_INIT_ERR;
+      }
+
+      if(strcmp(cl_platform_vendor, "Xilinx") == 0)
+      {
+         H264EMSG(printf("INFO: Selected platform %d from %s\n", iplat, cl_platform_vendor));
+         platform_id = platforms[iplat];
+         platform_found = 1;
+      }
+   }
+   if(!platform_found)
+   {
+      H264EMSG(printf("ERROR: Platform Xilinx not found\n"));
+        return H264E_CL_INIT_ERR;
+   }
+   cl_uint      num_devices;
+   unsigned int device_found = 0;
+   cl_device_id devices[16];
+   char         cl_device_name[1001];
+   err = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ACCELERATOR, 16, devices, &num_devices);
+   if(err != CL_SUCCESS)
+   {
+      H264EMSG(printf("ERROR: Failed to create a device group\n"));
+        return H264E_CL_INIT_ERR;
+   }
+   for(unsigned int i=0; i<num_devices; i++)
+   {
+      err = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, 1024, cl_device_name, 0);
+
+      if(err != CL_SUCCESS)
+      {
+         H264EMSG(printf("ERROR: Failed to get device name for device %d\n", i));
+            return H264E_CL_INIT_ERR;
+      }
+      else
+      {
+         H264EMSG(printf("INFO: CL_DEVICE_NAME[%d / %d] = %s\n", i, num_devices, cl_device_name));
+      }
+
+      if(strcmp(cl_device_name, target_device_name) == 0)
+      {
+         device_id = devices[i];
+         device_found = 1;
+         H264EMSG(printf("INFO: Selected %s as the target device\n", cl_device_name));
+      }
+   }
+   if(!device_found)
+   {
+      H264EMSG(printf("ERROR: Target device %s not found\n", target_device_name));
+        return H264E_CL_INIT_ERR;
+   }
+   context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
+   if(!context)
+   {
+      H264EMSG(printf("ERROR: Failed to create a compute context\n"));
+        return H264E_CL_INIT_ERR;
+   }
+   commands = clCreateCommandQueue(context, device_id, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
+   if(!commands)
+   {
+      H264EMSG(printf("ERROR: Failed to create commands queue\n"));
+      return H264E_CL_INIT_ERR;
+   }
+   pthread_mutex_lock(&g_cl_mutex);
+   cl_int status;
+   unsigned char *kernelbinary;
+   if (!g_is_loaded)
+   {
+                H264EMSG(printf("INFO: loading kernel binary %s\n", ph264eCL->pKernelFileName));
+        int n_i0 = load_file_to_memory(ph264eCL->pKernelBinFileName, (unsigned char **) &kernelbinary);
+        if(n_i0 < 0)
+        {
+            H264EMSG(printf("ERROR: failed to load kernel binary from: %s\n", ph264eCL->pKernelBinFileName));
+            return H264E_CL_INIT_ERR;
+        }
+        size_t n0 = n_i0;
+            g_cl_program = clCreateProgramWithBinary(context,
+                                                                                                 1,
+                                                                                                 &device_id,
+                                                                                                 &n0,
+                                                                                                 (const unsigned char **) &kernelbinary,
+                                                                                                 &status,
+                                                                                                 &err);
+                if((!g_cl_program) || (err!=CL_SUCCESS))
+                {
+                        H264EMSG(printf("ERROR: Failed to create compute program from binary %d\n", err));
+            return H264E_CL_INIT_ERR;
+                }
+            free(kernelbinary);
+                err = clBuildProgram(g_cl_program, 0, NULL, NULL, NULL, NULL);
+                if(err != CL_SUCCESS)
+                {
+                        size_t len;
+                        char   buffer[2048];
+
+                        H264EMSG(printf("ERROR: Failed to build program executable\n"));
+                        clGetProgramBuildInfo(g_cl_program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
+                        H264EMSG(printf("%s\n", buffer));
+            return H264E_CL_INIT_ERR;
+                }
+                        g_is_loaded = 1;
+                }
+                else
+                {
+                        clRetainProgram(g_cl_program);
+                }
+                pthread_mutex_unlock(&g_cl_mutex);
+            kernel = clCreateKernel(g_cl_program, ph264eCL->pKernelName, &err);
+   if(!kernel || err != CL_SUCCESS)
+   {
+      H264EMSG(printf("ERROR: Failed to create compute kernel\n"));
+      return -1;
+   }
+   ph264eCL->platform_id = platform_id;
+   ph264eCL->device_id   = device_id;
+   ph264eCL->context     = context;
+   ph264eCL->commands    = commands;
+   ph264eCL->program     = g_cl_program;
+   ph264eCL->kernel      = kernel;
+   ph264eCL->bCLIsOK     = true;
+    return H264E_OK;
+}
+
+int h264e_cl_end(H264E_CL* ph264eCL)
+{
+    cl_int err;
+    int    retval = H264E_OK;
+    err = clReleaseProgram(ph264eCL->program);
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to release program\n"));
+        retval = H264E_CL_END_ERR;
+    }
+    else
+    {
+        H264EMSG(printf("INFO: Released program\n"));
+    }
+   clReleaseKernel(ph264eCL->kernel);
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to release kernel\n"));
+        retval = H264E_CL_END_ERR;
+    }
+    else
+    {
+        H264EMSG(printf("INFO: Released kernel\n"));
+    }
+   clReleaseCommandQueue(ph264eCL->commands);
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to release commands queue\n"));
+        retval = H264E_CL_END_ERR;
+    }
+    else
+    {
+        H264EMSG(printf("INFO: Released commands queue\n"));
+    }
+   clReleaseContext(ph264eCL->context);
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to release context\n"));
+        retval = H264E_CL_END_ERR;
+    }
+    else
+    {
+        H264EMSG(printf("INFO: Released context\n"));
+    }
+    clReleaseDevice(ph264eCL->device_id);
+    if(err != CL_SUCCESS)
+    {
+        H264EMSG(printf("ERROR: Failed to release device\n"));
+        retval = H264E_CL_END_ERR;
+    }
+    else
+    {
+        H264EMSG(printf("INFO: Released device\n"));
+}
+    return retval;
+}
+
+int load_file_to_memory(const char *filename, unsigned char **result)
+{
+   int size = 0;
+   FILE *f = fopen(filename, "rb");
+   if(f == NULL)
+   {
+       *result = NULL;
+        return -1; // -1 means file opening fail
+   }
+   fseek(f, 0, SEEK_END);
+   size = ftell(f);
+   fseek(f, 0, SEEK_SET);
+   *result = malloc(size);
+   if(size != fread(*result, sizeof(char), size, f))
+   {
+      free(*result);
+      return -2; // -2 means file reading fail
+   }
+   fclose(f);
+   return size;
+ }

new file libavcodec/xlnx_h264/h264e_sdaccel_utils.h

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#ifndef _H264E_SDACCEL_UTILS_H_
+#define _H264E_SDACCEL_UTILS_H_
+
+#include "h264e_sdaccel_private.h"
+
+int h264e_cl_init(H264E_CL* ph264eCL);
+int h264e_cl_end (H264E_CL* ph264eCL);
+
+#endif // _H264E_SDACCEL_UTILS_H_

new file libavcodec/xlnx_h264/hw_h264.c

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#include <inttypes.h>
+#include <stdlib.h>
+#include <linux/limits.h>
+#include <unistd.h>
+#include <string.h>
+#include <assert.h>
+#include "libavutil/opencl.h"
+#include "libavutil/log.h"
+#include "h264e_sdaccel.h"
+#include "h264e_sdaccel_private.h"
+#include "h264e_csregs.h"
+#include "xlnx_h264_config.h"
+
+void ff_init_h264_xlnx_hw(AVCodecContext *avctx)
+{
+    H264E_USR           h264eUsr;
+        unsigned char* xcl_path = getenv("XLNX_XCLBIN_PATH");
+        if (NULL == xcl_path)
+        {
+       av_log(NULL, AV_LOG_INFO, " * XLNX_XCLBIN_PATH not found\n");
+       return;
+        }
+        char *xcl_fullpath = malloc(PATH_MAX);
+        snprintf(xcl_fullpath, PATH_MAX, "%s/h264_8k1c.xclbin", xcl_path);
+        h264eUsr.pKernelBinFileName   = xcl_fullpath;
+        h264eUsr.pDeviceName              = "xilinx:xil-accel-rd-vu9p:4ddr-xpr:4.1";
+        h264eUsr.pKernelName              = "ALMA_H264_E_KERNEL";
+    h264eUsr.bCLIsOK              = false;
+    h264eUsr.FrameX               = avctx->width;
+    h264eUsr.FrameY               = avctx->height;
+    h264eUsr.ScanFormat           = H264E_PLANAR_SEPARATED;
+    h264eUsr.PackPlanarSeparated  = 0;
+    h264eUsr.HWOptimizeScanFormat = 1;
+        h264eUsr.Timescale                = 60000;
+        h264eUsr.Nuit                 = 1000;
+        h264eUsr.Level                = 42;
+        h264eUsr.GOPLength                = avctx->gop_size;
+        h264eUsr.IntraRefreshCol      = -1;
+        h264eUsr.MBRowsPerSlice       = -1;
+        h264eUsr.DeblockingFilter     = -1;
+        h264eUsr.DBRestrict           = -1;
+        h264eUsr.MVRestrict           = -1;
+        h264eUsr.QPI                  = avctx->global_quality;
+        h264eUsr.QPP                  = avctx->global_quality;
+        h264eUsr.QPChrOffset          = -1;
+        h264eUsr.Bitrate              = (avctx->bit_rate)/1000;
+        h264eUsr.BUSize               = -1;
+        h264eUsr.RCFIFOSize           = -1;
+        h264eUsr.UseCABAC             = -1;
+    h264eUsr.bUseUserContextAsCfg = false;
+    avctx->priv_data = h264eInit(&h264eUsr);
+}
+
+void h264eSendEncodeFrame(H264E_KERNEL_HANDLE h264eKernelHandle,
+                                                  AVFrame *pic, char *pCurrentBuffer)
+{
+        H264E_KCFG *ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+        H264E_REQ_MSG reqMsg;
+        reqMsg.pHostYBuffer = pic->data[Y];
+        reqMsg.pHostUBuffer = pic->data[U];
+        reqMsg.pHostVBuffer = pic->data[V];
+        reqMsg.pHostNALBuffer = pCurrentBuffer;
+        reqMsg.pFrame           = pic;
+        reqMsg.endThread        = 0;
+        av_thread_message_queue_send(ph264eKCFG->h264eReqMsgQ, &reqMsg, 0);
+}
+
+int h264eRecvEncodeFrame(H264E_KERNEL_HANDLE h264eKernelHandle,
+                                                 H264E_RSP_MSG          *rsp)
+{
+        H264E_KCFG *ph264eKCFG = (H264E_KCFG*)h264eKernelHandle;
+        int rc;
+        rc = av_thread_message_queue_recv(ph264eKCFG->h264eRspMsgQ, rsp,
+                                                                  AV_THREAD_MESSAGE_NONBLOCK);
+        if (rc == AVERROR(EAGAIN))
+                return 0;
+        return 1;
+}
+
+void ff_uninit_h264_xlnx_hw(AVCodecContext *avctx)
+{
+   h264eEnd(avctx->priv_data);
+}
+
+

new file libavcodec/xlnx_h264/xlnx_h264_config.h

@@ +1 @@
+/*
+-- Copyright notice:
+-- Copyright (c) 2017 Alma Technologies S.A.
+-- www.alma-technologies.com
+-- All rights reserved worldwide.
+
+-- This copy is made available under the terms of the GNU General Public
+-- License version 2.1 as published by the Free Software Foundation; either
+-- version 2.1 of the License, or (at your option) any later version.
+--
+-- This copy 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. Please See the GNU
+-- General Public License for more details.
+--
+-- You should have received a copy of the GNU General Public
+-- License along with FFmpeg; if not, write to the Free Software
+-- Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+*/
+
+#ifndef _XLNX_H264_CONFIG_H_
+#define _XLNX_H264_CONFIG_H_
+
+#include "libavcodec/avcodec.h"
+
+#define OCLCHECK(method, ... )                                                 \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n",                   \
+               av_opencl_errstr(status));                                      \
+        return;                                                                \
+    }                                                                          \
+} while (0)
+
+#define OCLCHECK1(method, ... )                                                \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n",                   \
+               av_opencl_errstr(status));                                      \
+        return AVERROR_EXTERNAL;                                               \
+    }                                                                          \
+} while (0)
+
+#define CREATEBUF(out, flags, size)                                            \
+do {                                                                           \
+    out = clCreateBuffer(cl_env->context, flags, size, NULL, &status);         \
+    if (status != CL_SUCCESS) {                                                \
+        av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n");        \
+        return;                                                                \
+    }                                                                          \
+} while (0)
+
+typedef enum
+{
+        Y = 0,
+        U = 1,
+        V = 2
+}col_fmt_t;
+
+void ff_init_h264_xlnx_hw(AVCodecContext *c);
+void ff_uninit_h264_xlnx_hw(AVCodecContext *c);
+uint8_t *ff_h264_xlnx_hw(AVCodecContext *avctx, AVFrame *pic,
+                                           char *pCurrentBuffer, uint32_t *dstSize);
+
+#endif //_XLNX_H264_CONFIG_H_

new file libavcodec/xlnx_h264_enc.c

@@ +1 @@
+/*
+ * Xilinx H.264 Alma encoder
+ *
+ * Copyright (C) 2017 Xilinx Inc.
+ *
+ * This file is part of FFmpeg.
+ *
+ * FFmpeg is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * FFmpeg 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
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with FFmpeg; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+ */
+#include <stdio.h>
+#include "avcodec.h"
+#include "internal.h"
+#include "xlnx_h264/xlnx_h264_config.h"
+#include "xlnx_h264/h264e_sdaccel_private.h"
+
+static int xlnx_h264_encode_init(AVCodecContext *avctx)
+{
+    //printf("xlnx_h264_encode_init IN\n");
+    switch(avctx->pix_fmt)
+    {
+         case AV_PIX_FMT_YUV420P:
+            avctx->bits_per_raw_sample = 8;
+            break;
+         default:
+            av_log(NULL, AV_LOG_ERROR,"ERROR: Pixel format not supported...\n");
+            return AVERROR(EINVAL);
+            break;
+    }
+    ff_init_h264_xlnx_hw(avctx);
+    return 0;
+}
+
+static int xlnx_h264_encode_frame(AVCodecContext *avctx, AVPacket *pkt, const AVFrame *pic, int *got_packet)
+{
+
+    //av_log(NULL, AV_LOG_DEBUG, "xlnx_h264_encode_frame IN\n");
+    int EntrySize = avctx->bits_per_raw_sample == 8 ? 1 : 2;
+    int EntryCount = 0;
+    uint64_t EncodeBuffSize = 0;
+        uint8_t *pBuffer = NULL;
+        AVFrame *pFrame;
+        H264E_RSP_MSG   rspMsg;
+        int             rc;
+
+        if (!pic)
+        {
+                *got_packet = 0;
+                return 0;
+        }
+
+    switch(pic->format)
+    {
+        case AV_PIX_FMT_YUV420P:
+            EntryCount = 2;
+            break;
+        default:
+            av_log(NULL, AV_LOG_ERROR,"Error. Encoding unknown pixel format.\n");
+            return AVERROR_INVALIDDATA;
+            break;
+    }
+    
+        pFrame = av_frame_clone(pic);
+    EncodeBuffSize = EntryCount * (pic->width) * (pic->height) * EntrySize;
+
+        pBuffer = av_malloc(EncodeBuffSize);
+
+        // Send the encode request
+        h264eSendEncodeFrame(avctx->priv_data, pFrame, pBuffer);
+
+        // Check if the encode has completed
+        rc = h264eRecvEncodeFrame(avctx->priv_data, &rspMsg);
+
+        if (rc == 1)
+        {
+                *got_packet = 1;
+                av_packet_from_data(pkt, rspMsg.pHostNALBuffer,
+                                                        rspMsg.NalSizeInBytes);
+                pkt->pts = pkt->dts = (int64_t)rspMsg.pFrame->pts;
+                av_frame_unref(rspMsg.pFrame);
+                av_frame_free(&rspMsg.pFrame);
+        }
+        else
+        {
+        *got_packet = 0;
+        }
+        
+    return 0;
+}
+
+static int xlnx_h264_encode_close(AVCodecContext *avctx)
+{
+    printf("xlnx_h264_encode_close IN\n");
+    ff_uninit_h264_xlnx_hw(avctx);
+    return 0;
+}
+
+AVCodec ff_xlnx_h264_encoder = {
+    .name           = "xlnx_h264_enc",
+    .long_name      = NULL_IF_CONFIG_SMALL("Xilinx Advanced Video Codec (H.264) Encoder"),
+    .type           = AVMEDIA_TYPE_VIDEO,
+    .id             = AV_CODEC_ID_H264,
+        .priv_data_size = sizeof(H264E_KCFG),
+    .init           = xlnx_h264_encode_init,
+    .encode2        = xlnx_h264_encode_frame,
+    .close          = xlnx_h264_encode_close,
+    .capabilities   = CODEC_CAP_DR1 | CODEC_CAP_DELAY,
+    .pix_fmts       = (const enum AVPixelFormat[]) {
+                                      AV_PIX_FMT_YUV420P, AV_PIX_FMT_NONE
+                                  },
+};

new file libavcodec/xlnx_hevc/Makefile

@@ +1 @@
+OBJS                            += xlnx_hevc/hw_hevc.o

new file libavcodec/xlnx_hevc/hw_hevc.c

@@ +1 @@
+//
+// NOTE: currently using ICD for DSA 2.0 50Mhz ...
+// TODO: update this when new DSA is available
+//
+
+#include <inttypes.h>
+#include <stdlib.h>
+#include <linux/limits.h>
+#include <unistd.h>
+#include "libavutil/opencl.h"
+#include "libavutil/log.h"
+#include "xlnx_hevc_config.h"
+
+// Header based loading
+//#include "opencl_image.h"
+
+//Globals - for opencl
+static AVOpenCLExternalEnv *cl_env;
+static cl_kernel kernel = NULL;
+static cl_program program = NULL;
+
+// OpenCL Buffers
+static cl_mem CSR3_INPUT_FP_Y_V; 
+static cl_mem CSR4_INPUT_FP_U_V; 
+static cl_mem CSR5_INPUT_FP_V_V; 
+static cl_mem CSR6_NEXT_FP_Y_V; 
+static cl_mem CSR7_NEXT_FP_U_V; 
+static cl_mem CSR8_NEXT_FP_V_V;
+static cl_mem CSR12_RECON_BASE_V; 
+static cl_mem CSR13_BITSTREAM_BASE_V; 
+static cl_mem CSR14_BITSTREAM_LEN_V;
+
+static cl_int status;
+static unsigned int CurSourceImageSize[3] = { 0, 0, 0 };
+static unsigned int EncodedFrameSize = 0;
+static unsigned int EncodedFrameLengthSize = 0;
+static int frameNumber = 0;
+
+int load_file_to_memory_hevc(const char *filename, unsigned char **result);
+
+//
+// HEVC init function
+//
+void ff_init_hevc_xlnx_hw(AVCodecContext *avctx)
+{
+    unsigned char *kernelbinary = 0;
+    size_t kernel_len = 0;
+    AVOpenCLDeviceList *device_list;
+    AVOpenCLDeviceNode *device_node = NULL;
+    cl_context_properties cps[3];
+    int i;
+    int nb_devices;
+    int arg = 0;
+    // Header based loading (opencl_image.h)
+    //kernelbinary = hevc_xclbin;
+    //kernel_len = hevc_xclbin_len;
+
+    //
+    // Setup NGCodec HEVC core defaults
+    //
+    unsigned int CSR1_FRAME_WIDTH_V = avctx->width;
+    unsigned int CSR2_FRAME_HEIGHT_V = avctx->height;
+    unsigned int CSR9_BITRATE_V = 0;
+    unsigned int CSR10_FIXED_QP_V = 0;
+    unsigned int CSR11_INTRA_PERIOD_V = avctx->gop_size; 
+    unsigned int CSR15_RESET_ENCODER_V = 0;
+
+    av_log(NULL, AV_LOG_INFO, " * Xilinx HEVC HW accelerator init ***\n");
+    if (avctx->bit_rate > 0)
+    {
+         // Constant bitrate mode
+         CSR9_BITRATE_V = avctx->bit_rate;
+         av_log(NULL, AV_LOG_INFO, " * Encoding in constant bitrate mode***\n");
+    }
+    else 
+    {
+         // Fixed QP mode
+         CSR10_FIXED_QP_V = avctx->global_quality;
+         av_log(NULL, AV_LOG_INFO, " * Encoding in fixed QP mode***\n");
+    }
+
+    //
+    // Setup OpenCL context
+    //
+    cl_env = av_opencl_alloc_external_env();
+    cl_platform_id platforms[16] = { 0 };
+        cl_device_id devices[16];
+        char platformName[256];
+        char deviceName[256];
+        cl_uint platformCount = 0;
+        cl_int err = clGetPlatformIDs(0, 0, &platformCount);
+        err = clGetPlatformIDs(16, platforms, &platformCount);
+        if (err != CL_SUCCESS) {
+                //printf("%s\n", oclErrorCode(err));
+                return;// hardware;
+    }
+        cl_env->device_type = CL_DEVICE_TYPE_ACCELERATOR;
+    for (cl_uint i = 0; i < platformCount; i++) {
+                err = clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, 256, platformName, 0);
+                if (err != CL_SUCCESS) {
+                        //printf("%s\n", oclErrorCode(err));
+                        return;// hardware;
+                }
+                cl_uint deviceCount = 0;
+                err = clGetDeviceIDs(platforms[i], cl_env->device_type, 16, devices, &deviceCount);
+                if ((err != CL_SUCCESS) || (deviceCount == 0)) {
+                        continue;
+                }
+
+                err = clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 256, deviceName, 0);
+                if (err != CL_SUCCESS) {
+                        //printf("%s\n", oclErrorCode(err));
+                        return;// hardware;
+                }
+
+                cl_context_properties contextData[3] = {CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i], 0};
+                cl_context context = clCreateContextFromType(contextData, cl_env->device_type, 0, 0, &err);
+                if (err != CL_SUCCESS) {
+                        continue;
+                }
+                cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, &err);
+                if (err != CL_SUCCESS) {
+                        //printf("%s\n", oclErrorCode(err));
+                        return; //hardware;
+                }
+                cl_env->platform_id = platforms[i];
+                cl_env->context = context;
+                cl_env->command_queue = queue;
+                cl_env->device_type = CL_DEVICE_TYPE_ACCELERATOR;
+        cl_env->device_id = devices[0];
+        break;
+        }
+
+    //
+    // Create the program
+    //
+        unsigned char* xcl_path = getenv("XLNX_XCLBIN_PATH");
+        if (NULL == xcl_path)
+        {
+       av_log(NULL, AV_LOG_INFO, " * XLNX_XCLBIN_PATH not found\n");
+       return;
+        }
+    
+        char xcl_fullpath[PATH_MAX];
+        snprintf(xcl_fullpath, sizeof(xcl_fullpath), "%s/hevc_1k1c.xclbin", xcl_path); 
+    av_log(NULL, AV_LOG_INFO, " * loading kernel from file %s\n", xcl_fullpath);
+        load_file_to_memory_hevc(xcl_fullpath, &kernelbinary);
+    program = clCreateProgramWithBinary(cl_env->context, 1, &(cl_env->device_id), 
+                   &kernel_len, (const unsigned char**)&kernelbinary, &status, NULL);
+    if( status != CL_SUCCESS || !program )
+    {
+        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create fpga accelerator program\n");
+        return;
+    }
+
+    //
+    // Build the program
+    //
+    av_log(NULL, AV_LOG_INFO," * Building the program\n");
+    status = clBuildProgram(program, 1, &(cl_env->device_id), NULL, NULL, NULL);
+    if( status != CL_SUCCESS )
+    {
+        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to build program\n");
+        return;
+    }
+
+    //
+    // Create the kernel
+    //
+    av_log(NULL, AV_LOG_INFO," * Creating the HEVC kernel\n");
+    kernel = clCreateKernel(program, "krnl_ngcodec_hevc_0", &status);   //"HEVC_kernel"
+    if( status != CL_SUCCESS )
+    {
+        av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create kernel\n");
+        return;
+    }
+
+    //
+    // Setup Source and EncodedFrame Buffers
+    //
+    CurSourceImageSize[Y] = (avctx->width * avctx->height);
+    CurSourceImageSize[U] = (avctx->width * avctx->height) >> 2;
+    CurSourceImageSize[V] = (avctx->width * avctx->height) >> 2;
+        
+    // Encoded Frame data size - hard code to allocate host memory as this varies!
+    EncodedFrameSize = (avctx->width * avctx->height * 3) >> 1;
+    EncodedFrameLengthSize = 1;
+    // av_log(NULL, AV_LOG_INFO," * Image Sizes: Source[YUV]=[%d,%d,%d] EncodedFrame=%d\n", 
+    //          CurSourceImageSize[Y], CurSourceImageSize[U], CurSourceImageSize[V], EncodedFrameSize);
+
+    //
+    // Setup OpenCL Buffers and Kernel Parameters
+    //
+    // av_log(NULL, AV_LOG_INFO," * Creating the kernel buffers\n");
+    CREATEBUF(CSR3_INPUT_FP_Y_V, CL_MEM_READ_ONLY,  CurSourceImageSize[Y] * sizeof(unsigned char));
+    CREATEBUF(CSR4_INPUT_FP_U_V, CL_MEM_READ_ONLY,  CurSourceImageSize[U] * sizeof(unsigned char));
+    CREATEBUF(CSR5_INPUT_FP_V_V, CL_MEM_READ_ONLY,  CurSourceImageSize[V] * sizeof(unsigned char));
+    CREATEBUF(CSR6_NEXT_FP_Y_V, CL_MEM_READ_ONLY,  CurSourceImageSize[Y] * sizeof(unsigned char));
+    CREATEBUF(CSR7_NEXT_FP_U_V, CL_MEM_READ_ONLY,  CurSourceImageSize[U] * sizeof(unsigned char));
+    // Padding to make reference base pointer of reconstructed picture 128KB aligned
+    CREATEBUF(CSR8_NEXT_FP_V_V, CL_MEM_READ_ONLY,  (CurSourceImageSize[V]+ 48*128*1024-1920*1080*3) * sizeof(unsigned char));
+    // Reference frame requires at least 16MB space in DDR
+    CREATEBUF(CSR12_RECON_BASE_V, CL_MEM_READ_WRITE,  16*1024*1024);
+    CREATEBUF(CSR13_BITSTREAM_BASE_V, CL_MEM_WRITE_ONLY, EncodedFrameSize * sizeof(unsigned char));
+    CREATEBUF(CSR14_BITSTREAM_LEN_V, CL_MEM_WRITE_ONLY, EncodedFrameLengthSize * sizeof(unsigned long long));
+    
+    //
+    // Send Kernel Parameters
+    //
+    // av_log(NULL, AV_LOG_INFO," * Sending parameters to HEVC kernel\n"); 
+    // av_log(NULL, AV_LOG_INFO," * CSR1_FRAME_WIDTH_V:  0x%08x\n", CSR1_FRAME_WIDTH_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR2_FRAME_HEIGHT_V: 0x%08x\n", CSR2_FRAME_HEIGHT_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR3_INPUT_FP_Y_V addr %p\n", &CSR3_INPUT_FP_Y_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR4_INPUT_FP_U_V addr %p\n", &CSR4_INPUT_FP_U_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR5_INPUT_FP_V_V addr %p\n", &CSR5_INPUT_FP_V_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR6_NEXT_FP_Y_V addr %p\n", &CSR6_NEXT_FP_Y_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR7_NEXT_FP_U_V addr %p\n", &CSR7_NEXT_FP_U_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR8_NEXT_FP_V_V addr %p\n", &CSR8_NEXT_FP_V_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR9_BITRATE_V:  0x%08x\n", CSR9_BITRATE_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR10_FIXED_QP_V:  0x%08x\n", CSR10_FIXED_QP_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR11_INTRA_PERIOD_V:  0x%08x\n", CSR11_INTRA_PERIOD_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR12_RECON_BASE_V addr %p\n", &CSR12_RECON_BASE_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR13_BITSTREAM_BASE_V addr %p\n", &CSR13_BITSTREAM_BASE_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR14_BITSTREAM_LEN_V addr %p\n", &CSR14_BITSTREAM_LEN_V);
+    // av_log(NULL, AV_LOG_INFO," * CSR15_RESET_ENCODER_V:  0x%08x\n", CSR15_RESET_ENCODER_V);
+        
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(unsigned int), &CSR1_FRAME_WIDTH_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(unsigned int), &CSR2_FRAME_HEIGHT_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR3_INPUT_FP_Y_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR4_INPUT_FP_U_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR5_INPUT_FP_V_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR6_NEXT_FP_Y_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR7_NEXT_FP_U_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR8_NEXT_FP_V_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(unsigned int), &CSR9_BITRATE_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(unsigned int), &CSR10_FIXED_QP_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(unsigned int), &CSR11_INTRA_PERIOD_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR12_RECON_BASE_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR13_BITSTREAM_BASE_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &CSR14_BITSTREAM_LEN_V);
+    OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(unsigned int), &CSR15_RESET_ENCODER_V);
+    return;
+}
+
+//
+// HEVC opencl run function
+//
+int ff_hevc_xlnx_hw(AVCodecContext *avctx, AVFrame *pic, uint8_t *dst, uint32_t *dstSize)
+{
+    size_t global_work_size_1d[1] = {1};
+    size_t local_work_size_1d[1] = {1}; 
+         
+    if (frameNumber == 0)
+    {
+        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR3_INPUT_FP_Y_V, CL_TRUE, 0, CurSourceImageSize[Y] * sizeof(unsigned char), pic->data[Y], 0, NULL, NULL);
+        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR4_INPUT_FP_U_V, CL_TRUE, 0, CurSourceImageSize[U] * sizeof(unsigned char), pic->data[U], 0, NULL, NULL);
+        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR5_INPUT_FP_V_V, CL_TRUE, 0, CurSourceImageSize[V] * sizeof(unsigned char), pic->data[V], 0, NULL, NULL);  
+        // av_log(NULL, AV_LOG_INFO," * Loading frame %d\n", frameNumber);        
+        frameNumber++;
+        return 0;
+    }
+    else if (frameNumber == 1)
+    {
+        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR6_NEXT_FP_Y_V, CL_TRUE, 0, CurSourceImageSize[Y] * sizeof(unsigned char), pic->data[Y], 0, NULL, NULL);
+        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR7_NEXT_FP_U_V, CL_TRUE, 0, CurSourceImageSize[U] * sizeof(unsigned char), pic->data[U], 0, NULL, NULL);
+        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR8_NEXT_FP_V_V, CL_TRUE, 0, CurSourceImageSize[V] * sizeof(unsigned char), pic->data[V], 0, NULL, NULL);
+        // av_log(NULL, AV_LOG_INFO," * Loading frame %d\n", frameNumber);        
+        frameNumber++;
+    }
+    else 
+    {
+        // Ping-pong buffers used; switch cur and nxt buffer pointer for kernel
+        if (frameNumber%2 == 0)
+                {
+                        OCLCHECK1(clSetKernelArg, kernel, 2, sizeof(cl_mem), &CSR6_NEXT_FP_Y_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 3, sizeof(cl_mem), &CSR7_NEXT_FP_U_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 4, sizeof(cl_mem), &CSR8_NEXT_FP_V_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 5, sizeof(cl_mem), &CSR3_INPUT_FP_Y_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 6, sizeof(cl_mem), &CSR4_INPUT_FP_U_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 7, sizeof(cl_mem), &CSR5_INPUT_FP_V_V);
+                        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR3_INPUT_FP_Y_V, CL_TRUE, 0, CurSourceImageSize[Y] * sizeof(unsigned char), pic->data[Y], 0, NULL, NULL);
+                        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR4_INPUT_FP_U_V, CL_TRUE, 0, CurSourceImageSize[U] * sizeof(unsigned char), pic->data[U], 0, NULL, NULL);
+                        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR5_INPUT_FP_V_V, CL_TRUE, 0, CurSourceImageSize[V] * sizeof(unsigned char), pic->data[V], 0, NULL, NULL);
+                        // av_log(NULL, AV_LOG_INFO," * Loading frame %d\n", frameNumber); 
+                }
+                else 
+                {
+                        OCLCHECK1(clSetKernelArg, kernel, 2, sizeof(cl_mem), &CSR3_INPUT_FP_Y_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 3, sizeof(cl_mem), &CSR4_INPUT_FP_U_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 4, sizeof(cl_mem), &CSR5_INPUT_FP_V_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 5, sizeof(cl_mem), &CSR6_NEXT_FP_Y_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 6, sizeof(cl_mem), &CSR7_NEXT_FP_U_V);
+                        OCLCHECK1(clSetKernelArg, kernel, 7, sizeof(cl_mem), &CSR8_NEXT_FP_V_V);
+                        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR6_NEXT_FP_Y_V, CL_TRUE, 0, CurSourceImageSize[Y] * sizeof(unsigned char), pic->data[Y], 0, NULL, NULL); 
+                        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR7_NEXT_FP_U_V, CL_TRUE, 0, CurSourceImageSize[U] * sizeof(unsigned char), pic->data[U], 0, NULL, NULL); 
+                        OCLCHECK1(clEnqueueWriteBuffer, cl_env->command_queue, CSR8_NEXT_FP_V_V, CL_TRUE, 0, CurSourceImageSize[V] * sizeof(unsigned char), pic->data[V], 0, NULL, NULL);
+                        // av_log(NULL, AV_LOG_INFO," * Loading frame %d\n", frameNumber);  
+                }
+                frameNumber++;
+    }
+    OCLCHECK1(clEnqueueNDRangeKernel, cl_env->command_queue, kernel, 1, NULL, global_work_size_1d, local_work_size_1d, 0, NULL, NULL);
+
+    //
+    // Wait for the hw accelerator to complete
+    //
+    clFinish(cl_env->command_queue);
+    //av_log(NULL, AV_LOG_INFO," * Finished the xlnx hw accelerator kernel\n");
+
+    //
+    // Read the result
+    // Get the length of the Encoded Frame first as this varies!
+    //
+    // av_log(NULL, AV_LOG_INFO," * Enqueuing Read Buffer - CSR14_BITSTREAM_LEN_V\n");
+    OCLCHECK1(clEnqueueReadBuffer, cl_env->command_queue, CSR14_BITSTREAM_LEN_V, CL_TRUE, 0, EncodedFrameLengthSize * sizeof(unsigned int), dstSize, 0, NULL, NULL);
+    // av_log(NULL, AV_LOG_INFO," * Encoded Frame Length = %d\n", *dstSize);
+    //*dstSize = *dstSize * 4;
+
+    // av_log(NULL, AV_LOG_INFO," * Enqueuing Read Buffer - CSR13_BITSTREAM_BASE_V\n");
+    OCLCHECK1(clEnqueueReadBuffer, cl_env->command_queue, CSR13_BITSTREAM_BASE_V, CL_TRUE, 0, (*dstSize) * sizeof(unsigned char), dst, 0, NULL, NULL);
+
+    return 0;
+}
+
+//
+// HEVC exit function
+//
+void ff_uninit_hevc_xlnx_hw(AVCodecContext *avctx)
+{
+    // av_log(NULL, AV_LOG_INFO," * Tidying up ... releasing OpenCL Objects\n");
+    clReleaseKernel(kernel);
+    clReleaseProgram(program);
+    clReleaseMemObject(CSR3_INPUT_FP_Y_V);
+    clReleaseMemObject(CSR4_INPUT_FP_U_V);
+    clReleaseMemObject(CSR5_INPUT_FP_V_V);
+    clReleaseMemObject(CSR6_NEXT_FP_Y_V);
+    clReleaseMemObject(CSR7_NEXT_FP_U_V);
+    clReleaseMemObject(CSR8_NEXT_FP_V_V);
+    clReleaseMemObject(CSR13_BITSTREAM_BASE_V);
+    clReleaseMemObject(CSR14_BITSTREAM_LEN_V);  
+}
+
+int load_file_to_memory_hevc(const char *filename, unsigned char **result)
+{
+   int size = 0;
+   FILE *f = fopen(filename, "rb");
+   if(f == NULL)
+   {
+       *result = NULL;
+        return -1; // -1 means file opening fail
+   }
+   fseek(f, 0, SEEK_END);
+   size = ftell(f);
+   fseek(f, 0, SEEK_SET);
+   *result = malloc(size);
+   if(size != fread(*result, sizeof(char), size, f))
+   {
+      free(*result);
+      return -2; // -2 means file reading fail
+   }
+   fclose(f);
+   return size;
+ }
+

new file libavcodec/xlnx_hevc/xlnx_hevc_config.h

@@ +1 @@
+#ifndef _XLNX_HEVC_CONFIG_H_
+#define _XLNX_HEVC_CONFIG_H_
+
+#include "libavcodec/avcodec.h"
+
+//Macros
+#define OCLCHECK(method, ... )                                                 \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n",                   \
+               av_opencl_errstr(status));                                      \
+        return;                                                                \
+    }                                                                          \
+} while (0)
+
+#define OCLCHECK1(method, ... )                                                \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n",                   \
+               av_opencl_errstr(status));                                      \
+        return AVERROR_EXTERNAL;                                               \
+    }                                                                          \
+} while (0)
+
+#define CREATEBUF(out, flags, size)                                            \
+do {                                                                           \
+    out = clCreateBuffer(cl_env->context, flags, size, NULL, &status);         \
+    if (status != CL_SUCCESS) {                                                \
+        av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n");        \
+        return;                                                                \
+    }                                                                          \
+} while (0)
+
+//Typedefs
+typedef enum
+{
+        Y = 0,
+        U = 1,
+        V = 2
+}col_fmt_t;
+
+//Function Prototypes
+void ff_init_hevc_xlnx_hw(AVCodecContext *c);
+void ff_uninit_hevc_xlnx_hw(AVCodecContext *c);
+int ff_hevc_xlnx_hw(AVCodecContext *avctx, AVFrame *pic, uint8_t *dst, uint32_t *dstSize);
+
+#endif //_XLNX_HEVC_CONFIG_H_
+

new file libavcodec/xlnx_hevc_enc.c

@@ +1 @@
+/*
+ * Xilinx HEVC NGCodec encoder
+ *
+ * Copyright (C) 2015 Xilinx Inc.
+ *
+ * This file is part of FFmpeg.
+ *
+ * FFmpeg is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * FFmpeg 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
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with FFmpeg; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+ */
+#include <stdio.h>
+#include "avcodec.h"
+#include "internal.h"
+#include "xlnx_hw/xlnx_hw.h"
+#include "pthread.h"
+#include <sys/types.h>
+#include <sys/syscall.h>
+#include "xlnx_hw/common/inc/xlnx_queue.h"
+#include "xlnx_hw/common/inc/xlnx_hw_defines.h"
+
+typedef struct EncodeRequest {
+    xlnx_cmp_buf iBuf;
+    xlnx_hw_buf iBufHdr;
+    xlnx_data_buf out;
+    xlnx_hw_buf oBufHdr;
+    AVFrame* pic;
+}
+EncodeRequest;
+
+typedef struct XlnxHevcContext {
+    xlnx_hw hw;
+    HevcEncParam params;
+    xlnx_queue* rQ;
+    int frame;
+}
+XlnxHevcContext;
+
+static EncodeRequest* createEncReq(const AVFrame* pic, size_t aEncodeSize) {
+    EncodeRequest* req;
+    xlnx_data_buf* tmp;
+    req = av_mallocz(sizeof(EncodeRequest));
+    if (NULL == req) {
+        av_log(NULL, AV_LOG_ERROR,"Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+        return NULL;
+    }
+    //input
+    req->iBufHdr.type = ERAW;
+    req->oBufHdr.type = EDATA;
+    if (pic != NULL) {
+        req->pic = av_frame_clone(pic);
+        if (NULL == req->pic) {
+            av_log(NULL, AV_LOG_ERROR,"Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+            av_free(req);
+            return NULL;
+        }
+        tmp = (xlnx_data_buf*)&req->iBuf;
+        for (int i = 0; i < 3; i++) {
+            tmp[i].pDATA = req->pic->data[i];
+            tmp[i].offset = 0;
+            tmp[i].size = req->pic->linesize[i] * req->pic->height;
+        }
+        req->iBufHdr.ptr = &req->iBuf;
+        // output
+        req->out.pDATA = av_malloc(aEncodeSize);
+        if (NULL == req->out.pDATA) {
+            av_log(NULL, AV_LOG_ERROR,"Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+            av_frame_unref(req->pic);
+            return NULL;
+        }
+        req->out.size = aEncodeSize;
+        req->oBufHdr.ptr = &req->out;
+    } else {
+        req->iBufHdr.ptr = NULL;
+        req->oBufHdr.ptr = NULL;
+    }
+
+    return req;
+}
+
+static void destroyEncInput(EncodeRequest* aReq) {
+    if (aReq->pic) {
+        av_frame_unref(aReq->pic);
+        av_frame_free(&aReq->pic);
+    }
+}
+
+static void destroyEncOut(EncodeRequest* aReq) {
+    if(aReq->out.pDATA) av_free(aReq->out.pDATA);
+}
+
+static int destroyEncReq(EncodeRequest* aReq) {
+    destroyEncInput(aReq);
+    destroyEncOut(aReq);
+    av_free(aReq);
+    return 0;
+}
+
+static void hevc_init_params(AVCodecContext *avctx, HevcEncParam* params) {
+    params->width = avctx->width;
+    params->height = avctx->height;
+    params->ip = avctx->gop_size;
+    params->bitrate = avctx->bit_rate;
+    params->qp = avctx->global_quality;
+    if (params->qp > 51) {
+        // Constant bitrate mode
+        av_log(NULL, AV_LOG_INFO, " * Encoding in constant bitrate mode***\n");
+    } else {
+        // Fixed QP mode
+        av_log(NULL, AV_LOG_INFO, " * Encoding in fixed QP mode***\n");
+    }
+    params->reset = 0;
+    return;
+}
+static int isSupported(AVCodecContext *avctx) {
+    //Supported 1080p60, 8bits
+    // height and width should be a multiple of 4
+    // minimum resolution is 352x288
+    int width, height;
+    width = avctx->width;
+    height = avctx->height;
+    if (((width & 3) != 0) || ((height & 3) != 0)) {
+        av_log(NULL, AV_LOG_ERROR,"ERROR:height and width should be a multiple of 4\n");
+        return -1;
+    }
+    if ((width < 352 || height < 288) &&
+        (width < 288 || height < 352)) {
+        av_log(NULL, AV_LOG_ERROR,"ERROR: min resolution supported is 352x288\n");
+        return -1;
+    }
+    switch(avctx->pix_fmt) {
+    case AV_PIX_FMT_YUV420P:
+        avctx->bits_per_raw_sample = 8;
+        break;
+    default:
+        av_log(NULL, AV_LOG_ERROR,"ERROR: Pixel format not supported...\n");
+        return -1;
+        break;
+    }
+    return 0;
+}
+
+static int xlnx_hevc_encode_init(AVCodecContext *avctx) {
+    XlnxHevcContext* hevc;
+    HevcEncParam* params;
+    if (isSupported(avctx) < 0) {
+        return AVERROR(ENOSYS);
+    }
+    hevc = (XlnxHevcContext*)avctx->priv_data;
+    params = &hevc->params;
+    hevc_init_params(avctx, params);
+    hevc->rQ = createQueue(HEVC_ENC_OUT_Q_CAPACITY);
+    if (NULL == hevc->rQ) return -1;
+    hevc->hw = xlnx_hw_get(EHevcEnc, params);
+    if (NULL == hevc->hw) {
+        av_log(NULL, AV_LOG_ERROR,"ERROR: Failed to get EHevcEnc handle\n");
+        return AVERROR(EINVAL);
+    }
+    return 0;
+}
+
+static int xlnx_hevc_encode_frame(AVCodecContext *avctx, AVPacket *pkt, const AVFrame *pic, int *got_packet) {
+    XlnxHevcContext* hevc = (XlnxHevcContext*)avctx->priv_data;
+    int EntrySize = avctx->bits_per_raw_sample == 8 ? 1 : 2;
+    int ret;
+    int EntryCount = 0;
+    uint64_t EncodeBuffSize = 0;
+    int width, height, format;
+    EncodeRequest* req;
+    xlnx_hw_buf* outBuf;
+    void* nullptr;
+    HevcEncParam* params = &hevc->params;
+    if (NULL == pic) {
+        av_log(NULL, AV_LOG_DEBUG, "xlnx_hevc_encode_frame EOS Received\n");
+        if (params->reset == 0) {
+            params->reset = 1;
+        } else if (params->reset == 2) {
+            *got_packet = 0;
+            return 0;
+        }
+        width = avctx->width;
+        height = avctx->height;
+        format = avctx->pix_fmt;
+    } else {
+        width = pic->width;
+        height = pic->height;
+        format = pic->format;
+    }
+
+    switch(format) {
+    case AV_PIX_FMT_YUV420P:
+        EntryCount = 2;
+        break;
+    default:
+        av_log(NULL, AV_LOG_ERROR,"Error. Encoding unknown pixel format.\n");
+        return AVERROR_INVALIDDATA;
+        break;
+    }
+
+    EncodeBuffSize = EntryCount * width * height * EntrySize;
+
+    req = createEncReq(pic, EncodeBuffSize);
+    if (req == NULL) {
+        return -1;
+    }
+    outBuf = &req->oBufHdr;
+    nullptr = NULL;
+    ret = xlnx_hw_process_async(hevc->hw, &hevc->params, &req->iBufHdr, &outBuf, &nullptr);
+    if (ret < 0) {
+        *got_packet = 0;
+        destroyEncReq(req);
+        return AVERROR(EINVAL);
+    } else if (ret == 0) {
+        *got_packet = 0;
+        enqueue(hevc->rQ, (XItem)req);
+        return 0;
+    } else if (ret == 1) {
+        EncodeRequest* tmpReq = (EncodeRequest*)dequeue(hevc->rQ);
+        xlnx_data_buf* ptr = outBuf->ptr;
+        assert(tmpReq->out.pDATA == ptr->pDATA);
+        av_packet_from_data(pkt, ptr->pDATA, ptr->size);
+        pkt->pts = pkt->dts = (int64_t)tmpReq->pic->pts;
+        destroyEncInput(tmpReq);
+        av_free(tmpReq);
+        enqueue(hevc->rQ, (XItem)req);
+        *got_packet = 1;
+    } else if (ret == 2) {
+        *got_packet = 0;
+        params->reset = 2;
+    }
+    return 0;
+}
+
+static int xlnx_hevc_encode_close(AVCodecContext *avctx) {
+    XlnxHevcContext* hevc = (XlnxHevcContext*)avctx->priv_data;
+    xlnx_hw_release(hevc->hw);
+    if (hevc->rQ) {
+        while(!isEmpty(hevc->rQ)) {
+            EncodeRequest* req = (EncodeRequest*)dequeue(hevc->rQ);
+            destroyEncReq(req);
+        }
+        destroyQueue(hevc->rQ);
+    }
+    return 0;
+}
+
+AVCodec ff_xlnx_hevc_encoder = {
+    .name           = "xlnx_hevc_enc",
+    .long_name      = NULL_IF_CONFIG_SMALL("Xilinx High Efficiency Video Codec (H.265) Encoder"),
+    .type           = AVMEDIA_TYPE_VIDEO,
+    .id             = AV_CODEC_ID_HEVC,
+    .priv_data_size = sizeof(XlnxHevcContext),
+    .init           = xlnx_hevc_encode_init,
+    .encode2        = xlnx_hevc_encode_frame,
+    .close          = xlnx_hevc_encode_close,
+    .capabilities   = CODEC_CAP_DR1 | AV_CODEC_CAP_DELAY,
+    .pix_fmts       = (const enum AVPixelFormat[]) {
+                          AV_PIX_FMT_YUV420P, AV_PIX_FMT_NONE
+                      },
+};
+

new file libavcodec/xlnx_hw/Makefile

@@ +1 @@
+OBJS                            += xlnx_hw/xlnx_hw.o 
+OBJS                            += xlnx_hw/xlnx_hrm.o
+OBJS                            += xlnx_hw/xcl.o

new file libavcodec/xlnx_hw/common/inc/xcl.h

@@ +1 @@
+#ifndef XCL_H
+#define XCL_H
+
+#include <CL/opencl.h>
+#include "xlnx_types_internal.h"
+
+int xcl_GetPlatform(const char* platform, cl_platform_id* id);
+int xcl_CreateContextFromType(const cl_platform_id* id, cl_device_type aType,
+    cl_context* aContext);
+int xcl_ReleaseContext(cl_context* aContext);
+int xcl_GetDeviceList(const cl_platform_id* id, cl_device_type aType,
+    xlnx_device** aDeviceList, cl_uint* aNumDevices);
+int xcl_FreeDeviceList(xlnx_device* aDeviceList, cl_uint aNumDevices);
+int xcl_IsDeviceAvailable(cl_device_id aDevice);
+int xcl_ImportBinaryHeader(cl_context* aContext, cl_device_id* aDevice,
+    const char *xclbin, size_t aBinaryLength, cl_program* aProgram);
+int xcl_ImportBinary(cl_context* aContext, cl_device_id* aDevice,
+    const char *xclbin_file_name, cl_program* aProgram);
+int xcl_ReleaseProgram(cl_program* aProgram);
+int xcl_CreateKernels(cl_context* aContext, cl_device_id* aDevice,
+    cl_program* aProgram, const char** aKernelNames, uint32_t aNumKernels,
+        xlnx_kernel** aKernelList);
+int xcl_CreateKernels2(cl_context* aContext, cl_device_id* aDevice,
+    cl_program* aProgram, const KRNL_CONFIG* aKernelConfig, uint32_t aNumKernels,
+        xlnx_kernel** aKernelList);
+int xcl_FreeKernels(xlnx_kernel* aKernelList, uint32_t aNumKernels);
+
+#endif //XCL_H
+ No newline at end of file

new file libavcodec/xlnx_hw/common/inc/xlnx_hrm.h

@@ +1 @@
+#ifndef XLNX_HRM_H
+#define XLNX_HRM_H
+
+#include "../../xlnx_hw_types.h"
+#include "../inc/xlnx_types_internal.h"
+
+int xlnx_hrm_init(const char* aCfgFileName, xlnx_hw_id aHWID);
+int xlnx_hrm_deinit(void);
+xlnx_host* xlnx_hrm_get_host(xlnx_hw_id aHWID, xlnx_hw_params aParams);
+int xlnx_hrm_release_host(xlnx_host* aHost);
+
+#endif //XLNX_HRM_H

new file libavcodec/xlnx_hw/common/inc/xlnx_hw_defines.h

@@ +1 @@
+#ifndef XLNX_HW_DEFINES_H
+#define XLNX_HW_DEFINES_H
+
+#define RESET_UPPER_16 0x0000FFFF
+#define RESET_LOWER_16 0xFFFF0000
+
+#define FILE_PATH_LENGTH 1024
+#define LOAD_BINARY_FROM_HEADER 0
+#define VU9P_HEVC_ENC_2_K 1
+#define VU9P_HEVC_ENC_NUM_CH 2
+
+#define HEVC_KRNL_CH_Q_CAPACITY 3
+#define HEVC_ENC_IN_Q_CAPACITY 2
+#define HEVC_ENC_OUT_Q_CAPACITY (HEVC_ENC_IN_Q_CAPACITY + HEVC_KRNL_CH_Q_CAPACITY + 1)
+
+#define OCLCHECK(method, ... )                                                 \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        printf( # method " error '%d'\n", status);                             \
+        return -1;                                                             \
+    }                                                                          \
+} while (0)
+
+#define ENQUEUEBUF(method, dst)                                                \
+do {                                                                           \
+    status = method(cmdQ, dst, CL_TRUE, 0, size, iHBuf[i].pDATA, 0, NULL, NULL);\
+    if (CL_SUCCESS != status) {                                                \
+        printf(# method " failed, error %d\n", status);                        \
+        return -1;                                                             \
+    }                                                                          \
+} while (0)
+
+#endif //XLNX_HW_DEFINES_H
+ No newline at end of file

new file libavcodec/xlnx_hw/common/inc/xlnx_queue.h

@@ +1 @@
+#ifndef XLNX_QUEUE_H
+#define XLNX_QUEUE_H
+
+typedef void* XItem;
+typedef struct xlnx_queue xlnx_queue;
+
+xlnx_queue* createQueue(size_t capacity);
+void destroyQueue(xlnx_queue* queue);
+int isFull(xlnx_queue* queue);
+int isEmpty(xlnx_queue* queue);
+size_t getSize(xlnx_queue* queue);
+int enqueue(xlnx_queue* queue, XItem item);
+XItem dequeue(xlnx_queue* queue);
+XItem front(xlnx_queue* queue);
+XItem rear(xlnx_queue* queue);
+
+#endif // XLNX_QUEUE_H
+ No newline at end of file

new file libavcodec/xlnx_hw/common/inc/xlnx_ts_queue.h

@@ +1 @@
+#ifndef XLNX_TS_QUEUE_H
+#define XLNX_TS_QUEUE_H
+
+#include "xlnx_queue.h"
+
+typedef struct xlnx_ts_queue xlnx_ts_queue;
+
+xlnx_ts_queue* createTSQ(size_t capacity);
+void destroyTSQ(xlnx_ts_queue* queue);
+int isFullTSQ(xlnx_ts_queue* queue);
+int isEmptyTSQ(xlnx_ts_queue* queue);
+size_t getSizeTSQ(xlnx_ts_queue* queue);
+int enqueueTSQ(xlnx_ts_queue* queue, XItem item);
+XItem dequeueTSQ(xlnx_ts_queue* queue);
+XItem frontTSQ(xlnx_ts_queue* queue);
+XItem rearTSQ(xlnx_ts_queue* queue);
+
+#endif // XLNX_TS_QUEUE_H
+ No newline at end of file

new file libavcodec/xlnx_hw/common/inc/xlnx_types_internal.h

@@ +1 @@
+#ifndef XLNX_TYPES_INTERNAL_H
+#define XLNX_TYPES_INTERNAL_H
+
+#include <CL/opencl.h>
+#include "../../xlnx_hw_types.h"
+
+typedef struct {
+    char* devName;
+    cl_device_id device;
+    uint32_t majorV;
+    uint32_t minorV;
+    void* privData;
+} xlnx_device;
+
+typedef struct xlnx_kernel {
+    cl_context* context;
+    cl_command_queue cmdQ;
+    cl_kernel kernel;
+    U32 prfBank;
+    void* appData;
+    void* privData;
+} xlnx_kernel;
+
+typedef struct xlnx_host xlnx_host;
+typedef int (*Run)(xlnx_host* aHost, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf* aOutput);
+typedef int (*RunAsync)(xlnx_host* aHost, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf** aOutput, XContext* aContext);
+
+struct xlnx_host {
+    int ch;
+    Run process;
+    RunAsync processAsync;
+    int eosRcv;
+    void* appData;
+    void* privData;
+};
+
+typedef enum {
+    Y = 0,
+    U = 1,
+    V = 2
+} xlnx_col_fmt;
+
+typedef struct {
+    const char* krnl_name;
+    U32 prfBank;
+} KRNL_CONFIG;
+
+
+#endif //XLNX_TYPES_INTERNAL_H
+ No newline at end of file

new file libavcodec/xlnx_hw/common/src/xcl.c

@@ +1 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <unistd.h>
+#include <limits.h>
+#include <string.h>
+
+#include "../inc/xcl.h"
+#include "CL/cl_ext.h"
+
+// Wrap any OpenCL API calls that return error code(cl_int)
+// with the below macro to quickly check for an error
+#define OCL_CHECK(call)                                              \
+    do {                                                             \
+        cl_int err = call;                                           \
+        if (err != CL_SUCCESS) {                                     \
+            printf("Error from " #call ", error code is %d\n", err); \
+            return err;                                              \
+        }                                                            \
+    } while (0);
+
+static void* smalloc(size_t size) {
+    void* ptr;
+    ptr = malloc(size);
+    if (ptr == NULL) {
+        printf("Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+        return NULL;
+    }
+    return ptr;
+}
+
+static int load_file_to_memory(const char *filename, char **result) {
+    unsigned int size;
+    FILE* f;
+
+    f = fopen(filename, "rb");
+    if (f == NULL) {
+        *result = NULL;
+        printf("Error: Could not read file %s\n", filename);
+        return -1;
+    }
+
+    fseek(f, 0, SEEK_END);
+    size = ftell(f);
+    fseek(f, 0, SEEK_SET);
+
+    *result = (char *) smalloc(sizeof(char)*(size+1));
+
+    if (size != fread(*result, sizeof(char), size, f)) {
+        free(*result);
+        fclose(f);
+        printf("Error: read of kernel failed\n");
+        return -1;
+    }
+
+    fclose(f);
+    (*result)[size] = 0;
+
+    return size;
+}
+//
+// Get device version
+//
+static int getDeviceVersion(cl_device_id* aDevice, uint32_t* aMajor, uint32_t* aMinor)
+{
+    char versionString[512];
+    size_t size = 0;
+    uint32_t major = 0;
+    uint32_t minor = 0;
+    uint32_t state = 0;
+    cl_int err;
+    err = clGetDeviceInfo(*aDevice, CL_DEVICE_VERSION, 511, versionString, &size);
+    if ( err != CL_SUCCESS) {
+        printf("clGetDeviceInfo(CL_DEVICE_VERSION) Failed\n");
+        return -1;
+    }
+    size = strlen(versionString);
+    for (size_t i = 0; i < size; i++) {
+        if (versionString[i] == ' ') {
+            if (0 == state)
+                state++;
+            else
+                break;
+            continue;
+        }
+        if (versionString[i] == '.') {
+            state++;
+            continue;
+        }
+        if (state == 0) {
+            continue;
+        }
+        if (state == 1) {
+            major *= 10;
+            major += (versionString[i] - '0');
+            continue;
+        }
+        if (state == 2) {
+            minor *= 10;
+            minor += (versionString[i] - '0');
+            continue;
+        }
+        break;
+    }
+    *aMajor = major;
+    *aMinor = minor;
+    return -1;
+}
+
+int xcl_GetPlatform(const char* platform, cl_platform_id* id) {
+    int err = -1;
+    cl_uint num_platforms = 0;
+    cl_platform_id* platform_ids;
+    size_t i;
+    char* platform_name;
+
+    if ((platform == NULL) || (id == NULL)) return -1;
+
+    err = clGetPlatformIDs(0, NULL, &num_platforms);
+    if (err != CL_SUCCESS) {
+        printf("Error: no platforms available or OpenCL install broken\n");
+        return err;
+    }
+
+    platform_ids = (cl_platform_id *) malloc(num_platforms * sizeof(cl_platform_id));
+    if (platform_ids == NULL) {
+        printf("Error: Out of Memory %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+
+    err = clGetPlatformIDs(num_platforms, platform_ids, NULL);
+    if (err != CL_SUCCESS) {
+        printf("Error: Failed to find an OpenCL platform!\n");
+        free(platform_ids);
+        return err;
+    }
+
+    for(i = 0; i < num_platforms; i++) {
+        size_t platform_name_size;
+        err = clGetPlatformInfo(platform_ids[i], CL_PLATFORM_NAME,
+                                0, NULL, &platform_name_size);
+        if( err != CL_SUCCESS) {
+            printf("Error: Could not determine platform name!\n");
+            return err;
+        }
+
+        platform_name = (char*) malloc(sizeof(char)*platform_name_size);
+        if(platform_name == NULL) {
+            printf("Error: out of memory! %s at %d\n", __FILE__, __LINE__);
+            free(platform_ids);
+            return -1;
+        }
+
+        err = clGetPlatformInfo(platform_ids[i], CL_PLATFORM_NAME,
+                                platform_name_size, platform_name, NULL);
+        if(err != CL_SUCCESS) {
+            free(platform_ids);
+            free(platform_name);
+            printf("Error: could not determine platform name!\n");
+            return err;
+        }
+
+        if (!strcmp(platform_name, platform)) {
+            free(platform_name);
+            *id = platform_ids[i];
+            break;
+        }
+
+        free(platform_name);
+    }
+
+    free(platform_ids);
+
+    if (i == num_platforms) {
+        printf("Error: Failed to find Xilinx platform\n");
+        return -1;
+    }
+    return 0;
+}
+
+int xcl_CreateContextFromType(const cl_platform_id* id, cl_device_type aType, cl_context* aContext) {
+    int err;
+    cl_context_properties contextData[3];
+
+    if ((id == NULL) || (aContext == NULL)) return -1;
+    err = -1;
+    contextData[0] = CL_CONTEXT_PLATFORM;
+    contextData[1] = (cl_context_properties)*id;
+    contextData[2] = 0;
+    *aContext = clCreateContextFromType(contextData, aType, 0, 0, &err);
+    if (err != CL_SUCCESS) {
+        return -1;
+    }
+    return 0;
+}
+
+int xcl_ReleaseContext(cl_context* aContext) {
+    if (aContext == NULL) return -1;
+    OCL_CHECK(clReleaseContext(*aContext));
+    return 0;
+}
+
+int xcl_GetDeviceList(const cl_platform_id* id, cl_device_type aType, xlnx_device** aDeviceList, cl_uint* aNumDevices) {
+    int err;
+    int i;
+    cl_uint num_devices;
+    xlnx_device* devices;
+    cl_device_id* deviceIDs;
+    size_t device_name_size;
+
+    if ((id) == NULL || (aDeviceList == NULL) || (aNumDevices == NULL)) return -1;
+    err = -1;
+    num_devices = 0;
+
+    err = clGetDeviceIDs(*id, aType, 0, NULL, &num_devices);
+    if (err != CL_SUCCESS) {
+        printf("Error: no ACCELERATOR devices available\n");
+        return err;
+    }
+
+    deviceIDs = (cl_device_id *) malloc(sizeof(cl_device_id) * num_devices);
+    if (deviceIDs == NULL) {
+        printf("Error: Out of Memory %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+
+    err = clGetDeviceIDs(*id, aType, num_devices, deviceIDs, NULL);
+    if (err != CL_SUCCESS) {
+        free(deviceIDs);
+        printf("Error: could not get device ids\n");
+        return err;
+    }
+
+
+    devices = (xlnx_device *) calloc(num_devices, sizeof(xlnx_device));
+    if (devices == NULL) {
+        free(deviceIDs);
+        printf("Error: Out of Memory %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+
+    for (i = 0; i < num_devices; i++) {
+        err = clGetDeviceInfo(deviceIDs[i], CL_DEVICE_NAME,
+                              0, NULL, &device_name_size);
+        if(err != CL_SUCCESS) {
+            free(deviceIDs);
+            if (i>0) {
+                xcl_FreeDeviceList(devices, i);
+            }
+            printf("Error: could not determine device name\n");
+            return err;
+        }
+
+        devices[i].devName = (char*) malloc(sizeof(char)*device_name_size);
+        if(devices[i].devName == NULL) {
+            free(deviceIDs);
+            xcl_FreeDeviceList(devices, i+1);
+            printf("Error: Out of Memory! %s at %d\n", __FILE__, __LINE__);
+            return -1;
+        }
+
+        err = clGetDeviceInfo(deviceIDs[i], CL_DEVICE_NAME,
+                              device_name_size, devices[i].devName, NULL);
+        if(err != CL_SUCCESS) {
+            free(deviceIDs);
+            xcl_FreeDeviceList(devices, i+1);
+            printf("Error: could not determine device name\n");
+            return err;
+        }
+        getDeviceVersion(&deviceIDs[i], &devices[i].majorV, &devices[i].minorV);
+        printf("Found %s : version %d.%d\n", devices[i].devName, devices[i].majorV, devices[i].minorV);
+        devices[i].device = deviceIDs[i];
+    }
+    free(deviceIDs);
+    *aNumDevices = num_devices;
+    *aDeviceList = devices;
+    return 0;
+}
+
+int xcl_FreeDeviceList(xlnx_device* aDeviceList, cl_uint aNumDevices) {
+    if (aDeviceList == NULL) return -1;
+    for (int i=0; i<aNumDevices; i++) {
+        if (aDeviceList[i].devName)
+            free(aDeviceList[i].devName);
+    }
+    free(aDeviceList);
+    return 0;
+}
+
+int xcl_IsDeviceAvailable(cl_device_id aDevice) {
+    cl_bool avail = CL_TRUE;
+    OCL_CHECK(clGetDeviceInfo(aDevice, CL_DEVICE_AVAILABLE,
+                              sizeof(cl_bool), &avail, NULL));
+    if (avail == CL_FALSE) {
+        printf("Error: device busy!!\n");
+        return 0;
+    }
+    return 1;
+}
+
+int xcl_ImportBinaryHeader(cl_context* aContext, cl_device_id* aDevice,
+                        const char *xclbin, size_t aBinaryLength, cl_program* aProgram) {
+    int err;
+    cl_program program;
+    size_t len;
+    char buffer[2048];
+    if ((aContext == NULL) || (aDevice == NULL) || (xclbin == NULL ) ||
+        (aBinaryLength == 0) || (aProgram == NULL))
+        return -1;
+
+    err = -1;
+    program = clCreateProgramWithBinary(*aContext, 1,
+                                            aDevice, &aBinaryLength,
+                                            (const unsigned char **) &xclbin,
+                                            NULL, &err);
+    if ((!program) || (err!=CL_SUCCESS)) {
+        printf("Error: Failed to create compute program from binary %d!\n", err);
+        return -1;
+    }
+
+    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+    if (err != CL_SUCCESS) {
+        xcl_ReleaseProgram(&program);
+
+        clGetProgramBuildInfo(program, *aDevice, CL_PROGRAM_BUILD_LOG,
+                              sizeof(buffer), buffer, &len);
+        printf("%s\n", buffer);
+        printf("Error: Failed to build program executable!\n");
+        return -1;
+    }
+    *aProgram = program;
+    return 0;
+}
+
+int xcl_ImportBinary(cl_context* aContext, cl_device_id* aDevice,
+                            const char *xclbin_file_name, cl_program* aProgram) {
+    int err;
+    char* krnl_bin;
+    size_t krnl_size;
+    cl_program program;
+    size_t len;
+    char buffer[2048];
+
+    if ((aContext == NULL) || (aDevice == NULL) ||
+          (xclbin_file_name == NULL ) || (aProgram == NULL))
+        return -1;
+
+    if(access(xclbin_file_name, R_OK) != 0) {
+        printf("ERROR: %s xclbin not available please build\n", xclbin_file_name);
+        return -1;
+    }
+
+    krnl_size = load_file_to_memory(xclbin_file_name, &krnl_bin);
+    program = clCreateProgramWithBinary(*aContext, 1,
+                                        aDevice, &krnl_size,
+                                        (const unsigned char**) &krnl_bin,
+                                        NULL, &err);
+    if ((!program) || (err!=CL_SUCCESS)) {
+        free(krnl_bin);
+        printf("Error: Failed to create compute program from binary %d!\n",
+               err);
+        return -1;
+    }
+
+    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+    if (err != CL_SUCCESS) {
+        xcl_ReleaseProgram(&program);
+        free(krnl_bin);
+
+        clGetProgramBuildInfo(program, *aDevice, CL_PROGRAM_BUILD_LOG,
+                              sizeof(buffer), buffer, &len);
+        printf("%s\n", buffer);
+        printf("Error: Failed to build program executable!\n");
+        return -1;
+    }
+    *aProgram = program;
+    free(krnl_bin);
+
+    return 0;
+}
+
+int xcl_ReleaseProgram(cl_program* aProgram) {
+    if (aProgram == NULL) { return -1;}
+    OCL_CHECK(clReleaseProgram(*aProgram));
+    return 0;
+}
+
+int xcl_CreateKernels(cl_context* aContext, cl_device_id* aDevice,
+                      cl_program* aProgram, const char** aKernelNames, uint32_t aNumKernels,
+                      xlnx_kernel** aKernelList) {
+    int err = -1;
+    xlnx_kernel* kList;
+    int i;
+
+    if ((aContext == NULL) || (aDevice == NULL) ||
+          (aProgram == NULL ) || (aKernelNames == NULL) || (aNumKernels == 0)) {
+        return -1;
+    }
+    err = -1;
+    kList = (xlnx_kernel*)calloc(aNumKernels, sizeof(xlnx_kernel));
+    if (kList == NULL) {
+        printf("Error: Out of Memory! %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+    for (i = 0; i < aNumKernels; i++) {
+        kList[i].kernel = clCreateKernel(*aProgram, aKernelNames[i], &err);
+        if (err != CL_SUCCESS) {
+            xcl_FreeKernels(kList, i);
+            printf("Error: OpenCL unable to create kernel %s\n", aKernelNames[i]);
+            return -1;
+        }
+
+        if (!strcmp(aKernelNames[i], "krnl_ngcodec_hevc_0")) {
+            printf("%s bank %d\n", aKernelNames[i], XCL_MEM_DDR_BANK0);
+            kList[i].prfBank = XCL_MEM_DDR_BANK0;
+        } else if (!strcmp(aKernelNames[i], "krnl_ngcodec_hevc_1")) {
+            printf("%s bank %d\n", aKernelNames[i], XCL_MEM_DDR_BANK3);
+            kList[i].prfBank = XCL_MEM_DDR_BANK3;
+        }
+
+        kList[i].cmdQ = clCreateCommandQueue(*aContext, *aDevice, 0, &err);
+        if (err != CL_SUCCESS) {
+            xcl_FreeKernels(kList, i+1);
+            printf("Error: OpenCL unable to create CommandQueue\n");
+            return -1;
+        }
+        kList[i].context = aContext;
+    }
+    *aKernelList = kList;
+    return 0;
+}
+
+int xcl_CreateKernels2(cl_context* aContext, cl_device_id* aDevice,
+                      cl_program* aProgram, const KRNL_CONFIG* aKernelConfig, uint32_t aNumKernels,
+                      xlnx_kernel** aKernelList) {
+    int err;
+    xlnx_kernel* kList;
+    if ((aContext == NULL) || (aDevice == NULL) ||
+          (aProgram == NULL ) || (aKernelConfig == NULL) || (aNumKernels == 0)) {
+        return -1;
+    }
+    err = -1;
+    kList = (xlnx_kernel*)calloc(aNumKernels, sizeof(xlnx_kernel));
+    if (kList == NULL) {
+        printf("Error: Out of Memory! %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+    for (int i = 0; i < aNumKernels; i++) {
+        kList[i].kernel = clCreateKernel(*aProgram, aKernelConfig[i].krnl_name, &err);
+        if (err != CL_SUCCESS) {
+            xcl_FreeKernels(kList, i);
+            printf("Error: OpenCL unable to create kernel %s\n", aKernelConfig[i].krnl_name);
+            return -1;
+        }
+        kList[i].prfBank = aKernelConfig[i].prfBank;
+        kList[i].cmdQ = clCreateCommandQueue(*aContext, *aDevice, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
+        //kList[i].cmdQ = clCreateCommandQueue(*aContext, *aDevice, 0, &err);
+        if (err != CL_SUCCESS) {
+            xcl_FreeKernels(kList, i+1);
+            printf("Error: OpenCL unable to create CommandQueue\n");
+            return -1;
+        }
+        kList[i].context = aContext;
+    }
+    *aKernelList = kList;
+    return 0;
+}
+
+int xcl_FreeKernels(xlnx_kernel* aKernelList, uint32_t aNumKernels) {
+    if (aKernelList == NULL) return -1;
+    for (int i = 0; i < aNumKernels; i++) {
+        if (aKernelList[i].kernel)
+            clReleaseKernel(aKernelList[i].kernel);
+        if (aKernelList[i].cmdQ)
+            clReleaseCommandQueue(aKernelList[i].cmdQ);
+    }
+    free(aKernelList);
+    return 0;
+}

new file libavcodec/xlnx_hw/common/src/xlnx_hrm.c

@@ +1 @@
+#include <string.h>
+#include <stdio.h>
+#include <pthread.h>
+#include <assert.h>
+#include "../inc/xcl.h"
+#include "../inc/xlnx_hrm.h"
+#include "../inc/xlnx_hw_defines.h"
+#include "../../xlnx_hw_types.h"
+#include "../..//hevc/ku115_hevc_enc_config.h"
+#include "../../hevc/vu9p_hevc_enc_config.h"
+
+#include "../../hevc/enc_vu9p.h"
+#include "../../hevc/enc_vu9p_2c.h"
+
+// Wrap any OpenCL API calls that return error code(cl_int)
+// with the below macro to quickly check for an error
+#define XCL_CHECK(call)                                              \
+do {                                                                 \
+    int err = call;                                                  \
+    if (err != 0) {                                                  \
+        printf("Error from " #call ", error code is %d\n", err);     \
+        destroy_hrm();                                               \
+        return err;                                                  \
+    }                                                                \
+} while (0);
+
+static pthread_mutex_t * volatile atomic_hrm_lock = NULL;
+#define LOCK_HRM pthread_mutex_lock(atomic_hrm_lock)
+#define UNLOCK_HRM pthread_mutex_unlock(atomic_hrm_lock)
+
+static uint32_t xlnx_init_count = 0;
+
+typedef struct xlnx_krnl_info {
+    const char* kName;
+    xlnx_hw_id hwID;
+    uint32_t numCh;
+    // Can hold status for Max 32 channels
+    uint32_t chStatus;
+    uint32_t refCount;
+} xlnx_krnl_info;
+
+typedef struct xlnx_dev_info {
+    cl_program program;
+    uint32_t nKernels;
+    xlnx_kernel* kernelList;
+    xlnx_krnl_info* krnlInfo;
+} xlnx_dev_info;
+
+typedef enum xlnx_dev_type {
+    ENONE,
+    EKU115,
+    EVU9P
+} xlnx_dev_type;
+
+typedef struct xlnx_hrm {
+    cl_platform_id platform_id;
+    cl_context context;
+    cl_uint nDevices;
+    xlnx_device* deviceList;
+} xlnx_hrm;
+
+static xlnx_hrm* gHRM = NULL;
+
+static int clear_device_info(xlnx_dev_info* aDevInfo) {
+    if (aDevInfo == NULL) return -1;
+    if (aDevInfo->krnlInfo) free(aDevInfo->krnlInfo);
+    if (aDevInfo->kernelList) xcl_FreeKernels(aDevInfo->kernelList, aDevInfo->nKernels);
+    if (aDevInfo->program) xcl_ReleaseProgram(&aDevInfo->program);
+    free(aDevInfo);
+    return 0;
+}
+
+static int clear_device_list(void) {
+    int i;
+    xlnx_dev_info* temp;
+    if (gHRM->deviceList == NULL) return 0;
+    for (i = 0; i < gHRM->nDevices; i++) {
+        if (gHRM->deviceList[i].privData != NULL) {
+            temp = (xlnx_dev_info*)gHRM->deviceList[i].privData;
+            clear_device_info(temp);
+            gHRM->deviceList[i].privData = NULL;
+        }
+    }
+    xcl_FreeDeviceList(gHRM->deviceList, gHRM->nDevices);
+    return 0;
+}
+
+static int setup_KU115(xlnx_device* aDevice, xlnx_hw_id aHWID) {
+    //@ToDo parsing of config and loading of appropriate binary
+    //Assume the profile 1K1C/KU115 = kernel binary
+    int ret = -1;
+    xlnx_dev_info* devInfo;
+    int nK;
+    xlnx_krnl_info* info;
+    printf("setup_KU115\n");
+    devInfo = (xlnx_dev_info*) calloc(1, sizeof(xlnx_dev_info));
+    if (devInfo == NULL) {
+        printf("Error: OOM %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+    if (aHWID == EHevcEnc) {
+    // Header based loading
+#if LOAD_BINARY_FROM_HEADER
+        ret = xcl_ImportBinaryHeader(&gHRM->context, &aDevice->device, hevc_xclbin, hevc_xclbin_len, &devInfo->program);
+#else
+        ret = xcl_ImportBinary(&gHRM->context, &aDevice->device, KU115_HEVC_ENC_BINARY, &devInfo->program);
+#endif //LOAD_BINARY_FROM_HEADER
+        if (ret < 0) {
+            free(devInfo);
+            return -1;
+        }
+        nK =  sizeof(KU115_HEVC_KRNL_NAME)/sizeof(char*);
+        ret = xcl_CreateKernels(&gHRM->context, &aDevice->device, &devInfo->program, KU115_HEVC_KRNL_NAME, nK,
+                          &devInfo->kernelList);
+        if (ret < 0) {
+            clear_device_info(devInfo);
+            return -1;
+        }
+        devInfo->nKernels = nK;
+        devInfo->krnlInfo = (xlnx_krnl_info*) calloc(nK, sizeof(xlnx_krnl_info));
+        if (NULL == devInfo->krnlInfo) {
+            clear_device_info(devInfo);
+            return -1;
+        }
+        info = devInfo->krnlInfo;
+        for (int i = 0; i < nK; i++) {
+            info[i].kName = KU115_HEVC_KRNL_NAME[i];
+            info[i].hwID = EHevcEnc;
+            info[i].refCount = 0;
+        }
+    }
+    aDevice->privData = (void*)devInfo;
+    return 0;
+}
+
+static int setup_VU9P(xlnx_device* aDevice, xlnx_hw_id aHWID) {
+    int ret = -1;
+    xlnx_dev_info* devInfo;
+    char* xcl_fullpath;
+    unsigned char* xcl_path;
+    int nK;
+    int i;
+    xlnx_krnl_info* info;
+    uint32_t numCh = VU9P_HEVC_ENC_NUM_CH;
+    unsigned char* numch;
+
+    printf("Setup VU9P\n");
+    devInfo = (xlnx_dev_info*) calloc(1, sizeof(xlnx_dev_info));
+    if (devInfo == NULL) {
+        printf("Error: OOM %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+    if (aHWID == EHevcEnc) {
+    // Header based loading
+#if LOAD_BINARY_FROM_HEADER
+        ret = xcl_ImportBinaryHeader(&gHRM->context, &aDevice->device, hevc_enc_vu9p_2k2c_xclbin, hevc_enc_vu9p_2k2c_xclbin_len, &devInfo->program);
+#else
+        xcl_fullpath = malloc(FILE_PATH_LENGTH);
+        xcl_path = getenv("XLNX_XCLBIN_PATH");
+        if (NULL != xcl_path) {
+            snprintf(xcl_fullpath, FILE_PATH_LENGTH, "%s/%s", xcl_path, VU9P_HEVC_ENC_BINARY);
+        } else {
+            printf(" * XLNX_XCLBIN_PATH not found\n");
+            snprintf(xcl_fullpath, FILE_PATH_LENGTH, "%s", VU9P_HEVC_ENC_BINARY);
+        }
+        ret = xcl_ImportBinary(&gHRM->context, &aDevice->device, xcl_fullpath, &devInfo->program);
+        free(xcl_fullpath);
+#endif //LOAD_BINARY_FROM_HEADER
+        if (ret < 0) {
+            free(devInfo);
+            return -1;
+        }
+        nK =  sizeof(VU9P_HEVC_CONFIG)/sizeof(KRNL_CONFIG);
+        ret = xcl_CreateKernels2(&gHRM->context, &aDevice->device, &devInfo->program, VU9P_HEVC_CONFIG, nK,
+                          &devInfo->kernelList);
+        if (ret < 0) {
+            clear_device_info(devInfo);
+            return -1;
+        }
+        devInfo->nKernels = nK;
+        devInfo->krnlInfo = (xlnx_krnl_info*) calloc(nK, sizeof(xlnx_krnl_info));
+        if (NULL == devInfo->krnlInfo) {
+            clear_device_info(devInfo);
+            return -1;
+        }
+        info = devInfo->krnlInfo;
+        for (i = 0; i < nK; i++) {
+            // @Todo load info from some source
+            // for now, hardcore infor for hevc 2k2c
+            info[i].kName = VU9P_HEVC_CONFIG[i].krnl_name;
+            info[i].hwID = EHevcEnc;
+            numCh = VU9P_HEVC_ENC_NUM_CH;
+            numch = getenv("VU9P_HEVC_ENC_NUM_CH");
+            if (NULL != numch) {
+                numCh = atoi(numch);
+                printf(" * VU9P_HEVC_ENC_NUM_CH set to = %d\n", VU9P_HEVC_ENC_NUM_CH);
+            } else {
+                printf(" * VU9P_HEVC_ENC_NUM_CH not set. Default = %d\n", VU9P_HEVC_ENC_NUM_CH);
+            }
+            info[i].numCh = numCh;
+            info[i].chStatus = 0;
+            info[i].refCount = 0;
+            devInfo->kernelList[i].privData = (void*)&info[i];
+        }
+    }
+    aDevice->privData = (void*)devInfo;
+    return 0;
+}
+
+static xlnx_dev_type get_device_type(const xlnx_device* device) {
+    assert(device != NULL);
+    if (!strncmp(device->devName, KU115_BOARD, strlen(KU115_BOARD))) {
+        return EKU115;
+    } else if (!strncmp(device->devName, VU9P_BOARD, strlen(VU9P_BOARD))) {
+        return EVU9P;
+    }
+    return ENONE;
+}
+
+static int setup_devices(const char* aCfgFileName, xlnx_hw_id aHWID) {
+    //@Todo depending upon the config file or environment variable settings,
+    //create context and load binaries here or at a later stage
+    // for now just assume HEVC encoder config only.
+    int ret = -1;
+    xlnx_device* device = NULL;
+    /* seems CL_DEVICE_AVAILABLE not implemented in runtime
+    for (int i = 0; i < gHRM->nDevices; i++) {
+        if (xcl_IsDeviceAvailable(gHRM->deviceList[i].device)) {
+            device = &gHRM->deviceList[i];
+            break;
+        }
+    }*/
+    //So choose the first available device
+    // @TODO setup for all the available devices
+
+    device = &gHRM->deviceList[0];
+    if (NULL == device) {
+        return -1;
+    }
+    switch (get_device_type(device)) {
+      case EKU115:
+          ret = setup_KU115(device, aHWID);
+          break;
+      case EVU9P:
+          ret = setup_VU9P(device, aHWID);
+          break;
+      case ENONE:
+      default:
+          return -1;
+    }
+    return ret;
+}
+
+static void destroy_hrm(void) {
+    printf("destroy_hrm\n");
+    assert(gHRM != NULL);
+    if (gHRM->deviceList)
+        clear_device_list();
+    gHRM->deviceList = NULL;
+    if (gHRM->context)
+        xcl_ReleaseContext(&gHRM->context);
+    gHRM->context = 0;
+    free(gHRM);
+    gHRM = NULL;
+    printf("destroyed hrm\n");
+}
+
+static int create_hrm(const char* aCfgFileName, xlnx_hw_id aHWID) {
+    //@Todo parse aCfgFileName
+    // For now aCfgFileName is the URI of the Kernal binary file
+    printf("create_hrm\n");
+    assert(gHRM == NULL);
+    gHRM = calloc(1, sizeof(xlnx_hrm));
+    if (!gHRM) {
+        printf("Error: OOM %s at %d\n", __FILE__, __LINE__);
+        return -1;
+    }
+    XCL_CHECK(xcl_GetPlatform("Xilinx", &gHRM->platform_id));
+    XCL_CHECK(xcl_CreateContextFromType(&gHRM->platform_id, CL_DEVICE_TYPE_ACCELERATOR, &gHRM->context));
+    XCL_CHECK(xcl_GetDeviceList(&gHRM->platform_id, CL_DEVICE_TYPE_ACCELERATOR, &gHRM->deviceList, &gHRM->nDevices));
+    //If aHWID is known, then setup the devices here
+    if (ENone != aHWID) {
+        XCL_CHECK(setup_devices(aCfgFileName, aHWID));
+    }
+    printf("created hrm\n");
+    return 0;
+}
+
+static inline int init_hrm_mtx(void) {
+    if (!atomic_hrm_lock) {
+        int err;
+        pthread_mutex_t *tmp = malloc(sizeof(pthread_mutex_t));
+        if (!tmp)
+            return -1;
+        if ((err = pthread_mutex_init(tmp, NULL))) {
+            free(tmp);
+            return err;
+        }
+        if (__sync_val_compare_and_swap((void * volatile *)&atomic_hrm_lock, NULL, tmp)) {
+            pthread_mutex_destroy(tmp);
+            free(tmp);
+        }
+    }
+    return 0;
+}
+
+static inline int deinit_hrm_mtx(void) {
+    if (atomic_hrm_lock) {
+        pthread_mutex_t *tmp = atomic_hrm_lock;
+        UNLOCK_HRM;
+        pthread_mutex_destroy(atomic_hrm_lock);
+        atomic_hrm_lock = NULL;
+        free(tmp);
+    }
+    return 0;
+}
+
+int xlnx_hrm_init(const char* aCfgFileName, xlnx_hw_id aHWID) {
+    int ret = init_hrm_mtx();
+    if (ret < 0)
+        return ret;
+
+    LOCK_HRM;
+
+    if (xlnx_init_count == 0) {
+        ret = create_hrm(aCfgFileName, aHWID);
+        if (ret < 0) {
+            deinit_hrm_mtx();
+            return -1;
+        }
+    }
+    xlnx_init_count++;
+
+    UNLOCK_HRM;
+    return xlnx_init_count;
+}
+
+int xlnx_hrm_deinit() {
+    LOCK_HRM;
+    xlnx_init_count--;
+    if (xlnx_init_count == 0) {
+        destroy_hrm();
+    } else {
+      uint32_t count = xlnx_init_count;
+        UNLOCK_HRM;
+        return count;
+    }
+    deinit_hrm_mtx();
+    return 0;
+}
+
+static int allocChannel(int aNumCh, int* aChStatus) {
+    for(int i = 0; i < aNumCh; i++) {
+        if ((*aChStatus & (1 << i)) == 0) {
+            *aChStatus = *aChStatus | (1<<i);
+            return i;
+        }
+    }
+    return -1;
+}
+
+static int alloc_xKernel(xlnx_dev_info* devInfo, xlnx_hw_id aHWID, int* aIndex, int* aChannel) {
+    for (int i = 0; i < devInfo->nKernels; i++) {
+        xlnx_krnl_info* kInfo = &devInfo->krnlInfo[i];
+        if ((kInfo->hwID == aHWID) && (kInfo->refCount < kInfo->numCh)) {
+            int ch = allocChannel(kInfo->numCh, &kInfo->chStatus);
+            if (ch < 0) return ch;
+            kInfo->refCount++;
+            *aChannel = ch;
+            *aIndex = i;
+            return 0;
+        }
+    }
+    return -1;
+}
+
+static int free_xKernel(xlnx_krnl_info* aKrnlInfo, int aChannel) {
+    if (aKrnlInfo) {
+        aKrnlInfo->refCount--;
+        assert(aKrnlInfo->chStatus & (1 << aChannel));
+        aKrnlInfo->chStatus &= ~(1 << aChannel);
+    }
+    else
+        return -1;
+    return 0;
+}
+
+xlnx_host* xlnx_hrm_get_host(xlnx_hw_id aHWID, xlnx_hw_params aParams) {
+    // @TODO
+    // Device setup can be delayed to this point as well
+    // It will either be done as part of hrm init(if aHWID is known) or delayed till this point.
+    // hrm is being initialized from here in the current case, so it does not make a difference
+    xlnx_host* host = NULL;
+    xlnx_dev_info* devInfo;
+    int kIndex, ch, ret;
+
+    if (aHWID == ENONE) return host;
+    if (xlnx_hrm_init(NULL, aHWID) < 0) {
+        return host;
+    }
+    LOCK_HRM;
+    //@TODO check the whole devicelist
+    devInfo = (xlnx_dev_info*)gHRM->deviceList[0].privData;
+    if (devInfo == NULL) {
+        // Should not happen as hrm init i sdone with xlnx_hw_id
+        // Other wise setup_devices() should be done here
+        UNLOCK_HRM;
+        return host; 
+    }
+
+    ret = alloc_xKernel(devInfo, aHWID, &kIndex, &ch);
+    if (ret < 0) {
+        UNLOCK_HRM;
+        return host;
+    }
+    // ONLY HEVC ENc host available right now!!
+    //@Todo update when other hosts are available
+    if (aHWID == EHevcEnc) {
+        host = createHevcEnc2C(&devInfo->kernelList[kIndex], ch, aParams);
+        if (host == NULL) {
+            free_xKernel(&devInfo->krnlInfo[kIndex], ch);
+            UNLOCK_HRM;
+            printf("Error: Cannot allocate memory for xlnx_host %s at %d\n", __FILE__, __LINE__);
+            return host;
+        }
+    } else {
+        UNLOCK_HRM;
+        return host;
+    }
+
+    host->appData = (void*) &devInfo->krnlInfo[kIndex];
+    UNLOCK_HRM;
+    return host;
+}
+
+int xlnx_hrm_release_host(xlnx_host* aHost) {
+    xlnx_krnl_info* krnlInfo;
+    if (NULL == aHost) return -1;
+    krnlInfo = (xlnx_krnl_info*)aHost->appData;
+    if (free_xKernel(krnlInfo, aHost->ch) < 0) return -1;
+    destroyHevcEnc2C(aHost);
+    xlnx_hrm_deinit();
+    return 0;
+}

new file libavcodec/xlnx_hw/common/src/xlnx_queue.c

@@ +1 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <assert.h>
+#include "../inc/xlnx_queue.h"
+
+struct xlnx_queue {
+    int front, rear;
+    size_t size, capacity;
+    XItem* pItem;
+};
+
+xlnx_queue* createQueue(size_t capacity) {
+    xlnx_queue* queue = (xlnx_queue*) malloc(sizeof(xlnx_queue));
+    if (NULL == queue) {
+        printf("Error: Out of Memory! %s at %d\n", __FILE__, __LINE__);
+        return queue;
+    }
+    queue->capacity = capacity;
+    queue->front = queue->size = 0;
+    queue->rear = -1;
+    queue->pItem = (XItem*)malloc(queue->capacity * sizeof(XItem));
+    if (NULL == queue->pItem) {
+        printf("Error: Out of Memory! %s at %d\n", __FILE__, __LINE__);
+        free(queue);
+        return NULL;
+    }
+    return queue;
+}
+
+void destroyQueue(xlnx_queue* queue) {
+    if (NULL == queue) return;
+    free(queue->pItem);
+    free(queue);
+    return;
+}
+
+int isFull(xlnx_queue* queue) {
+    assert(queue != NULL);
+    return (queue->size == queue->capacity);
+}
+
+int isEmpty(xlnx_queue* queue) {
+    assert(queue != NULL);
+    return (queue->size == 0);
+}
+
+size_t getSize(xlnx_queue* queue) {
+    assert(queue != NULL);
+    return queue->size;
+}
+
+int enqueue(xlnx_queue* queue, XItem item) {
+    assert(queue != NULL);
+    if (isFull(queue))
+        return -1;
+    queue->rear = (queue->rear + 1)%queue->capacity;
+    queue->pItem[queue->rear] = item;
+    queue->size++;
+    return 0;
+}
+
+XItem dequeue(xlnx_queue* queue) {
+    XItem item;
+    assert(queue != NULL);
+    if (isEmpty(queue))
+        return NULL;
+    item = queue->pItem[queue->front];
+    queue->pItem[queue->front] = 0;
+    queue->front = (queue->front + 1)%queue->capacity;
+    queue->size = queue->size - 1;
+    return item;
+}
+
+XItem front(xlnx_queue* queue) {
+    assert(queue != NULL);
+    if (isEmpty(queue))
+        return NULL;
+    return queue->pItem[queue->front];
+}
+
+XItem rear(xlnx_queue* queue) {
+    assert(queue != NULL);
+    if (isEmpty(queue))
+        return NULL;
+    return queue->pItem[queue->rear];
+}
+ No newline at end of file

new file libavcodec/xlnx_hw/common/src/xlnx_ts_queue.c

@@ +1 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <pthread.h>
+
+#include "../inc/xlnx_ts_queue.h"
+
+#define LOCK_Q pthread_mutex_lock(&queue->lock)
+#define UNLOCK_Q pthread_mutex_unlock(&queue->lock)
+
+struct xlnx_ts_queue {
+    pthread_mutex_t lock;
+    xlnx_queue* q;
+};
+
+xlnx_ts_queue* createTSQ(size_t capacity) {
+    xlnx_queue* queue;
+    xlnx_ts_queue* q = calloc(1, sizeof(xlnx_ts_queue));
+    if (NULL == q) {
+        printf("Error: Out of Memory! %s at %d\n", __FILE__, __LINE__);
+        return q;
+    }
+    queue = createQueue(capacity);
+    if (NULL == queue) {
+        printf("Error: Out of Memory! %s at %d\n", __FILE__, __LINE__);
+        free(q);
+        return NULL;
+    }
+    q->q = queue;
+    if (pthread_mutex_init(&q->lock, NULL) < 0) {
+        destroyQueue(queue);
+        free(q);
+        return NULL;
+    }
+    return q;
+}
+
+void destroyTSQ(xlnx_ts_queue* queue) {
+    if (NULL == queue) return;
+    pthread_mutex_destroy(&queue->lock);
+    destroyQueue(queue->q);
+    free(queue);
+    return;
+}
+
+int isFullTSQ(xlnx_ts_queue* queue) {
+    int ret;
+    assert(queue != NULL);
+    LOCK_Q;
+    ret = isFull(queue->q);
+    UNLOCK_Q;
+    return ret;
+}
+
+int isEmptyTSQ(xlnx_ts_queue* queue) {
+    int ret;
+    assert(queue != NULL);
+    LOCK_Q;
+    ret = isEmpty(queue->q);
+    UNLOCK_Q;
+    return ret;
+}
+
+size_t getSizeTSQ(xlnx_ts_queue* queue) {
+    size_t size;
+    assert(queue != NULL);
+    LOCK_Q;
+    size = getSize(queue->q);
+    UNLOCK_Q;
+    return size;
+}
+
+int enqueueTSQ(xlnx_ts_queue* queue, XItem item) {
+    int ret;
+    assert(queue != NULL);
+    LOCK_Q;
+    ret = enqueue(queue->q, item);
+    UNLOCK_Q;
+    return ret;
+}
+
+XItem dequeueTSQ(xlnx_ts_queue* queue) {
+    XItem item;
+    assert(queue != NULL);
+    LOCK_Q;
+    item = dequeue(queue->q);
+    UNLOCK_Q;
+    return item;
+}
+
+XItem frontTSQ(xlnx_ts_queue* queue) {
+    XItem item;
+    assert(queue != NULL);
+    LOCK_Q;
+    item = front(queue->q);
+    UNLOCK_Q;
+    return item;
+}
+
+XItem rearTSQ(xlnx_ts_queue* queue) {
+    XItem item;
+    assert(queue != NULL);
+    LOCK_Q;
+    item = rear(queue->q);
+    UNLOCK_Q;
+    return item;
+}
+ No newline at end of file

new file libavcodec/xlnx_hw/hevc/Makefile

@@ +1 @@
+OBJS                            += xlnx_hw/hevc/enc_vu9p.o

new file libavcodec/xlnx_hw/hevc/enc_vu9p.c

@@ +1 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <unistd.h>
+#include <limits.h>
+#include <string.h>
+#include <assert.h>
+#include <CL/opencl.h>
+#include "enc_vu9p.h"
+#include "../common/inc/xlnx_hw_defines.h"
+
+//Macros
+
+#define OCLCHECKI(method, ... )                                                \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        printf( # method " error '%d'\n", status);                             \
+        KDD_deinit(aKDD);                                                      \
+        return -1;                                                             \
+    }                                                                          \
+} while (0)
+
+#define CREATEBUF(out, flags, size)                                            \
+do {                                                                           \
+    out = clCreateBuffer(*context, flags, size, NULL, &status);                \
+    if (status != CL_SUCCESS) {                                                \
+        printf("Could not create OpenCL buffer\n");                            \
+        KDD_deinit(aKDD);                                                      \
+        return -1;                                                             \
+    }                                                                          \
+} while (0)
+
+typedef struct KernelDeviceData {
+    // CSR3/CSR4/CSR5
+    cl_mem  InCur[3];
+    // CSR6/CSR7/CSR8
+    cl_mem  inNxt[3];
+    cl_mem  CSR12_RECON_BASE_V;
+    cl_mem  CSR13_BITSTREAM_BASE_V;
+    cl_mem  CSR14_BITSTREAM_LEN_V;
+    U32 CurImgSize[3];
+    U32 EncodedFrameSize;
+    U32 EncodedFrameLengthSize;
+    HevcEncParam* params;
+} KernelDeviceData;
+typedef KernelDeviceData KDD;
+
+typedef struct HevcEnc {
+    //KHD khd;
+    KDD kdd;
+    HevcEncParam params;
+    xlnx_kernel* xKernel;
+    U32 frameNumber;
+} HevcEnc;
+
+static int KDD_deinit(KDD* aKDD) {
+    if (aKDD->InCur[Y]) clReleaseMemObject(aKDD->InCur[Y]);
+    if (aKDD->InCur[U]) clReleaseMemObject(aKDD->InCur[U]);
+    if (aKDD->InCur[V]) clReleaseMemObject(aKDD->InCur[V]);
+    if (aKDD->inNxt[Y]) clReleaseMemObject(aKDD->inNxt[Y]);
+    if (aKDD->inNxt[U]) clReleaseMemObject(aKDD->inNxt[U]);
+    if (aKDD->inNxt[V]) clReleaseMemObject(aKDD->inNxt[V]);
+    if (aKDD->CSR12_RECON_BASE_V) clReleaseMemObject(aKDD->CSR12_RECON_BASE_V);
+    if (aKDD->CSR13_BITSTREAM_BASE_V) clReleaseMemObject(aKDD->CSR13_BITSTREAM_BASE_V);
+    if (aKDD->CSR14_BITSTREAM_LEN_V) clReleaseMemObject(aKDD->CSR14_BITSTREAM_LEN_V);
+    memset(aKDD, 0, sizeof(KDD));
+    return 0;
+}
+
+static int KDD_init(KDD* aKDD, HevcEnc* aEnc) {
+    cl_int status;
+    U32 width, height;
+    cl_kernel kernel;
+    cl_context* context;
+    int arg = 0;
+    if ((aKDD == NULL) || (aEnc == NULL)) return -1;
+    status = -1;
+    aKDD->params = &aEnc->params;
+    width = aKDD->params->width;
+    height = aKDD->params->height;
+    aKDD->CurImgSize[Y] = width * height;
+    aKDD->CurImgSize[U] = (width * height) >> 2;
+    aKDD->CurImgSize[V] = (width * height) >> 2;
+
+    // Encoded Frame data size - hard code to allocate host memory as this varies!
+    aKDD->EncodedFrameSize = (width * height * 3) >> 1;
+    aKDD->EncodedFrameLengthSize = 1;
+
+    // Setup the CSR registers
+    kernel = aEnc->xKernel->kernel;
+    context = aEnc->xKernel->context;
+    CREATEBUF(aKDD->InCur[Y], CL_MEM_READ_ONLY,  aKDD->CurImgSize[Y] * sizeof(U8));
+    CREATEBUF(aKDD->InCur[U], CL_MEM_READ_ONLY,  aKDD->CurImgSize[U] * sizeof(U8));
+    CREATEBUF(aKDD->InCur[V], CL_MEM_READ_ONLY,  aKDD->CurImgSize[V] * sizeof(U8));
+    CREATEBUF(aKDD->inNxt[Y], CL_MEM_READ_ONLY,  aKDD->CurImgSize[Y] * sizeof(U8));
+    CREATEBUF(aKDD->inNxt[U], CL_MEM_READ_ONLY,  aKDD->CurImgSize[U] * sizeof(U8));
+    // Padding to make reference base pointer of reconstructed picture 128KB aligned
+    CREATEBUF(aKDD->inNxt[V], CL_MEM_READ_ONLY,  (aKDD->CurImgSize[V]+ 48*128*1024-1920*1080*3) * sizeof(U8));
+    // Reference frame requires at least 16MB space in DDR
+    CREATEBUF(aKDD->CSR12_RECON_BASE_V, CL_MEM_READ_WRITE,  16*1024*1024);
+    CREATEBUF(aKDD->CSR13_BITSTREAM_BASE_V, CL_MEM_WRITE_ONLY, aKDD->EncodedFrameSize * sizeof(U8));
+    CREATEBUF(aKDD->CSR14_BITSTREAM_LEN_V, CL_MEM_WRITE_ONLY, aKDD->EncodedFrameLengthSize * sizeof(U64));
+
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(U32), &aKDD->params->width);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(U32), &aKDD->params->height);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->InCur[Y]);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->InCur[U]);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->InCur[V]);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->inNxt[Y]);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->inNxt[U]);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->inNxt[V]);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(U32), &aKDD->params->bitrate);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(U32), &aKDD->params->qp);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(U32), &aKDD->params->ip);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->CSR12_RECON_BASE_V);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->CSR13_BITSTREAM_BASE_V);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(cl_mem),       &aKDD->CSR14_BITSTREAM_LEN_V);
+    OCLCHECKI(clSetKernelArg, kernel, arg++, sizeof(U32), &aKDD->params->reset);
+    return 0;
+}
+
+static int hevc_encode(xlnx_host* aHost, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf* aOutput) {
+    cl_command_queue cmdQ;
+    cl_kernel kernel;
+    xlnx_data_buf* iHBuf;
+    xlnx_data_buf* oBuf;
+    size_t global_work_size_1d[1];
+    size_t local_work_size_1d[1];
+    int size;
+    cl_int status = -1;
+    HevcEnc* enc = (HevcEnc*)aHost->privData;
+    KDD kdd = enc->kdd;
+    HevcEncParam* params = (HevcEncParam*) aParams;
+    xlnx_kernel* xKernel = enc->xKernel;
+    U32 frameNumber = enc->frameNumber;
+
+    if ((frameNumber == 0) && (aInput->ptr == NULL)) { return -1;}
+
+    cmdQ = xKernel->cmdQ;
+    kernel = xKernel->kernel;
+
+    assert(aInput->type == ERAW);
+    assert(aOutput->type == EDATA);
+    iHBuf = aInput->ptr;
+    oBuf = aOutput->ptr;
+
+    global_work_size_1d[0] = 1;
+    local_work_size_1d[0] = 1;
+    size = 0;
+
+    if (frameNumber == 0) {
+        for (int i = 0; i < 3; i++) {
+            size = kdd.CurImgSize[i] * sizeof(U8);
+            ENQUEUEBUF(clEnqueueWriteBuffer, kdd.InCur[i]);
+        }
+        enc->frameNumber++;
+        oBuf->size = 0;
+        return 0;
+    } else if (frameNumber == 1) {
+        if (aInput->ptr != NULL) {
+            for (int i = 0; i < 3; i++) {
+                size = kdd.CurImgSize[i] * sizeof(U8);
+                ENQUEUEBUF(clEnqueueWriteBuffer, kdd.inNxt[i]);
+            }
+        } else {
+            OCLCHECK(clSetKernelArg, kernel, 14, sizeof(U32), &params->reset);
+        }
+        // av_log(NULL, AV_LOG_INFO," * Loading frame %d\n", frameNumber);
+        enc->frameNumber++;
+    } else {
+        // Ping-pong buffers used; switch cur and nxt buffer pointer for kernel
+        if (frameNumber%2 == 0) {
+            OCLCHECK(clSetKernelArg, kernel, 2, sizeof(cl_mem), &kdd.inNxt[Y]);
+            OCLCHECK(clSetKernelArg, kernel, 3, sizeof(cl_mem), &kdd.inNxt[U]);
+            OCLCHECK(clSetKernelArg, kernel, 4, sizeof(cl_mem), &kdd.inNxt[V]);
+            OCLCHECK(clSetKernelArg, kernel, 5, sizeof(cl_mem), &kdd.InCur[Y]);
+            OCLCHECK(clSetKernelArg, kernel, 6, sizeof(cl_mem), &kdd.InCur[U]);
+            OCLCHECK(clSetKernelArg, kernel, 7, sizeof(cl_mem), &kdd.InCur[V]);
+            if (aInput->ptr != NULL) {
+                for (int i = 0; i < 3; i++) {
+                    size = kdd.CurImgSize[i] * sizeof(U8);
+                    ENQUEUEBUF(clEnqueueWriteBuffer, kdd.InCur[i]);
+                }
+            } else {
+                OCLCHECK(clSetKernelArg, kernel, 14, sizeof(U32), &params->reset);
+            }
+        } else {
+            OCLCHECK(clSetKernelArg, kernel, 2, sizeof(cl_mem), &kdd.InCur[Y]);
+            OCLCHECK(clSetKernelArg, kernel, 3, sizeof(cl_mem), &kdd.InCur[U]);
+            OCLCHECK(clSetKernelArg, kernel, 4, sizeof(cl_mem), &kdd.InCur[V]);
+            OCLCHECK(clSetKernelArg, kernel, 5, sizeof(cl_mem), &kdd.inNxt[Y]);
+            OCLCHECK(clSetKernelArg, kernel, 6, sizeof(cl_mem), &kdd.inNxt[U]);
+            OCLCHECK(clSetKernelArg, kernel, 7, sizeof(cl_mem), &kdd.inNxt[V]);
+            if (aInput->ptr != NULL) {
+                for (int i = 0; i < 3; i++) {
+                    size = kdd.CurImgSize[i] * sizeof(U8);
+                    ENQUEUEBUF(clEnqueueWriteBuffer, kdd.inNxt[i]);
+                }
+            } else {
+                OCLCHECK(clSetKernelArg, kernel, 14, sizeof(U32), &params->reset);
+            }
+        }
+        enc->frameNumber++;
+    }
+    OCLCHECK(clEnqueueNDRangeKernel, cmdQ, kernel, 1, NULL, global_work_size_1d, local_work_size_1d, 0, NULL, NULL);
+
+    //
+    // Wait for the hw accelerator to complete
+    //
+    clFinish(cmdQ);
+    //av_log(NULL, AV_LOG_INFO," * Finished the xlnx hw accelerator kernel\n");
+
+    //
+    // Read the result
+    // Get the length of the Encoded Frame first as this varies!
+    //
+    // av_log(NULL, AV_LOG_INFO," * Enqueuing Read Buffer - CSR14_BITSTREAM_LEN_V\n");
+    OCLCHECK(clEnqueueReadBuffer, cmdQ, kdd.CSR14_BITSTREAM_LEN_V, CL_TRUE, 0, kdd.EncodedFrameLengthSize * sizeof(U64), &oBuf->size, 0, NULL, NULL);
+    // av_log(NULL, AV_LOG_INFO," * Encoded Frame Length = %d\n", *dstSize);
+    //*dstSize = *dstSize * 4;
+    printf(" * Encoded Frame Length = %llu enc->frameNumber=%u\n", oBuf->size, enc->frameNumber);
+
+    // av_log(NULL, AV_LOG_INFO," * Enqueuing Read Buffer - CSR13_BITSTREAM_BASE_V\n");
+    OCLCHECK(clEnqueueReadBuffer, cmdQ, kdd.CSR13_BITSTREAM_BASE_V, CL_TRUE, 0, oBuf->size * sizeof(U8), oBuf->pDATA, 0, NULL, NULL);
+    return 0;
+}
+
+xlnx_host* createHevcEnc(xlnx_kernel* aKernel, xlnx_hw_params aParams) {
+    HevcEnc* enc;
+    xlnx_host* host;
+    int err;
+    if (aKernel == NULL) return NULL;
+    enc = (HevcEnc*)calloc(1, sizeof(HevcEnc));
+    if (enc == NULL) {
+        printf("Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+        return NULL;
+    }
+    memcpy(&enc->params, aParams, sizeof(HevcEncParam));
+    enc->xKernel = aKernel;
+    err = KDD_init(&enc->kdd, enc);
+    if (err < 0) {
+        printf("Error: createHevcEnc Failed\n");
+        free(enc);
+        return NULL;
+    }
+    enc->frameNumber = 0;
+    host = (xlnx_host*) calloc(1, sizeof(xlnx_host));
+    if (host == NULL) {
+        printf("Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+        return NULL;
+    }
+    host->privData = (void*) enc;
+    host->process = &hevc_encode;
+    return host;
+}
+
+int destroyHevcEnc(xlnx_host* aHost) {
+    HevcEnc* enc;
+    if (aHost == NULL) return -1;
+    enc = (HevcEnc*)aHost->privData;
+    if (enc == NULL) return -1;
+    KDD_deinit(&enc->kdd);
+    free(enc);
+    free(aHost);
+    return 0;
+}

new file libavcodec/xlnx_hw/hevc/enc_vu9p.h

@@ +1 @@
+#ifndef HEVC_ENC_VU9P_H
+#define HEVC_ENC_VU9P_H
+
+#include "../common/inc/xlnx_types_internal.h"
+
+xlnx_host* createHevcEnc(xlnx_kernel* aKernel, xlnx_hw_params aParams);
+int destroyHevcEnc(xlnx_host* aHost);
+
+#endif //HEVC_ENC_VU9P_H
+ No newline at end of file

new file libavcodec/xlnx_hw/hevc/enc_vu9p_2c.c

@@ +1 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <unistd.h>
+#include <limits.h>
+#include <string.h>
+#include <assert.h>
+#include <CL/opencl.h>
+#include "pthread.h"
+
+//@TODO implement own msg queue and remove libavutil dependency
+#include "libavutil/atomic.h"
+#include "libavutil/threadmessage.h"
+#include "libavutil/error.h"
+
+#include "../common/inc/xlnx_hw_defines.h"
+#include "../common/inc/xlnx_ts_queue.h"
+#include "enc_vu9p_2c.h"
+
+#define OCLCHECKI(method, ... )                                                \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        printf( # method " error '%d'\n", status);                             \
+        return -1;                                                             \
+    }                                                                          \
+} while (0)
+
+#define OCLCHECKI2(method, ... )                                               \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        printf( # method " error '%d'\n", status);                             \
+        KDD2_deinit(kdd, ch, refcount);                                        \
+        return -1;                                                             \
+    }                                                                          \
+} while (0)
+
+#define CREATEBUF2(out, flags, size)                                               \
+do {                                                                               \
+    out = clCreateBuffer(*context, flags, size, &ext, &status);                    \
+    if (status != CL_SUCCESS) {                                                    \
+        printf("Could not create OpenCL buffer\n");                                \
+        KDD2_deinit(aKDD, ch, aRefcount);                                          \
+        return -1;                                                                 \
+    }                                                                              \
+    OCLCHECKI2(clEnqueueMigrateMemObjects, cmdQ, 1, &out, migflags, 0, NULL, NULL);\
+} while (0)
+
+#define OCLCHECKCB(method, ... )                                               \
+do {                                                                           \
+    status = method(__VA_ARGS__);                                              \
+    if (status != CL_SUCCESS) {                                                \
+        printf( # method " error '%d'\n", status);                             \
+        break;                                                                 \
+    }                                                                          \
+} while (0)
+
+#define WAIT_IF_EMPTY(Q) while (isEmptyTSQ(Q)) usleep(1000)
+
+
+typedef struct HevcEnc2 HevcEnc2;
+
+typedef enum HevcEncState {
+    EInit,
+    ESendIp,
+    EEncode,
+    EReadOpSize,
+    ECopyOp
+} HevcEncState;
+
+typedef struct HevcKDCtx {
+    cl_mem* inD;
+    cl_mem* outD;
+    xlnx_hw_buf* outH;
+    void* appCtx;
+} HevcKDCtx;
+
+typedef struct HevcEncCtx {
+    HevcEnc2* enc;
+    HevcKDCtx kdCtx[VU9P_HEVC_ENC_NUM_CH];
+    HevcEncState state;
+    cl_event lEncode;
+} HevcEncCtx;
+
+typedef struct HevcEncParam2 {
+    U32 wh[2];
+    U32 br;
+    U32 qp;
+    U32 ip;
+    U32 reset;
+} HevcEncParam2;
+
+typedef struct KernelDeviceData2 {
+    // CSR3/CSR4/CSR5
+    // CSR6/CSR7/CSR8
+    cl_mem  InCur[VU9P_HEVC_ENC_NUM_CH][3];
+    // CSR9/CSR10/CSR11
+    cl_mem  inNxt[3];
+    xlnx_ts_queue* inQ[VU9P_HEVC_ENC_NUM_CH];
+    cl_mem  CSR15_RECON_BASE_V;
+    // CSR 16 & 18 + CSR 17 & 19
+    xlnx_ts_queue* outQ[VU9P_HEVC_ENC_NUM_CH];
+    U32 CurImgSize[VU9P_HEVC_ENC_NUM_CH][3];
+    U32 EncodedFrameSize[VU9P_HEVC_ENC_NUM_CH];
+    //CSR20
+    U32 isDualCh;
+    HevcEncParam2* params;
+} KernelDeviceData2;
+
+typedef KernelDeviceData2 KDD2;
+
+struct HevcEnc2 {
+    KDD2 kdd;
+    HevcEncParam2 params;
+    xlnx_kernel* xKernel;
+    U32 frameNumber[VU9P_HEVC_ENC_NUM_CH];
+    int refCount;
+    pthread_t encThread;
+    AVThreadMessageQueue* reqQ[VU9P_HEVC_ENC_NUM_CH];
+    AVThreadMessageQueue* rspQ[VU9P_HEVC_ENC_NUM_CH];
+    xlnx_ts_queue* ctxQ;
+    HevcEncCtx* curCtx;
+};
+
+typedef enum {
+    EProcess,
+    EEOS
+} HevcMsgType;
+
+typedef struct HevcEncReq {
+    HevcMsgType type;
+    xlnx_host* host;
+    xlnx_hw_params params;
+    xlnx_hw_buf* in;
+    xlnx_hw_buf* out;
+    void* appCtx;
+} HevcEncReq;
+
+typedef struct HevcEncRsp {
+    HevcMsgType type;
+    xlnx_hw_buf* out;
+    void* appCtx;
+} HevcEncRsp;
+
+static int KDD2_deinit(KDD2* aKDD, uint32_t aChannel,  int aRefcount);
+
+static int hevc_prepare(HevcEnc2* aEnc, HevcEncReq* aReq) {
+    cl_int status = -1;
+    int ch;
+    HevcEnc2* enc = aEnc;
+    KDD2* kdd = &enc->kdd;
+    int refcount = enc->refCount;
+    xlnx_kernel* xKernel = enc->xKernel;
+    cl_command_queue cmdQ;
+    cl_kernel kernel;
+    int size;
+    xlnx_data_buf* iHBuf;
+    xlnx_host* host = aReq->host;
+    xlnx_hw_buf* inHost = aReq->in;
+    HevcEncCtx* ctx;
+    HevcKDCtx* kdCtx;
+    int arg;
+    U32 frameNumber;
+    cl_mem* inFBuf = NULL;
+
+    iHBuf = inHost->ptr;
+    ch = host->ch;
+    frameNumber = enc->frameNumber[ch];
+    if ((frameNumber == 0) && (iHBuf == NULL)) { return -1;}
+    assert(inHost->type == ERAW);
+
+    cmdQ = xKernel->cmdQ;
+    kernel = xKernel->kernel;
+
+    size = 0;
+    if (ch == 0) {
+        if (frameNumber == 0) {
+            arg = 2;
+            inFBuf = (cl_mem*)dequeueTSQ(kdd->inQ[ch]);
+            for (int i = 0; i < 3; i++) {
+                size = kdd->CurImgSize[ch][i] * sizeof(U8);
+                ENQUEUEBUF(clEnqueueWriteBuffer, inFBuf[i]);
+                OCLCHECKI2(clSetKernelArg, kernel, arg, sizeof(cl_mem),&inFBuf[i]);
+                arg++;
+            }
+            enc->frameNumber[ch]++;
+            ctx = enc->curCtx = dequeueTSQ(enc->ctxQ);
+            ctx->state = ESendIp;
+            kdCtx = &ctx->kdCtx[ch];
+            kdCtx->inD = inFBuf;
+            kdCtx->outH = aReq->out;
+            kdCtx->appCtx = aReq->appCtx;
+            return 0;
+        } else if (frameNumber == 1) {
+            if (iHBuf != NULL) {
+                inFBuf = (cl_mem*)dequeueTSQ(kdd->inQ[ch]);
+                arg = 8;
+                for (int i = 0; i < 3; i++) {
+                    size = kdd->CurImgSize[ch][i] * sizeof(U8);
+                    ENQUEUEBUF(clEnqueueWriteBuffer, inFBuf[i]);
+                    OCLCHECKI2(clSetKernelArg, kernel, arg, sizeof(cl_mem), &inFBuf[i]);
+                    arg++;
+                }
+            }
+            enc->frameNumber[ch]++;
+            ctx = frontTSQ(enc->ctxQ);
+            kdCtx = &ctx->kdCtx[ch];
+            kdCtx->inD = inFBuf;
+            kdCtx->outH = aReq->out;
+            kdCtx->appCtx = aReq->appCtx;
+        } else {
+            WAIT_IF_EMPTY(enc->ctxQ);
+            ctx = enc->curCtx = dequeueTSQ(enc->ctxQ);
+            ctx->state = ESendIp;
+            kdCtx = &ctx->kdCtx[ch];
+            inFBuf = kdCtx->inD;
+            arg = 2;
+            for (int i = 0; i < 3; i++) {
+                OCLCHECKI2(clSetKernelArg, kernel, arg, sizeof(cl_mem),&inFBuf[i]);
+                arg++;
+            }
+            if (iHBuf != NULL) {
+                WAIT_IF_EMPTY(kdd->inQ[ch]);
+                inFBuf = (cl_mem*)dequeueTSQ(kdd->inQ[ch]);
+                arg = 8;
+                for (int i = 0; i < 3; i++) {
+                    size = kdd->CurImgSize[ch][i] * sizeof(U8);
+                    ENQUEUEBUF(clEnqueueWriteBuffer, inFBuf[i]);
+                    OCLCHECKI2(clSetKernelArg, kernel, arg, sizeof(cl_mem), &inFBuf[i]);
+                    arg++;
+                }
+            }
+            enc->frameNumber[ch]++;
+            WAIT_IF_EMPTY(enc->ctxQ);
+            ctx = frontTSQ(enc->ctxQ);
+            kdCtx = &ctx->kdCtx[ch];
+            kdCtx->inD = inFBuf;
+            kdCtx->outH = aReq->out;
+            kdCtx->appCtx = aReq->appCtx;
+        }
+    } else if (ch == 1) {
+        if (iHBuf != NULL) {
+            ctx = enc->curCtx;
+            assert(ctx->state == ESendIp);
+            kdCtx = &ctx->kdCtx[ch];
+            WAIT_IF_EMPTY(kdd->inQ[ch]);
+            inFBuf = (cl_mem*)dequeueTSQ(kdd->inQ[ch]);
+            arg = 5;
+            for (int i = 0; i < 3; i++) {
+                size = kdd->CurImgSize[ch][i] * sizeof(U8);
+                ENQUEUEBUF(clEnqueueWriteBuffer, inFBuf[i]);
+                OCLCHECKI2(clSetKernelArg, kernel, arg, sizeof(cl_mem), &inFBuf[i]);
+                arg++;
+            }
+            kdCtx = &ctx->kdCtx[ch];
+            kdCtx->inD = inFBuf;
+            kdCtx->outH = aReq->out;
+            kdCtx->appCtx = aReq->appCtx;
+        } else {
+            U32 wh = 0;
+            OCLCHECKI2(clSetKernelArg, kernel, 1, sizeof(U32),          &wh);
+            OCLCHECKI2(clSetKernelArg, kernel, 5, sizeof(cl_mem),       NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 6, sizeof(cl_mem),       NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 7, sizeof(cl_mem),       NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 17, sizeof(cl_mem),      NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 18, sizeof(cl_mem),      NULL);
+        }
+        enc->frameNumber[ch]++;
+    }
+    return 0;
+}
+
+static void CL_CALLBACK notify_callback(cl_event event, cl_int event_command_exec_status, void* user_data) {
+    cl_int status = -1;
+    HevcEncCtx* ctx;
+    HevcKDCtx* kdCtx;
+    HevcEnc2* enc;
+    KDD2* kdd;
+    int isDualCh;
+    cl_command_queue cmdQ;
+    cl_event eventRead;
+    xlnx_hw_buf* outH;
+    xlnx_data_buf* oHBuf;
+    cl_mem* oFBuf;
+    HevcEncRsp rspMsg;
+
+    ctx = (HevcEncCtx*)user_data;
+    kdCtx = &ctx->kdCtx[0];
+    outH = kdCtx->outH;
+    oHBuf = outH->ptr;
+    oFBuf = kdCtx->outD;
+    enc = ctx->enc;
+    cmdQ = enc->xKernel->cmdQ;
+    kdd = &enc->kdd;
+    isDualCh = kdd->isDualCh;
+
+    switch (ctx->state) {
+        case EEncode:
+            oHBuf->size = 0;
+            OCLCHECKCB(clEnqueueReadBuffer, cmdQ, oFBuf[1], CL_FALSE, 0, sizeof(U64), &oHBuf->size, 0, NULL, &eventRead);
+            enqueueTSQ(kdd->inQ[0], kdCtx->inD);
+            kdCtx->inD = NULL;
+            ctx->state = EReadOpSize;
+            if (isDualCh) {
+                kdCtx = &ctx->kdCtx[1];
+                oFBuf = kdCtx->outD;
+                outH = kdCtx->outH;
+                oHBuf = outH->ptr;
+                OCLCHECKCB(clEnqueueReadBuffer, cmdQ, oFBuf[1], CL_FALSE, 0, sizeof(U64), &oHBuf->size, 1, &eventRead, &ctx->lEncode);
+                enqueueTSQ(kdd->inQ[1], kdCtx->inD);
+                kdCtx->inD = NULL;
+                OCLCHECKCB(clSetEventCallback, ctx->lEncode, CL_COMPLETE, &notify_callback, ctx);
+            } else {
+                OCLCHECKCB(clSetEventCallback, eventRead, CL_COMPLETE, &notify_callback, ctx);
+            }
+            break;
+        case EReadOpSize:
+            OCLCHECKCB(clEnqueueReadBuffer, cmdQ, oFBuf[0], CL_FALSE, 0, oHBuf->size * sizeof(U8), oHBuf->pDATA, 0, NULL, &eventRead);
+            ctx->state = ECopyOp;
+            if (isDualCh) {
+                kdCtx = &ctx->kdCtx[1];
+                oFBuf = kdCtx->outD;
+                outH = kdCtx->outH;
+                oHBuf = outH->ptr;
+                OCLCHECKCB(clEnqueueReadBuffer, cmdQ, oFBuf[0], CL_FALSE, 0, oHBuf->size * sizeof(U8), oHBuf->pDATA, 1, &eventRead, &ctx->lEncode);
+                OCLCHECKCB(clSetEventCallback, ctx->lEncode, CL_COMPLETE, &notify_callback, ctx);
+            } else {
+                OCLCHECKCB(clSetEventCallback, eventRead, CL_COMPLETE, &notify_callback, ctx);
+            }
+            break;
+        case ECopyOp:
+            rspMsg.type = EProcess;
+            rspMsg.out = kdCtx->outH;
+            rspMsg.appCtx = kdCtx->appCtx;
+            av_thread_message_queue_send(enc->rspQ[0], &rspMsg, 0);
+            kdd = &enc->kdd;
+            enqueueTSQ(kdd->outQ[0], oFBuf);
+            kdCtx->outD = NULL;
+            kdCtx->outH = NULL;
+            if (isDualCh) {
+                kdCtx = &ctx->kdCtx[1];
+                rspMsg.out = kdCtx->outH;
+                rspMsg.appCtx = kdCtx->appCtx;
+                oFBuf = kdCtx->outD;
+                av_thread_message_queue_send(enc->rspQ[1], &rspMsg, 0);
+                enqueueTSQ(kdd->outQ[1], oFBuf);
+                kdCtx->outD = NULL;
+                kdCtx->outH = NULL;
+            }
+            ctx->state = EInit;
+            enqueueTSQ(enc->ctxQ, ctx);
+            break;
+        default:
+            printf("Error: Unknown STATE!!\n");
+            break;
+    }
+}
+
+static int hevc_encode(HevcEnc2* aEnc) {
+    cl_int status = -1;
+    HevcEnc2* enc;
+    HevcEncCtx* ctx;
+    xlnx_hw_buf* outH0;
+    xlnx_data_buf* oHBuf0;
+    //xlnx_data_buf* oHBuf1;
+    KDD2* kdd;
+    int isDualCh;
+    xlnx_kernel* xKernel;
+    cl_command_queue cmdQ;
+    cl_kernel kernel;
+    size_t global_work_size_1d[1];
+    size_t local_work_size_1d[1];
+    cl_event event;
+    HevcKDCtx* kdCtx;
+    cl_mem* oFBuf;
+    enc = aEnc;
+    ctx = enc->curCtx;
+    assert(ctx != NULL);
+    kdCtx = &ctx->kdCtx[0];
+    outH0 = kdCtx->outH;
+    assert(outH0->type == EDATA);
+    oHBuf0 = outH0->ptr;
+    kdd = &enc->kdd;
+    xKernel = enc->xKernel;
+    cmdQ = xKernel->cmdQ;
+    kernel = xKernel->kernel;
+    isDualCh = kdd->isDualCh;
+
+    WAIT_IF_EMPTY(kdd->outQ[0]);
+    oFBuf = dequeueTSQ(kdd->outQ[0]);
+    OCLCHECKI(clSetKernelArg, kernel, 15, sizeof(cl_mem), &oFBuf[0]);
+    OCLCHECKI(clSetKernelArg, kernel, 16, sizeof(cl_mem), &oFBuf[1]);
+    kdCtx->outD = oFBuf;
+    if(isDualCh) {
+        kdCtx = &ctx->kdCtx[1];
+        WAIT_IF_EMPTY(kdd->outQ[1]);
+        oFBuf = dequeueTSQ(kdd->outQ[1]);
+        OCLCHECKI(clSetKernelArg, kernel, 17, sizeof(cl_mem), &oFBuf[0]);
+        OCLCHECKI(clSetKernelArg, kernel, 18, sizeof(cl_mem), &oFBuf[1]);
+        kdCtx->outD = oFBuf;
+    }
+    global_work_size_1d[0] = 1;
+    local_work_size_1d[0] = 1;
+    if (ctx->lEncode != 0) {
+        OCLCHECK(clEnqueueNDRangeKernel, cmdQ, kernel, 1, NULL, global_work_size_1d, local_work_size_1d, 1, &ctx->lEncode, &event);
+    } else {
+        OCLCHECK(clEnqueueNDRangeKernel, cmdQ, kernel, 1, NULL, global_work_size_1d, local_work_size_1d, 0, NULL, &event);
+    }
+    ctx->state = EEncode;
+    OCLCHECK(clSetEventCallback, event, CL_COMPLETE, &notify_callback, ctx);
+    enc->curCtx = NULL;
+    WAIT_IF_EMPTY(enc->ctxQ);
+    ctx = frontTSQ(enc->ctxQ);
+    ctx->lEncode = event;
+    if (kdd->params->reset == 1) {
+        kdd->params->reset = 0;
+        OCLCHECKI(clSetKernelArg, kernel, 20, sizeof(U32), &kdd->params->reset);
+    }
+    return 0;
+}
+
+
+static void* process(void* aEnc) {
+    HevcEnc2* enc;
+    AVThreadMessageQueue* reqQ;
+    AVThreadMessageQueue* rspQ;
+    AVThreadMessageQueue* reqQ1;
+    AVThreadMessageQueue* rspQ1;
+    HevcEncReq reqMsg;
+    HevcEncRsp rspMsg;
+    HevcEncReq reqMsg1;
+    HevcEncRsp rspMsg1;
+    int ret, isDualCh;
+
+    printf("Thread started\n");
+    enc = (HevcEnc2*)aEnc;
+    reqQ = enc->reqQ[0];
+    rspQ = enc->rspQ[0];
+    rspMsg.type = EProcess;
+
+    reqQ1 = NULL;
+    rspQ1 = NULL;
+    rspMsg1.out = NULL;
+    rspMsg1.appCtx = NULL;
+    isDualCh = 0;
+
+    av_thread_message_queue_recv(reqQ, &reqMsg, 0);
+    ret = hevc_prepare(enc, &reqMsg);
+    if (ret < 0) {
+        printf("Error: Prepare failed, Exit Thread aEnc=%p\n", aEnc);
+        pthread_exit(NULL);
+    }
+    while(1) {
+        av_thread_message_queue_recv(reqQ, &reqMsg, 0);
+        ret = hevc_prepare(enc, &reqMsg);
+        if (ret < 0) break;
+
+        isDualCh = enc->kdd.isDualCh;
+        if (isDualCh) {
+            reqQ1 = enc->reqQ[1];
+            rspQ1 = enc->rspQ[1];
+            av_thread_message_queue_recv(reqQ1, &reqMsg1, 0);
+            if (reqMsg1.type == EEOS) {
+                while (!isFullTSQ(enc->ctxQ)) usleep(1000);
+                break;
+            }
+            ret = hevc_prepare(enc, &reqMsg1);
+            if (ret < 0) break;
+        }
+        ret = hevc_encode(aEnc);
+        if (ret < 0) break;
+        if (reqMsg.type == EEOS) {
+            while (!isFullTSQ(enc->ctxQ)) usleep(1000);
+            break;
+        }
+    }
+    rspMsg.type = EEOS;
+    av_thread_message_queue_send(rspQ, &rspMsg, 0);
+    av_thread_message_queue_recv(reqQ, &reqMsg, AV_THREAD_MESSAGE_NONBLOCK);
+    if (isDualCh) {
+        av_thread_message_queue_recv(reqQ1, &reqMsg1, 0);
+        assert(reqMsg1.type == EEOS);
+        rspMsg1.type = EEOS;
+        av_thread_message_queue_send(rspQ1, &rspMsg1, 0);
+    }
+    printf("Exit Thread aEnc=%p\n", aEnc);
+    pthread_exit(NULL);
+}
+
+int KDD2_deinit(KDD2* aKDD, uint32_t aChannel,  int aRefcount) {
+    int ch = aChannel;
+    cl_mem* mem;
+    xlnx_ts_queue* inQ = aKDD->inQ[ch];
+    xlnx_ts_queue* outQ = aKDD->outQ[ch];
+    if (inQ) {
+        while(!isEmptyTSQ(inQ)) {
+            mem = (cl_mem*)dequeueTSQ(inQ);
+            if (mem) {
+                for (int i = 0; i < 3; i++) {
+                    clReleaseMemObject(mem[i]);
+                }
+                free(mem);
+            }
+        }
+        destroyTSQ(inQ);
+        aKDD->inQ[ch] = NULL;
+    }
+    if (outQ) {
+        while(!isEmptyTSQ(outQ)) {
+            mem = (cl_mem*)dequeueTSQ(outQ);
+            if (mem) {
+                for (int i = 0; i < 2; i++) {
+                    clReleaseMemObject(mem[i]);
+                }
+                free(mem);
+            }
+        }
+        destroyTSQ(outQ);
+        aKDD->outQ[ch] = NULL;
+    }
+    for (int i = 0; i < 3; i++) {
+        if (aKDD->InCur[ch][i]) clReleaseMemObject(aKDD->InCur[ch][i]);
+        if ((ch == 0) && (aKDD->inNxt[i])) clReleaseMemObject(aKDD->inNxt[i]);
+        aKDD->CurImgSize[ch][i] = 0;
+    }
+    if ((aRefcount == 0) && aKDD->CSR15_RECON_BASE_V) clReleaseMemObject(aKDD->CSR15_RECON_BASE_V);
+    aKDD->EncodedFrameSize[ch] = 0;
+    aKDD->isDualCh = 0;
+    return 0;
+}
+
+static int KDD2_init(KDD2* aKDD, HevcEnc2* aEnc, uint32_t aChannel, int aRefcount) {
+    cl_int status = -1;
+    KDD2* kdd = aKDD;
+    int refcount = aRefcount;
+    int ch = aChannel;
+    U32 width, height;
+    cl_kernel kernel;
+    cl_context* context;
+    U32 ddrBank;
+    cl_mem_ext_ptr_t ext;
+    cl_mem* mem;
+    xlnx_ts_queue* inQ;
+    xlnx_ts_queue* outQ;
+    cl_command_queue cmdQ;
+    cl_mem_migration_flags  migflags;
+
+    if ((aKDD == NULL) || (aEnc == NULL)) return -1;
+
+    aKDD->params = &aEnc->params;
+    width = aKDD->params->wh[ch] >> 16;
+    height = aKDD->params->wh[ch] & RESET_UPPER_16;
+    aKDD->CurImgSize[ch][Y] = width * height;
+    aKDD->CurImgSize[ch][U] = (width * height) >> 2;
+    aKDD->CurImgSize[ch][V] = (width * height) >> 2;
+
+    // Encoded Frame data size - hard code to allocate host memory as this varies!
+    aKDD->EncodedFrameSize[ch] = (width * height * 3) >> 1;
+
+    // Setup the CSR registers
+    kernel = aEnc->xKernel->kernel;
+    context = aEnc->xKernel->context;
+    ddrBank = aEnc->xKernel->prfBank;
+    cmdQ = aEnc->xKernel->cmdQ;
+    ext.obj = NULL;
+    ext.param = 0;
+    ext.flags = ddrBank;
+    migflags = 0;
+
+    if (aRefcount == 1) {
+        CREATEBUF2(aKDD->CSR15_RECON_BASE_V, CL_MEM_READ_WRITE  | CL_MEM_EXT_PTR_XILINX | CL_MEM_HOST_NO_ACCESS,  32*1024*1024 * sizeof(U8));
+    } else {
+        aKDD->isDualCh = 1;
+    }
+    inQ = createTSQ(HEVC_KRNL_CH_Q_CAPACITY);
+    outQ = createTSQ(HEVC_KRNL_CH_Q_CAPACITY);
+    if ((NULL == aKDD->outQ) || (NULL == aKDD->inQ)) {
+        if (aKDD->inQ) free(aKDD->inQ);
+        return -1;
+    }
+    aKDD->inQ[ch] = inQ;
+    aKDD->outQ[ch] = outQ;
+    for (int j = 0; j < HEVC_KRNL_CH_Q_CAPACITY; j++) {
+        mem = calloc(3, sizeof(cl_mem));
+        if (NULL == mem) return -1;
+        CREATEBUF2(mem[Y], CL_MEM_READ_ONLY | CL_MEM_EXT_PTR_XILINX,  aKDD->CurImgSize[ch][Y] * sizeof(U8));
+        CREATEBUF2(mem[U], CL_MEM_READ_ONLY | CL_MEM_EXT_PTR_XILINX,  aKDD->CurImgSize[ch][U] * sizeof(U8));
+        CREATEBUF2(mem[V], CL_MEM_READ_ONLY | CL_MEM_EXT_PTR_XILINX,  aKDD->CurImgSize[ch][V] * sizeof(U8));
+        enqueueTSQ(inQ, mem);
+        mem = calloc(2, sizeof(cl_mem));
+        if (NULL == mem) return -1;
+        CREATEBUF2(mem[0], CL_MEM_WRITE_ONLY  | CL_MEM_EXT_PTR_XILINX, aKDD->EncodedFrameSize[ch] * sizeof(U8));
+        CREATEBUF2(mem[1], CL_MEM_WRITE_ONLY  | CL_MEM_EXT_PTR_XILINX, sizeof(U64));
+        enqueueTSQ(outQ, mem);
+    }
+
+    if (ch == 0) {
+        OCLCHECKI2(clSetKernelArg, kernel, 0, sizeof(U32),          &aKDD->params->wh[0]);
+        OCLCHECKI2(clSetKernelArg, kernel, 14, sizeof(cl_mem),       &aKDD->CSR15_RECON_BASE_V);
+        if (aRefcount == 1) {
+            U32 wh = 0;
+            OCLCHECKI2(clSetKernelArg, kernel, 1, sizeof(U32),          &wh);
+            OCLCHECKI2(clSetKernelArg, kernel, 5, sizeof(cl_mem),       NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 6, sizeof(cl_mem),       NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 7, sizeof(cl_mem),       NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 17, sizeof(cl_mem),      NULL);
+            OCLCHECKI2(clSetKernelArg, kernel, 18, sizeof(cl_mem),      NULL);
+        }
+    } else if (ch == 1) {
+        OCLCHECKI2(clSetKernelArg, kernel, 1, sizeof(U32),          &aKDD->params->wh[1]);
+    }
+
+    OCLCHECKI2(clSetKernelArg, kernel, 11, sizeof(U32),          &aKDD->params->br);
+    OCLCHECKI2(clSetKernelArg, kernel, 12, sizeof(U32),          &aKDD->params->qp);
+    OCLCHECKI2(clSetKernelArg, kernel, 13, sizeof(U32),          &aKDD->params->ip);
+    OCLCHECKI2(clSetKernelArg, kernel, 19, sizeof(U32), &aKDD->isDualCh);
+    OCLCHECKI2(clSetKernelArg, kernel, 20, sizeof(U32), &aKDD->params->reset);
+    return 0;
+}
+
+static int hevc_process(xlnx_host* aHost, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf** aOutput, XContext* aContext) {
+    HevcEncReq reqMsg;
+    HevcEncRsp rspMsg;
+    int ret;
+    HevcEnc2* enc;
+    int ch = aHost->ch;
+
+    reqMsg.host = aHost;
+    reqMsg.params = aParams;
+    reqMsg.in = aInput;
+    reqMsg.out = *aOutput;
+    reqMsg.appCtx = *aContext;
+    reqMsg.type = EProcess;
+    if ((aInput->ptr == NULL) && (aHost->eosRcv == 0)) {
+        reqMsg.type = EEOS;
+    }
+
+    enc = (HevcEnc2*)aHost->privData;
+
+    // Send a request to encode the frame
+    if (aHost->eosRcv == 0) {
+        av_thread_message_queue_send(enc->reqQ[ch], &reqMsg, 0);
+    }
+
+    // Check if there is a response to return from a previous invocation
+    if (aInput->ptr == NULL) {
+        aHost->eosRcv = 1;
+        ret = av_thread_message_queue_recv(enc->rspQ[ch], &rspMsg, 0);
+    } else {
+        ret = av_thread_message_queue_recv(enc->rspQ[ch], &rspMsg, AV_THREAD_MESSAGE_NONBLOCK);
+    }
+    if (ret == AVERROR(EAGAIN)) {
+        *aOutput = NULL;
+        return 0;
+    }
+    if (rspMsg.type == EEOS) {
+        *aOutput = NULL;
+        return 2;
+    }
+    *aOutput = rspMsg.out;
+    *aContext = rspMsg.appCtx;
+    return 1;
+}
+
+static int setEnc2CParams(HevcEncParam2* dParames, uint32_t aChannel, HevcEncParam* sParams) {
+    int ch = aChannel;
+    dParames->wh[ch] = (sParams->width << 16) + sParams->height;
+    if (ch == 0) {
+        dParames->qp = (dParames->qp & RESET_LOWER_16) + sParams->qp;
+        dParames->ip = (dParames->ip & RESET_LOWER_16) + sParams->ip;
+        dParames->br = (dParames->br & RESET_LOWER_16) + sParams->bitrate;
+    } else {
+        dParames->qp = (dParames->qp & RESET_UPPER_16) + (sParams->qp << 16);
+        dParames->ip = (dParames->ip & RESET_UPPER_16) + (sParams->ip << 16);
+        dParames->br = (dParames->br & RESET_UPPER_16) + (sParams->bitrate << 16);
+    }
+    dParames->reset = 1;
+    return 0;
+}
+
+xlnx_host* createHevcEnc2C(xlnx_kernel* aKernel, uint32_t aChannel, xlnx_hw_params aParams) {
+    HevcEnc2* enc;
+    int err;
+    xlnx_host* host;
+    xlnx_ts_queue* ctxQ;
+    HevcEncCtx* ctx;
+
+    if (aKernel == NULL) return NULL;
+    assert((aChannel == 0) || ((aChannel == 1) && (aKernel->appData != NULL)));
+
+    if (aChannel == 0) {
+        enc = (HevcEnc2*)calloc(1, sizeof(HevcEnc2));
+        ctxQ = createTSQ(HEVC_KRNL_CH_Q_CAPACITY);
+        if ((enc == NULL) || (ctxQ == NULL)) {
+            if (enc) free(enc);
+            printf("Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+            return NULL;
+        }
+        enc->refCount = 0;
+        for (int i = 0; i < HEVC_KRNL_CH_Q_CAPACITY; i++) {
+            ctx = calloc(1, sizeof(HevcEncCtx));
+            if (ctx == NULL) {
+                free(enc);
+                while(!isEmptyTSQ(enc->ctxQ)) {
+                    free(dequeueTSQ(enc->ctxQ));
+                }
+                destroyTSQ(ctxQ);
+                printf("Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+                return NULL;
+            }
+            ctx->enc = enc;
+            ctx->state = EInit;
+            ctx->lEncode = 0;
+            enqueueTSQ(ctxQ, ctx);
+        }
+        enc->ctxQ = ctxQ;
+    } else if (aChannel == 1) {
+        enc = (HevcEnc2*)aKernel->appData;
+    }
+
+    setEnc2CParams(&enc->params, aChannel, aParams);
+    enc->xKernel = aKernel;
+    enc->refCount++;
+    err = KDD2_init(&enc->kdd, enc, aChannel, enc->refCount);
+    if (err < 0) {
+        printf("Error: createHevcEnc Failed\n");
+        free(enc);
+        return NULL;
+    }
+    enc->frameNumber[aChannel] = 0;
+    host = (xlnx_host*) calloc(1, sizeof(xlnx_host));
+    if (host == NULL) {
+        printf("Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+        return NULL;
+    }
+    host->privData = (void*) enc;
+    host->ch = aChannel;
+    host->processAsync = &hevc_process;
+    host->eosRcv = 0;
+    // Create the message queues
+    printf("Creating request message queue ch[%d]\n", aChannel);
+    err = av_thread_message_queue_alloc(&enc->reqQ[aChannel],
+                                       HEVC_ENC_IN_Q_CAPACITY, sizeof(HevcEncReq));
+    printf("Creating response message queue ch[%d]\n", aChannel);
+    err |= av_thread_message_queue_alloc(&enc->rspQ[aChannel],
+                                       HEVC_ENC_OUT_Q_CAPACITY, sizeof(HevcEncRsp));
+
+    if (err < 0)
+    {
+        printf("Error: Cannot allocate memory %s at %d\n", __FILE__, __LINE__);
+        destroyHevcEnc2C(host);
+        return NULL;
+    }
+
+    if (aChannel == 0) {
+        printf("create Thread for enc=%p host=%p\n", enc, host);
+        aKernel->appData = (void*) enc;
+        if (pthread_create(&enc->encThread, NULL, process, enc) < 0) {
+            printf("Error: Thread Creation Failed!!\n");
+            destroyHevcEnc2C(host);
+            return NULL;
+        }
+    }
+    return host;
+}
+
+int destroyHevcEnc2C(xlnx_host* aHost) {
+    HevcEnc2* enc;
+    if (aHost == NULL) return -1;
+    enc = (HevcEnc2*)aHost->privData;
+    if (enc == NULL) return -1;
+    enc->refCount--;
+    KDD2_deinit(&enc->kdd, aHost->ch, enc->refCount);
+    if (enc->refCount == 0) {
+        // @Todo Terminate the thread for sure
+        pthread_join(enc->encThread, NULL);
+
+        // Free the message queues
+        for (int i = 0; i < VU9P_HEVC_ENC_NUM_CH; i++) {
+            av_thread_message_queue_free(&enc->reqQ[i]);
+            av_thread_message_queue_free(&enc->rspQ[i]);
+        }
+        if (enc->ctxQ) {
+            while(!isEmptyTSQ(enc->ctxQ)) {
+                free(dequeueTSQ(enc->ctxQ));
+            }
+            destroyTSQ(enc->ctxQ);
+        }
+        free(enc);
+    }
+    free(aHost);
+    return 0;
+}

new file libavcodec/xlnx_hw/hevc/enc_vu9p_2c.h

@@ +1 @@
+#ifndef HEVC_ENC_VU9P_2C_H
+#define HEVC_ENC_VU9P_2C_H
+
+#include "../common/inc/xlnx_types_internal.h"
+
+/*
+Restrictions for dual channel mode:
+ 1. Both the channels must have same framerate
+ 2. Parameters like height/width do not change once encoder is initialized
+ 3. Both the channels need to be setup (createHevcEnc2C), before the second process
+    request arrives for the first channel(0)
+*/
+xlnx_host* createHevcEnc2C(xlnx_kernel* aKernel, uint32_t aChannel, xlnx_hw_params aParams);
+int destroyHevcEnc2C(xlnx_host* aHost);
+
+#endif //HEVC_ENC_VU9P_2C_H
+ No newline at end of file

new file libavcodec/xlnx_hw/hevc/ku115_hevc_enc_config.h

@@ +1 @@
+#ifndef KU115_HEVC_ENC_CONFIG
+#define KU115_HEVC_ENC_CONFIG
+
+#if LOAD_BINARY_FROM_HEADER
+#include "hevc_enc_ku115_bin.h"
+#else
+static const char* KU115_HEVC_ENC_BINARY = "sdxl_hevc.xclbin";
+#endif //LOAD_BINARY_FROM_HEADER
+// KU115
+static const char* KU115_BOARD = "xilinx:xil-accel-rd-ku115:4ddr-xpr";
+static const char* KU115_HEVC_KRNL_NAME[] = {"sdxl_hevc"};
+
+#endif //KU115_HEVC_ENC_CONFIG
+ No newline at end of file

new file libavcodec/xlnx_hw/hevc/vu9p_hevc_enc_config.h

@@ +1 @@
+#ifndef VU9P_HEVC_ENC_CONFIG
+#define VU9P_HEVC_ENC_CONFIG
+
+#include "../common/inc/xlnx_types_internal.h"
+
+#if LOAD_BINARY_FROM_HEADER
+#include "hevc_enc_vu9p_bin.h"
+#include "hevc_enc_vu9p_2k2c_bin.h"
+#else
+static const char* VU9P_HEVC_ENC_BINARY = "hevc_enc_vu9p_2k2c.aws.awsxclbin";
+#endif //LOAD_BINARY_FROM_HEADER
+
+// VU9P
+static const char* VU9P_BOARD = "xilinx:aws-vu9p-f1:4ddr-xpr-2pr";
+
+static KRNL_CONFIG VU9P_HEVC_CONFIG[] = {
+                                            {"krnl_ngcodec_hevc_f1_0", XCL_MEM_DDR_BANK3},
+#if 0
+                                            {"krnl_ngcodec_hevc_1", XCL_MEM_DDR_BANK3}
+#endif //VU9P_HEVC_ENC_2_K
+};
+
+
+#endif //VU9P_HEVC_ENC_CONFIG

new file libavcodec/xlnx_hw/xlnx_hw.c

@@ +1 @@
+#include "stdio.h"
+#include "stddef.h"
+#include "xlnx_hw.h"
+#include "common/inc/xlnx_hrm.h"
+#include "common/inc/xlnx_types_internal.h"
+
+int xlnx_hw_init(const char* aConfigFile, xlnx_hw_id aHWID) {
+    return xlnx_hrm_init(aConfigFile, aHWID);
+}
+
+int xlnx_hw_deinit(void) {
+    return xlnx_hrm_deinit();
+}
+
+xlnx_hw xlnx_hw_get(xlnx_hw_id aHW, xlnx_hw_params aParams) {
+    return (xlnx_hw)xlnx_hrm_get_host(aHW, aParams);
+}
+
+int xlnx_hw_process(xlnx_hw aHW, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf* aOutput) {
+    int err;
+    xlnx_host* host;
+    err = -1;
+    host = (xlnx_host*)aHW;
+    if (host->process) {
+        err = host->process(host, aParams, aInput, aOutput);
+    } else {
+        printf("Error: Not Implemented");
+    }
+    return err;
+}
+
+int xlnx_hw_process_async(xlnx_hw aHW, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf** aOutput, XContext* aContext) {
+    int err;
+    xlnx_host* host;
+    err = -1;
+    host = (xlnx_host*)aHW;
+    if (host->processAsync) {
+        err = host->processAsync(host, aParams, aInput, aOutput, aContext);
+    } else {
+        printf("Error: Not Implemented");
+    }
+    return err;
+}
+
+int xlnx_hw_release(xlnx_hw aHW) {
+    return xlnx_hrm_release_host((xlnx_host*) aHW);
+}
+
+

new file libavcodec/xlnx_hw/xlnx_hw.h

@@ +1 @@
+#ifndef XLNX_HW_H
+#define XLNX_HW_H
+#include "xlnx_hw_types.h"
+
+int xlnx_hw_init(const char* aConfigFile, xlnx_hw_id aHWID);
+int xlnx_hw_deinit(void);
+
+xlnx_hw xlnx_hw_get(xlnx_hw_id aHWID, xlnx_hw_params aParams);
+int xlnx_hw_process(xlnx_hw aHW, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf* aOutput);
+//@TODO Async API added for POC. Aync API with callback to be added later
+int xlnx_hw_process_async(xlnx_hw aHW, xlnx_hw_params aParams, xlnx_hw_buf* aInput, xlnx_hw_buf** aOutput, XContext* aContext);
+int xlnx_hw_release(xlnx_hw aHW);
+
+#endif // XLNX_HW_H

new file libavcodec/xlnx_hw/xlnx_hw_types.h

@@ +1 @@
+#ifndef XLNX_HW_TYPES_H
+#define XLNX_HW_TYPES_H
+
+//@TODO define informative error codes
+#define XLNX_MAX_COMPONENTS 3
+typedef unsigned char       U8;
+typedef unsigned int       U32;
+typedef unsigned long long U64;
+
+
+typedef void* xlnx_hw;
+typedef void* xlnx_hw_params;
+typedef void* XBUF;
+typedef void* XDATA;
+typedef void* XContext;
+
+typedef enum {
+    EDATA,
+    ERAW
+} xlnx_buf_type;
+
+typedef struct xlnx_hw_buf {
+    xlnx_buf_type type;
+    XBUF ptr;
+    void* user_priv;
+    void* priv;
+} xlnx_hw_buf;
+
+//EDATA
+typedef struct xlnx_data {
+    XDATA pDATA;
+    U64 offset;
+    U64 size;
+} xlnx_data_buf;
+
+//ERAW (YUV, RGB etc)
+typedef struct {
+    xlnx_data_buf data[XLNX_MAX_COMPONENTS];
+} xlnx_cmp_buf;
+
+typedef enum xlnx_hw_id {
+    ENone,
+    EHevcEnc
+} xlnx_hw_id;
+
+typedef struct HevcEncParam {
+    U32 width;
+    U32 height;
+    U32 bitrate;
+    U32 qp; // quantization parameter
+    U32 ip; // intra period
+    U32 reset;
+} HevcEncParam;
+
+#endif //XLNX_HW_TYPES_H
+ No newline at end of file

libavutil/opencl.c

@@ -88 +88 @@
 
 static OpenclContext opencl_ctx = {&openclutils_class};
 
-static const cl_device_type device_type[] = {CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_CPU};
+static const cl_device_type device_type[] = {CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_ACCELERATOR};
 
 typedef struct {
     int err_code;