Upload 72 files

.gitattributes

@@ -34,3 +34,16 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
 *.zst filter=lfs diff=lfs merge=lfs -text
 *tfevents* filter=lfs diff=lfs merge=lfs -text
 facewrapper/libs/libttvfaceengine6.so filter=lfs diff=lfs merge=lfs -text
+openvino/libgna.so filter=lfs diff=lfs merge=lfs -text
+openvino/libgna.so.2 filter=lfs diff=lfs merge=lfs -text
+openvino/libgna.so.3.0.0.1455 filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino_gapi_preproc.so filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino_intel_cpu_plugin.so filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino_intel_gna_plugin.so filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino_intel_hddl_plugin.so filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino_intel_myriad_plugin.so filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino_onnx_frontend.so filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino_tensorflow_fe.so filter=lfs diff=lfs merge=lfs -text
+openvino/libopenvino.so filter=lfs diff=lfs merge=lfs -text
+openvino/pcie-ma2x8x.mvcmd filter=lfs diff=lfs merge=lfs -text
+openvino/usb-ma2x8x.mvcmd filter=lfs diff=lfs merge=lfs -text

openvino/libgna.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:22441d86dca92b00ae7fb9d315bcb1c6a8a213ac4fe86396489753ebe76f869e
+size 3120536

openvino/libgna.so.2

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:22441d86dca92b00ae7fb9d315bcb1c6a8a213ac4fe86396489753ebe76f869e
+size 3120536

openvino/libgna.so.3.0.0.1455

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:22441d86dca92b00ae7fb9d315bcb1c6a8a213ac4fe86396489753ebe76f869e
+size 3120536

openvino/libopenvino.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:fd216848c1ba78e62360c12c9684df0c160f6962f3d900e5918cc042b42b2b46
+size 13495416

openvino/libopenvino_gapi_preproc.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:3ac5ce0a8f3acefb41e8aa8161f78035dafff25c4b8c3485ebc541573b2b15f0
+size 1312920

openvino/libopenvino_intel_cpu_plugin.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:afe05ada6d5b11495a21787fa6ab0162fc40f7a9ab97be78f7b7185126d15b18
+size 33299880

openvino/libopenvino_intel_gna_plugin.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:ef15b623e7f81788160c4056ccd5e887a8184affe381e84a906646ef36cae1ab
+size 4067016

openvino/libopenvino_intel_hddl_plugin.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:96362327fbc404e88583bdcd2a526ccbf4ca26d4ecdb8898234be7986d9b8b2b
+size 5894680

openvino/libopenvino_intel_myriad_plugin.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:e596436002565356b80400e0d7e50093d53d338f623b171f658de527477852de
+size 6120168

openvino/libopenvino_onnx_frontend.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:0770ed09d471b20bffcf4ef57ab1fb002db04c4404598bd5c52a4418a67f5441
+size 3781640

openvino/libopenvino_tensorflow_fe.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:c2dadbcd8ba32cec02873caf8dcc644d1d8856cdcd2978c603e5bac169e01bb9
+size 2723864

openvino/pcie-ma2x8x.mvcmd

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:f03146453508f2bcab1589907bccaa429b48db6123a7b8a428d6ce221d1fbb4d
+size 2099248

openvino/plugins.xml

@@ -0,0 +1,27 @@
+<ie>
+    <plugins>
+        <plugin name="AUTO" location="libopenvino_auto_plugin.so">
+            <properties>
+                <property key="MULTI_WORK_MODE_AS_AUTO" value="YES"/>
+            </properties>
+        </plugin>
+        <plugin name="BATCH" location="libopenvino_auto_batch_plugin.so">
+        </plugin>
+        <plugin name="CPU" location="libopenvino_intel_cpu_plugin.so">
+        </plugin>
+        <plugin name="GNA" location="libopenvino_intel_gna_plugin.so">
+        </plugin>
+        <plugin name="GPU" location="libopenvino_intel_gpu_plugin.so">
+        </plugin>
+        <plugin name="HETERO" location="libopenvino_hetero_plugin.so">
+        </plugin>
+        <plugin name="MULTI" location="libopenvino_auto_plugin.so">
+        </plugin>
+        <plugin name="MYRIAD" location="libopenvino_intel_myriad_plugin.so">
+        </plugin>
+        <plugin name="HDDL" location="libopenvino_intel_hddl_plugin.so">
+        </plugin>
+        <plugin name="VPUX" location="libopenvino_intel_vpux_plugin.so">
+        </plugin>
+    </plugins>
+</ie>

openvino/usb-ma2x8x.mvcmd

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:faf33388b88708177a358fcb4704eba04b1cf9e88d6a047f90c833d686140a2e
+size 2298632

openvino/vpu_custom_kernels/binarization.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:3e0de6082c7bacca2ff5ad131f0afc44304fc792a6d99e7829399eb61491a0ac
+size 19632

openvino/vpu_custom_kernels/binarization.cl

@@ -0,0 +1,67 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__kernel void binarization(
+    const __global half *__restrict src_data,
+    const __global half *__restrict input_low_high,
+    const __global half *__restrict dst_data,
+    int switch_out,
+    int input_low_high_size,
+    int W,
+    int H)
+{
+    __local half local_src[15 * 1024];
+    __local half local_dst[15 * 1024];
+
+    event_t e1 = async_work_group_copy(local_src, src_data + get_group_id(2) * W * H, W * H, 0);
+    wait_group_events(1, &e1);
+
+    int c = get_global_id(2);
+    int C = get_global_size(2);
+
+    half dst_low  = switch_out ? 1.h : -1.h;
+    half dst_high = switch_out ? -1.h : 1.h;
+
+    half s_ilow_ihigh = input_low_high_size == 1 ? input_low_high[0] : input_low_high[c];
+
+    for (int h = 0; h < H; h++) {
+
+        __local const half *__restrict addr_src = local_src + h * W;
+        __local half *__restrict addr_dst       = local_dst + h * W;
+
+#if 1
+        for (int w = 0; w < W / 8; w++) {
+
+            half8 h_src_val8 = (*((__local half8 *)addr_src + w));
+
+            short8 cond1;
+            cond1.s0 = (h_src_val8.s0 <= s_ilow_ihigh);
+            cond1.s1 = (h_src_val8.s1 <= s_ilow_ihigh);
+            cond1.s2 = (h_src_val8.s2 <= s_ilow_ihigh);
+            cond1.s3 = (h_src_val8.s3 <= s_ilow_ihigh);
+            cond1.s4 = (h_src_val8.s4 <= s_ilow_ihigh);
+            cond1.s5 = (h_src_val8.s5 <= s_ilow_ihigh);
+            cond1.s6 = (h_src_val8.s6 <= s_ilow_ihigh);
+            cond1.s7 = (h_src_val8.s7 <= s_ilow_ihigh);
+
+            cond1 = ~(cond1 - (short8)1);
+
+            short8 res = cond1 & as_short8((half8)dst_low) | ~cond1 & as_short8((half8)dst_high);
+
+            *((__local half8 *)addr_dst + w) = as_half8(res);
+        }
+#endif
+        for (int w = W & (~0x7); w < W; w++) {
+            addr_dst[w] = (addr_src[w] <= s_ilow_ihigh) ? dst_low : dst_high;
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy(dst_data + get_group_id(2) * W * H, local_dst, W * H, 0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/binary_convolution.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:12c349d6f73c233b158e1d67af31715c7b8bda79f191b1e759476e01e65bb64a
+size 10764

openvino/vpu_custom_kernels/binary_convolution.cl

@@ -0,0 +1,95 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+int extract_weights(uchar val, int bit) { return ((val >> bit) & 1); }
+
+__kernel void binary_convolution(
+    const __global half *restrict src_data,
+    const __global uchar *restrict weights_data,
+    __global half *restrict dst_data,
+    float pad_value,
+
+    int IW,
+    int IH,
+    int IC,
+
+    int DW,
+    int DH,
+
+    int GC,
+
+    int KW,
+    int KH,
+
+    int PW,
+    int PH,
+
+    int SW,
+    int SH)
+{
+    int ipad_value = ((pad_value > 0.f) ? 1 : 0);
+    int c          = get_global_id(2);
+    int y          = get_global_id(1);
+    int x          = get_global_id(0);
+
+    int OC = get_global_size(2);
+    int OH = get_global_size(1);
+    int OW = get_global_size(0);
+
+    int KD = 1;
+    int SD = 0;
+    int DD = 0;
+    int PD = 0;
+    int ID = 1;
+    int OD = 1;
+
+    int nbits = 8;
+
+    int g  = c % GC;
+    int oc = c / GC;
+    int oh = y;
+    int ow = x;
+
+    for (int od = 0; od < OD; od++) {
+        int oidx = g * OC / GC * OD * OH * OW + oc * OD * OH * OW + od * OH * OW + oh * OW + ow;
+
+        int res = 0;
+
+        for (int ic = 0; ic < IC / GC; ic++) {
+            for (int kd = 0; kd < KD; kd++) {
+                for (int kh = 0; kh < KH; kh++) {
+                    for (int kw = 0; kw < KW; kw++) {
+                        int widx = g * OC / GC * IC / GC * KD * KH * KW
+                                   + oc * IC / GC * KD * KH * KW + ic * KD * KH * KW + kd * KH * KW
+                                   + kh * KW + kw;
+
+                        int w = extract_weights(weights_data[widx / nbits], (widx % nbits));
+
+                        int s;
+
+                        int iw = ow * SW - PW + kw * DW;
+                        int ih = oh * SH - PH + kh * DH;
+                        int id = od * SD - PD + kd * DD;
+
+                        if (iw < 0 || iw >= (int)IW || ih < 0 || ih >= (int)IH || id < 0
+                            || id >= (int)ID) {
+                            s = ipad_value;
+                        } else {
+                            int iidx = g * IC / GC * ID * IH * IW + ic * ID * IH * IW + id * IH * IW
+                                       + ih * IW + iw;
+
+                            s = ((src_data[iidx] > 0.f) ? 1 : 0);
+                        }
+
+                        res += s ^ w;
+                    }
+                }
+            }
+        }
+
+        dst_data[oidx] = (half)(IC / GC * KD * KH * KW - 2 * res);
+    }
+}

openvino/vpu_custom_kernels/binary_convolution1x1.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:6deff31d62aa84c643fbeba77e7dcd4ae5d9b488c1c98e07fffeb58ff8e9b945
+size 76316

openvino/vpu_custom_kernels/binary_convolution1x1.cl

@@ -0,0 +1,117 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+ushort extract_weights(uchar val, int bit) { return ((val >> bit) & 1); }
+
+__kernel void binary_convolution(
+    const __global half *restrict src_data,
+    const __global uchar *restrict weights_data,
+    __global half *restrict dst_data,
+    float pad_value,
+
+    int IW,
+    int IH,
+    int IC,
+
+    int DW,
+    int DH,
+
+    int GC,
+
+    int KW,
+    int KH,
+
+    int PW,
+    int PH,
+
+    int SW,
+    int SH,
+
+    int OW)
+{
+    __local half src_local[32 * 1024];
+    __local half dst_local[2 * 1024];
+
+    const int oh = get_group_id(0);
+    const int oc = get_group_id(1);
+    const int OH = get_global_size(0);
+    const int OC = get_global_size(1);
+
+    const int gc = oc / (OC / GC);
+
+    if (oh * SH >= 0 && oh * SH <= IH - 1) {
+        const __global half *src = src_data + (gc * IC / GC) * IW * IH + (SH * oh) * IW;
+
+        event_t e1 = async_work_group_copy_2D2D(
+            src_local, // dst
+            src, // src
+            IW, // num_elements_per_line,
+            IC / GC, // num_lines,
+            IH * IW - IW, // src_line_stride,
+            0, // dst_line_stride,
+            0);
+        wait_group_events(1, &e1);
+    }
+
+    half pad_value_half = convert_half(pad_value);
+
+    //padding row
+    if (oh * SH > IH - 1) {
+        __local half *dst = src_local;
+        for (int c = 0; c < IC / GC; c++) {
+            #pragma unroll 8
+            for (int j = 0; j < IW; j++) {
+                dst[j] = pad_value_half;
+            }
+            dst += IW;
+        }
+    }
+
+    int OWS = SW * OW;
+    ushort8 in;
+
+    for (int ows8 = 0; ows8 < (OWS + 7) / 8; ows8++) {
+        ushort8 val = {0, 0, 0, 0, 0, 0, 0, 0};
+        for (int ic = 0; ic < IC / GC; ++ic) {
+            __local half *src = (__local half *)((__local half8 *)(src_local + ic * IW) + ows8);
+            int weight_pos    = oc * IC / GC + ic;
+            ushort w =
+                extract_weights(weights_data[((weight_pos + 0)) / 8], ((weight_pos + 0) % 8));
+
+            if ((ows8 * 8) <= IW - 1) {
+                in = *((__local ushort8 *)(src));
+            }
+
+            //padding column
+            if (ows8 * 8 + 7 > IW - 1) {
+                int boundary = (IW - 1) - ows8 * 8 + 1;
+                boundary     = boundary < 0 ? 0 : boundary;
+                for (int offset = boundary; offset < 8; offset++) {
+                    *((half *)(&in) + offset) = pad_value_half;
+                }
+            }
+
+            ushort8 w8 = (ushort8)(w);
+
+            ushort8 cond =
+                (((in) < (ushort8)0x8000) && (in > (ushort8)0x0000)) ? (ushort8)(1) : (ushort8)(0);
+
+            val += (cond ^ w8);
+        }
+
+        ushort8 val_shift = val << 1;
+        int boundary      = (ows8 * 8 + 7) / SW < OW - 1 ? (ows8 * 8 + 7) / SW : OW - 1;
+        for (int ow = (ows8 * 8 + SW - 1) / SW; ow <= boundary; ow++) {
+            *(dst_local + ow) = (half)(IC / GC - *((ushort *)(&val_shift) + ow * SW - ows8 * 8));
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy(dst_data + oc * OW * OH + oh * OW, dst_local, OW, 0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/binary_convolution3x3.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:55e3c3f8863ff7a3583bcc7340d1e226775f5f14cfb11dd32bd671764570f7cb
+size 104136

openvino/vpu_custom_kernels/binary_convolution3x3.cl

@@ -0,0 +1,278 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+ushort extract_weights(uchar val, int bit) { return ((val >> bit) & 1); }
+
+__kernel void binary_convolution(
+    const __global half *restrict src_data,
+    const __global uchar *restrict weights_data,
+    const __global half *restrict dst_data,
+    float pad_value,
+
+    int IW,
+    int IH,
+    int IC,
+
+    int DW,
+    int DH,
+
+    int GC,
+
+    int KW,
+    int KH,
+
+    int PW,
+    int PH,
+
+    int SW,
+    int SH,
+
+    int OW)
+{
+    __local half src_local[32 * 1024];
+    __local half dst_local[2 * 1024];
+
+    const int oh = get_group_id(0);
+    const int oc = get_group_id(1);
+    const int OH = get_global_size(0);
+    const int OC = get_global_size(1);
+
+    const int gc = oc / (OC / GC);
+
+    if (oh * SH - 1 >= 0 && oh * SH + DH + DH - 1 <= IH - 1) //dma for 3 rows
+    {
+        event_t e = async_work_group_copy_3D3D(
+            src_local, // dst
+            src_data + (gc * IC / GC) * IW * IH + (SH * oh - 1) * IW, // src
+            IW, // num_elements_per_line
+            3, // num_lines
+            DH * IW - IW, // src_line_stride
+            0, // dst_line_stride
+            IC / GC, // num planes
+            IH * IW - 3 * DH * IW, // src plane stride
+            0, // dst plane stride
+            0);
+        wait_group_events(1, &e);
+    } else {
+        int ih = oh * SH - 1;
+        if (ih >= 0 && ih <= IH - 1) //dma for first row
+        {
+            event_t e = async_work_group_copy_2D2D(
+                src_local, // dst
+                src_data + (gc * IC / GC) * IW * IH + ih * IW, // src
+                IW, // num_elements_per_line,
+                IC / GC, // num_lines,
+                IH * IW - IW, // src_line_stride,
+                2 * IW, // dst_line_stride,
+                0);
+
+            wait_group_events(1, &e);
+        }
+        ih = oh * SH - 1 + DH;
+        if (ih >= 0 && ih <= IH - 1) //dma for second row
+        {
+            event_t e = async_work_group_copy_2D2D(
+                src_local + IW, // dst
+                src_data + (gc * IC / GC) * IW * IH + ih * IW, // src
+                IW, // num_elements_per_line,
+                IC / GC, // num_lines,
+                IH * IW - IW, // src_line_stride,
+                2 * IW, // dst_line_stride,
+                0);
+            wait_group_events(1, &e);
+        }
+        ih = oh * SH - 1 + 2 * DH;
+        if (ih >= 0 && ih <= IH - 1) //dma for third row
+        {
+            event_t e = async_work_group_copy_2D2D(
+                src_local + 2 * IW, // dst
+                src_data + (gc * IC / GC) * IW * IH + ih * IW, // src
+                IW, // num_elements_per_line,
+                IC / GC, // num_lines,
+                IH * IW - IW, // src_line_stride,
+                2 * IW, // dst_line_stride,
+                0);
+            wait_group_events(1, &e);
+        }
+    }
+
+    half pad_value_half = convert_half(pad_value);
+
+    //padding row
+    if (oh * SH - 1 < 0 || oh * SH - 1 > IH - 1) {
+        __local half *dst = src_local;
+        for (int c = 0; c < IC / GC; c++) {
+            #pragma unroll 8
+            for (int j = 0; j < IW; j++) {
+                dst[j] = pad_value_half;
+            }
+            dst += 3 * IW;
+        }
+    }
+    if (oh * SH + DH - 1 > IH - 1) {
+        __local half *dst = src_local + IW;
+        for (int c = 0; c < IC / GC; c++) {
+            #pragma unroll 8
+            for (int j = 0; j < IW; j++) {
+                dst[j] = pad_value_half;
+            }
+            dst += 3 * IW;
+        }
+    }
+    if (oh * SH + DH + DH - 1 > IH - 1) {
+        __local half *dst = src_local + 2 * IW;
+        for (int c = 0; c < IC / GC; c++) {
+            #pragma unroll 8
+            for (int j = 0; j < IW; j++) {
+                dst[j] = pad_value_half;
+            }
+            dst += 3 * IW;
+        }
+    }
+
+    int OWS = SW * OW;
+
+    ushort8 in00;
+    ushort8 in01;
+    ushort8 in02;
+    ushort8 in10;
+    ushort8 in11;
+    ushort8 in12;
+    ushort8 in20;
+    ushort8 in21;
+    ushort8 in22;
+
+    for (int ows8 = 0; ows8 < (OWS + 7) / 8; ows8++) {
+        ushort8 val = {0, 0, 0, 0, 0, 0, 0, 0};
+        for (int ic = 0; ic < IC / GC; ++ic) {
+            __local half *src =
+                (__local half *)((__local half8 *)(src_local + ic * IW * 3 + IW + DW - 1) + ows8);
+            int weight_pos = oc * IC / GC * 3 * 3 + ic * 3 * 3;
+            ushort w0 = extract_weights(weights_data[((weight_pos + 0)) / 8], ((weight_pos + 0) % 8));
+            ushort w1 = extract_weights(weights_data[((weight_pos + 1)) / 8], ((weight_pos + 1) % 8));
+            ushort w2 = extract_weights(weights_data[((weight_pos + 2)) / 8], ((weight_pos + 2) % 8));
+            ushort w3 = extract_weights(weights_data[((weight_pos + 3)) / 8], ((weight_pos + 3) % 8));
+            ushort w4 = extract_weights(weights_data[((weight_pos + 4)) / 8], ((weight_pos + 4) % 8));
+            ushort w5 = extract_weights(weights_data[((weight_pos + 5)) / 8], ((weight_pos + 5) % 8));
+            ushort w6 = extract_weights(weights_data[((weight_pos + 6)) / 8], ((weight_pos + 6) % 8));
+            ushort w7 = extract_weights(weights_data[((weight_pos + 7)) / 8], ((weight_pos + 7) % 8));
+            ushort w8 = extract_weights(weights_data[((weight_pos + 8)) / 8], ((weight_pos + 8) % 8));
+
+            if ((ows8 * 8) - 1 <= IW - 1) {
+                in00 = *((__local ushort8 *)(src - IW - DW));
+                in01 = *((__local ushort8 *)(src - IW));
+                in02 = *((__local ushort8 *)(src - IW + DW));
+
+                in10 = *((__local ushort8 *)(src - DW));
+                in11 = *((__local ushort8 *)(src));
+                in12 = *((__local ushort8 *)(src + DW));
+
+                in20 = *((__local ushort8 *)(src + IW - DW));
+                in21 = *((__local ushort8 *)(src + IW));
+                in22 = *((__local ushort8 *)(src + IW + DW));
+            }
+
+            //padding column
+            if (ows8 * 8 - 1 < 0) {
+                int boundary = 1 - ows8 * 8;
+                boundary     = boundary > 8 ? 8 : boundary;
+                for (int offset = 0; offset < boundary; offset++) {
+                    *((half *)(&in00) + offset) = pad_value_half;
+                    *((half *)(&in10) + offset) = pad_value_half;
+                    *((half *)(&in20) + offset) = pad_value_half;
+                }
+            }
+            if ((ows8 * 8 + 7) + DW + DW - 1 > IW - 1) {
+                int boundary = (IW - DW - 1 - DW + 1) - ows8 * 8 + 1;
+                boundary     = boundary < 0 ? 0 : boundary;
+                for (int offset = boundary; offset < 8; offset++) {
+                    *((half *)(&in02) + offset) = pad_value_half;
+                    *((half *)(&in12) + offset) = pad_value_half;
+                    *((half *)(&in22) + offset) = pad_value_half;
+                }
+            }
+            if ((ows8 * 8 + 7) + DW - 1 > IW - 1) {
+                int boundary = (IW - 1 - DW + 1) - ows8 * 8 + 1;
+                boundary     = boundary < 0 ? 0 : boundary;
+                for (int offset = boundary; offset < 8; offset++) {
+                    *((half *)(&in01) + offset) = pad_value_half;
+                    *((half *)(&in11) + offset) = pad_value_half;
+                    *((half *)(&in21) + offset) = pad_value_half;
+                }
+            }
+            if ((ows8 * 8 + 7) - 1 > IW - 1) {
+                int boundary = (IW - 1 + 1) - ows8 * 8 + 1;
+                boundary     = boundary < 0 ? 0 : boundary;
+                for (int offset = boundary; offset < 8; offset++) {
+                    *((half *)(&in00) + offset) = pad_value_half;
+                    *((half *)(&in10) + offset) = pad_value_half;
+                    *((half *)(&in20) + offset) = pad_value_half;
+                }
+            }
+
+            ushort8 w00 = (ushort8)(w0);
+            ushort8 w01 = (ushort8)(w1);
+            ushort8 w02 = (ushort8)(w2);
+            ushort8 w10 = (ushort8)(w3);
+            ushort8 w11 = (ushort8)(w4);
+            ushort8 w12 = (ushort8)(w5);
+            ushort8 w20 = (ushort8)(w6);
+            ushort8 w21 = (ushort8)(w7);
+            ushort8 w22 = (ushort8)(w8);
+
+            ushort8 cond0 = (((in00) < (ushort8)0x8000) && (in00 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond1 = (((in01) < (ushort8)0x8000) && (in01 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond2 = (((in02) < (ushort8)0x8000) && (in02 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond3 = (((in10) < (ushort8)0x8000) && (in10 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond4 = (((in11) < (ushort8)0x8000) && (in11 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond5 = (((in12) < (ushort8)0x8000) && (in12 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond6 = (((in20) < (ushort8)0x8000) && (in20 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond7 = (((in21) < (ushort8)0x8000) && (in21 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+            ushort8 cond8 = (((in22) < (ushort8)0x8000) && (in22 > (ushort8)0x0000)) ?
+                                (ushort8)(1) :
+                                (ushort8)(0);
+
+            val += (cond0 ^ w00);
+            val += (cond1 ^ w01);
+            val += (cond2 ^ w02);
+            val += (cond3 ^ w10);
+            val += (cond4 ^ w11);
+            val += (cond5 ^ w12);
+            val += (cond6 ^ w20);
+            val += (cond7 ^ w21);
+            val += (cond8 ^ w22);
+        }
+
+        ushort8 val_shift = val << 1;
+        int boundary      = (ows8 * 8 + 7) / SW <= OW - 1 ? (ows8 * 8 + 7) / SW : OW - 1;
+        for (int ow = (ows8 * 8 + SW - 1) / SW; ow <= boundary; ow++) {
+            *(dst_local + ow) =
+                (half)(IC / GC * KH * KW - *((ushort *)(&val_shift) + ow * SW - ows8 * 8));
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy(dst_data + oc * OW * OH + oh * OW, dst_local, OW, 0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/convolution1x1_chw.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:8717c8429d41a69337007871137f06a9e6b38c685b5b3fecc634fade0eaa7e7f
+size 9220

openvino/vpu_custom_kernels/convolution1x1_chw.cl

@@ -0,0 +1,114 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__kernel void Convolution1x1_NCHW(
+    const __global half *in,
+    const __global half *out,
+    const __global half *w,
+    int IW,
+    int IH,
+    int IC,
+    int OW,
+    int OH,
+    int OC)
+{
+    __local half in_local[8 * 1024];
+    __local half out_local[8 * 1024];
+
+    event_t e1 = async_work_group_copy_2D2D(
+        in_local, // dst
+        in + get_group_id(0) * IW, // src
+        IW, // num_elements_per_line,
+        IC, // num_lines,
+        IW * IH - IW, // src_line_stride,
+        0, // dst_line_stride,
+        0);
+    wait_group_events(1, &e1);
+
+    int oh = get_global_id(0);
+    int oc = get_global_id(1);
+
+    int stride;
+    int write_output = 0;
+    __global half *src;
+
+    __global half8 *w8 = (__global half8 *)(&w[oc * IC]);
+    __global half *w1  = (__global half *)(&w[oc * IC]);
+
+    for (uint ow = 0; ow < (OW & (~0x7)); ow += 8) {
+        uint iw = ow;
+        uint ih = oh;
+
+        half8 val8_0 = 0.0f;
+
+        __local half8 *in8_0 = (__local half8 *)(&in_local[iw + 0 * IW]);
+        __local half8 *in8_1 = (__local half8 *)(&in_local[iw + 1 * IW]);
+        __local half8 *in8_2 = (__local half8 *)(&in_local[iw + 2 * IW]);
+        __local half8 *in8_3 = (__local half8 *)(&in_local[iw + 3 * IW]);
+        __local half8 *in8_4 = (__local half8 *)(&in_local[iw + 4 * IW]);
+        __local half8 *in8_5 = (__local half8 *)(&in_local[iw + 5 * IW]);
+        __local half8 *in8_6 = (__local half8 *)(&in_local[iw + 6 * IW]);
+        __local half8 *in8_7 = (__local half8 *)(&in_local[iw + 7 * IW]);
+
+        for (uint ic = 0; ic < IC / 8; ic++) {
+            val8_0 += (in8_0[ic * IW]) * ((half8)w8[ic].s0);
+            val8_0 += (in8_1[ic * IW]) * ((half8)w8[ic].s1);
+            val8_0 += (in8_2[ic * IW]) * ((half8)w8[ic].s2);
+            val8_0 += (in8_3[ic * IW]) * ((half8)w8[ic].s3);
+            val8_0 += (in8_4[ic * IW]) * ((half8)w8[ic].s4);
+            val8_0 += (in8_5[ic * IW]) * ((half8)w8[ic].s5);
+            val8_0 += (in8_6[ic * IW]) * ((half8)w8[ic].s6);
+            val8_0 += (in8_7[ic * IW]) * ((half8)w8[ic].s7);
+        }
+
+        for (uint ic = (IC & (~0x7)); ic < IC; ++ic) {
+            val8_0 += *((__local half8 *)(&in_local[iw + ic * IW])) * ((half8)w1[ic]);
+        }
+        *((__local half8 *)&out_local[ow + 0]) = (val8_0);
+    }
+
+    uint iw = (OW & (~0x7));
+    uint ih = oh;
+
+    half8 val8_0 = 0.0f;
+
+    __local half8 *in8_0 = (__local half8 *)(&in_local[iw + 0 * IW]);
+    __local half8 *in8_1 = (__local half8 *)(&in_local[iw + 1 * IW]);
+    __local half8 *in8_2 = (__local half8 *)(&in_local[iw + 2 * IW]);
+    __local half8 *in8_3 = (__local half8 *)(&in_local[iw + 3 * IW]);
+    __local half8 *in8_4 = (__local half8 *)(&in_local[iw + 4 * IW]);
+    __local half8 *in8_5 = (__local half8 *)(&in_local[iw + 5 * IW]);
+    __local half8 *in8_6 = (__local half8 *)(&in_local[iw + 6 * IW]);
+    __local half8 *in8_7 = (__local half8 *)(&in_local[iw + 7 * IW]);
+
+    for (uint ic = 0; ic < IC / 8; ic++) {
+        val8_0 += (in8_0[ic * IW]) * ((half8)w8[ic].s0);
+        val8_0 += (in8_1[ic * IW]) * ((half8)w8[ic].s1);
+        val8_0 += (in8_2[ic * IW]) * ((half8)w8[ic].s2);
+        val8_0 += (in8_3[ic * IW]) * ((half8)w8[ic].s3);
+        val8_0 += (in8_4[ic * IW]) * ((half8)w8[ic].s4);
+        val8_0 += (in8_5[ic * IW]) * ((half8)w8[ic].s5);
+        val8_0 += (in8_6[ic * IW]) * ((half8)w8[ic].s6);
+        val8_0 += (in8_7[ic * IW]) * ((half8)w8[ic].s7);
+    }
+
+    for (uint ic = (IC & (~0x7)); ic < IC; ++ic) {
+        val8_0 += *((__local half8 *)(&in_local[iw + ic * IW])) * ((half8)w1[ic]);
+    }
+    for (uint ow = (OW & (~0x7)); ow < OW; ow++) {
+        out_local[ow + 0] = (val8_0[ow % 8]);
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy(
+        out + get_group_id(1) * OW * OH + get_group_id(0) * OW,
+        out_local,
+        OW,
+        0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/convolution1x1_hwc.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:5b6122a6bf6f50d2c7fc612d4e286559f9c96746e166892d192e1264e1ce5a2c
+size 4304

openvino/vpu_custom_kernels/convolution1x1_hwc.cl

@@ -0,0 +1,126 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__kernel void Convolution1x1_NHWC(
+    const __global half *in,
+    const __global half *out,
+    const __global half *w,
+    int IW,
+    int IH,
+    int IC,
+    int OW,
+    int OH,
+    int OC)
+{
+
+    __local half in_local[8 * 1024];
+    __local half out_local[8 * 1024];
+
+    const int sizeAct = IW * IC;
+
+    event_t e1 = async_work_group_copy(in_local, in + get_group_id(0) * sizeAct, sizeAct, 0);
+    wait_group_events(1, &e1);
+
+    int oh = get_global_id(0);
+    int oc = get_global_id(1);
+
+    int stride;
+    int write_output = 0;
+    __global half *src;
+
+    __global half8 *w8 = (__global half8 *)(&w[oc * IC]);
+    __global half *w1  = (__global half *)(&w[oc * IC]);
+
+    for (uint ow = 0; ow < (OW & (~0x7)); ow += 8) {
+        uint iw = ow;
+        uint ih = oh;
+
+        half8 val8_0 = 0.0f;
+        half8 val8_1 = 0.0f;
+        half8 val8_2 = 0.0f;
+        half8 val8_3 = 0.0f;
+        half8 val8_4 = 0.0f;
+        half8 val8_5 = 0.0f;
+        half8 val8_6 = 0.0f;
+        half8 val8_7 = 0.0f;
+
+        __local half8 *in8_0 = (__local half8 *)(&in_local[(iw + 0) * IC]);
+        __local half8 *in8_1 = (__local half8 *)(&in_local[(iw + 1) * IC]);
+        __local half8 *in8_2 = (__local half8 *)(&in_local[(iw + 2) * IC]);
+        __local half8 *in8_3 = (__local half8 *)(&in_local[(iw + 3) * IC]);
+        __local half8 *in8_4 = (__local half8 *)(&in_local[(iw + 4) * IC]);
+        __local half8 *in8_5 = (__local half8 *)(&in_local[(iw + 5) * IC]);
+        __local half8 *in8_6 = (__local half8 *)(&in_local[(iw + 6) * IC]);
+        __local half8 *in8_7 = (__local half8 *)(&in_local[(iw + 7) * IC]);
+
+        for (uint ic = 0; ic < IC / 8; ++ic) {
+            val8_0 += (in8_0[ic]) * (w8[ic]);
+            val8_1 += (in8_1[ic]) * (w8[ic]);
+            val8_2 += (in8_2[ic]) * (w8[ic]);
+            val8_3 += (in8_3[ic]) * (w8[ic]);
+            val8_4 += (in8_4[ic]) * (w8[ic]);
+            val8_5 += (in8_5[ic]) * (w8[ic]);
+            val8_6 += (in8_6[ic]) * (w8[ic]);
+            val8_7 += (in8_7[ic]) * (w8[ic]);
+        }
+
+        half val_0 = 0.0f;
+        half val_1 = 0.0f;
+        half val_2 = 0.0f;
+        half val_3 = 0.0f;
+        half val_4 = 0.0f;
+        half val_5 = 0.0f;
+        half val_6 = 0.0f;
+        half val_7 = 0.0f;
+        for (uint ic = IC & (~0x7); ic < IC; ++ic) {
+            val_0 += *((__local half *)in8_0 + ic) * (*((__global half *)w8 + ic));
+            val_1 += *((__local half *)in8_1 + ic) * (*((__global half *)w8 + ic));
+            val_2 += *((__local half *)in8_2 + ic) * (*((__global half *)w8 + ic));
+            val_3 += *((__local half *)in8_3 + ic) * (*((__global half *)w8 + ic));
+            val_4 += *((__local half *)in8_4 + ic) * (*((__global half *)w8 + ic));
+            val_5 += *((__local half *)in8_5 + ic) * (*((__global half *)w8 + ic));
+            val_6 += *((__local half *)in8_6 + ic) * (*((__global half *)w8 + ic));
+            val_7 += *((__local half *)in8_7 + ic) * (*((__global half *)w8 + ic));
+        }
+        out_local[ow + 0] = __builtin_shave_sau_sumx_f16_r(val8_0) + val_0;
+        out_local[ow + 1] = __builtin_shave_sau_sumx_f16_r(val8_1) + val_1;
+        out_local[ow + 2] = __builtin_shave_sau_sumx_f16_r(val8_2) + val_2;
+        out_local[ow + 3] = __builtin_shave_sau_sumx_f16_r(val8_3) + val_3;
+        out_local[ow + 4] = __builtin_shave_sau_sumx_f16_r(val8_4) + val_4;
+        out_local[ow + 5] = __builtin_shave_sau_sumx_f16_r(val8_5) + val_5;
+        out_local[ow + 6] = __builtin_shave_sau_sumx_f16_r(val8_6) + val_6;
+        out_local[ow + 7] = __builtin_shave_sau_sumx_f16_r(val8_7) + val_7;
+    }
+    for (uint ow = (OW & (~0x7)); ow < OW; ow++) {
+
+        uint iw = ow;
+        uint ih = oh;
+
+        half8 val8 = 0.0f;
+
+        __local half8 *in8 = (__local half8 *)(&in_local[iw * IC]);
+
+        for (uint ic = 0; ic < IC / 8; ++ic) {
+            val8 += (in8[ic]) * (w8[ic]);
+        }
+
+        half val = 0.0f;
+        for (uint ic = (IC & (~0x7)); ic < IC; ++ic) {
+            val += (*((__local half *)in8 + ic)) * (*((__global half *)w8 + ic));
+        }
+        out_local[ow] = __builtin_shave_sau_sumx_f16_r(val8) + val;
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy(
+        out + get_group_id(1) * OW * OH + get_group_id(0) * OW,
+        out_local,
+        OW,
+        0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/convolution3x3.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:021bb40840ff35506972e6f6a7dea1b5f40a8db0927aaa9a6c116b152e386851
+size 5748

openvino/vpu_custom_kernels/convolution3x3.cl

@@ -0,0 +1,158 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__kernel void Convolution3x3(
+    const __global half *in_param,
+    const __global half *out,
+    const __global half *w,
+    int IW,
+    int IH,
+    int IC,
+    int OW,
+    int OH,
+    int OC,
+    int KX,
+    int KY,
+    int stride_x,
+    int stride_y,
+    int pad_x,
+    int pad_y,
+    int dilation_x,
+    int dilation_y)
+{
+    __local half in_local[8 * 1024];
+    __local half out_local[8 * 1024];
+    __local half w_local[8 * 1024];
+
+    const int sizePlane = IW * IH;
+    event_t e1          = async_work_group_copy_2D2D(
+        in_local, // dst
+        in_param + get_group_id(0) * stride_y * IW, // src
+        3 * IW, // num_elements_per_line,
+        IC, // num_lines,
+        IW * IH - 3 * IW, // src_line_stride,
+        0, // dst_line_stride,
+        0);
+    wait_group_events(1, &e1);
+
+    const int sizeWeight = IC * 3 * 3;
+    e1 = async_work_group_copy(w_local, w + get_group_id(1) * sizeWeight, sizeWeight, 0);
+    wait_group_events(1, &e1);
+
+    int oh = get_global_id(0);
+    int oc = get_global_id(1);
+
+    __local half *in = (__local half *)in_local + 1;
+
+    int stride;
+    int write_output = 0;
+    __local half *src;
+
+    if ((stride_x == 1) && (stride_y == 1)) {
+        stride       = OW / 8;
+        write_output = 1;
+    }
+    if ((stride_x == 2) && (stride_y == 2)) {
+        stride       = OW / 4;
+        write_output = 2;
+    }
+
+    for (int ow = 0; ow < stride; ow++) {
+        float8 val = {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f};
+        for (int ic = 0; ic < IC; ++ic) {
+            src             = (__local half *)((__local half8 *)(in + ic * IW * 3) + ow);
+            __local half *k = (__local half *)(w_local + ic * 3 * 3);
+
+            half8 aux_in00 = *((__local half8 *)src - 1);
+            half8 aux_in01 = *((__local half8 *)src + 0);
+            half8 aux_in02 = *((__local half8 *)src + 1);
+            half8 aux_in10 = *((__local half8 *)(src + IW) - 1);
+            half8 aux_in11 = *((__local half8 *)(src + IW) + 0);
+            half8 aux_in12 = *((__local half8 *)(src + IW) + 1);
+            half8 aux_in20 = *((__local half8 *)(src + IW * 2) - 1);
+            half8 aux_in21 = *((__local half8 *)(src + IW * 2) + 0);
+            half8 aux_in22 = *((__local half8 *)(src + IW * 2) + 1);
+
+            short8 in00 = *((short8 *)&aux_in00);
+            short8 in01 = *((short8 *)&aux_in01);
+            short8 in02 = *((short8 *)&aux_in02);
+            short8 in10 = *((short8 *)&aux_in10);
+            short8 in11 = *((short8 *)&aux_in11);
+            short8 in12 = *((short8 *)&aux_in12);
+            short8 in20 = *((short8 *)&aux_in20);
+            short8 in21 = *((short8 *)&aux_in21);
+            short8 in22 = *((short8 *)&aux_in22);
+
+            short8 aux_aux00 = __builtin_shave_cmu_alignvec_rri_short8(in00, in01, 14);
+            short8 aux_aux01 = in01;
+            short8 aux_aux02 = __builtin_shave_cmu_alignvec_rri_short8(in01, in02, 2);
+            short8 aux_aux10 = __builtin_shave_cmu_alignvec_rri_short8(in10, in11, 14);
+            short8 aux_aux11 = in11;
+            short8 aux_aux12 = __builtin_shave_cmu_alignvec_rri_short8(in11, in12, 2);
+            short8 aux_aux20 = __builtin_shave_cmu_alignvec_rri_short8(in20, in21, 14);
+            short8 aux_aux21 = in21;
+            short8 aux_aux22 = __builtin_shave_cmu_alignvec_rri_short8(in21, in22, 2);
+
+            half8 aux00 = *((half8 *)&aux_aux00);
+            half8 aux01 = *((half8 *)&aux_aux01);
+            half8 aux02 = *((half8 *)&aux_aux02);
+            half8 aux10 = *((half8 *)&aux_aux10);
+            half8 aux11 = *((half8 *)&aux_aux11);
+            half8 aux12 = *((half8 *)&aux_aux12);
+            half8 aux20 = *((half8 *)&aux_aux20);
+            half8 aux21 = *((half8 *)&aux_aux21);
+            half8 aux22 = *((half8 *)&aux_aux22);
+
+            half8 w00 = (half8)(*(k + 0));
+            half8 w01 = (half8)(*(k + 1));
+            half8 w02 = (half8)(*(k + 2));
+            half8 w10 = (half8)(*(k + 3));
+            half8 w11 = (half8)(*(k + 4));
+            half8 w12 = (half8)(*(k + 5));
+            half8 w20 = (half8)(*(k + 6));
+            half8 w21 = (half8)(*(k + 7));
+            half8 w22 = (half8)(*(k + 8));
+
+            val += convert_float8(aux00) * convert_float8(w00);
+            val += convert_float8(aux01) * convert_float8(w01);
+            val += convert_float8(aux02) * convert_float8(w02);
+            val += convert_float8(aux10) * convert_float8(w10);
+            val += convert_float8(aux11) * convert_float8(w11);
+            val += convert_float8(aux12) * convert_float8(w12);
+            val += convert_float8(aux20) * convert_float8(w20);
+            val += convert_float8(aux21) * convert_float8(w21);
+            val += convert_float8(aux22) * convert_float8(w22);
+        }
+        if (write_output == 2) *((__local half4 *)(out_local) + ow) = convert_half4(val.s0246);
+        if (write_output == 1) *((__local half8 *)(out_local) + ow) = convert_half8(val);
+    }
+
+    for (int ow = OW & ~(0x7); ow < OW; ow++) {
+        float val = 0.0f;
+        for (int ic = 0; ic < IC; ++ic) {
+            for (int ky = 0; ky < 3; ++ky) {
+                for (int kx = 0; kx < 3; ++kx) {
+                    int iw = ow * stride_x - pad_x + kx * dilation_x;
+                    int ih = oh * stride_y - pad_y + ky * dilation_y;
+
+                    val += convert_float(in[ic * IW * 3 + (ky * dilation_y) * IW + iw])
+                           * convert_float(w_local[ic * 3 * 3 + ky * 3 + kx]);
+                }
+            }
+        }
+        out_local[ow] = convert_half(val);
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy(
+        out + get_group_id(1) * OW * OH + get_group_id(0) * OW,
+        out_local,
+        OW,
+        0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/correlate.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:e2b24b1b5bfd1786128682ee814230653b4b63aad5b472feec9c6f4a4c833e2f
+size 14336

openvino/vpu_custom_kernels/correlate.cl

@@ -0,0 +1,453 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define MAX_OPENCL_BUFF_SIZE 64 * 1024
+
+#define USE_DMA 1
+
+#if defined(USE_DMA)
+void dmacpyLineSrcStrideStart(global half *from, private half *to, int size, int src_width, int src_stride)
+{
+    item_dma_event_t copyEvent =
+        WorkItemDmaCreateStrideTransaction(from, to, src_width, src_width, src_stride, src_width, size, 0);
+    WaitWorkItemDmaEvents(1, &copyEvent);
+}
+
+void dmacpyLineDstStrideStart(private half *from, global half *to, int size, int src_width, int src_stride)
+{
+    item_dma_event_t copyEvent =
+        WorkItemDmaCreateStrideTransaction(from, to, src_width, src_width, src_width, src_stride, size, 0);
+    WaitWorkItemDmaEvents(1, &copyEvent);
+}
+#endif
+
+void memzero(void *ptr, size_t num)
+{
+    float4 *line0_ = (float4 *)ptr;
+    #pragma unroll 16
+    for (int i = 0; i < num / 16; i++) {
+        line0_[i] = (float4){0.f, 0.f, 0.f, 0.f};
+    }
+    uchar *ptr_ = (uchar *)ptr;
+    for (int i = num / 16 * 16; i < num; i++) {
+        ptr_[i] = 0;
+    }
+}
+
+void __attribute__((noinline)) crosscorrh(
+    __private const half *restrict line0,
+    __private const half *restrict line1,
+    __private half *restrict dline,
+    int topwidth,
+    int max_displacement,
+    int neighborhood_grid_radius,
+    int kernel_size,
+    int padding,
+    int bottomwidth,
+    int stride1,
+    int stride2,
+    int max_channels,
+    int cur_subchannels)
+{
+    if (max_channels == 64) {
+        for (int i = 0; i < kernel_size; i++) {
+            int x1      = max_displacement - padding + i;
+            int offset1 = x1 >= 0 ? 0 : (-x1 + stride1 - 1) / stride1;
+            x1 += offset1 * stride1;
+
+            for (int blockIdx_x = offset1; blockIdx_x < topwidth && x1 < bottomwidth; blockIdx_x++, x1 += stride1) {
+                int x2      = x1 - neighborhood_grid_radius * stride2;
+                int offset2 = x2 >= 0 ? 0 : (-x2 + stride2 - 1) / stride2;
+                x2 += offset2 * stride2;
+
+                for (int top_channel_x = offset2 - neighborhood_grid_radius;
+                     top_channel_x <= neighborhood_grid_radius && x2 < bottomwidth;
+                     top_channel_x++, x2 += stride2) {
+                    half8 sum4 = (half8){0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
+
+                    half8 *src0 = (half8 *)(line0 + x1 * max_channels);
+                    half8 *src1 = (half8 *)(line1 + x2 * max_channels);
+
+                    #pragma unroll 8
+                    for (int ch = 0; ch < max_channels / 8; ch++) sum4 += (src0[ch]) * (src1[ch]);
+
+                    half sum = __builtin_shave_sau_sumx_f16_r(sum4);
+                    dline[(top_channel_x + neighborhood_grid_radius) * topwidth + blockIdx_x] += (sum);
+                }
+            }
+        }
+    } else {
+        int neighborhood_grid_width = 2 * neighborhood_grid_radius + 1;
+
+        for (int blockIdx_x = 0; blockIdx_x < topwidth; blockIdx_x++) {
+            for (int i = 0; i < kernel_size; i++) {
+                int x1 = blockIdx_x * stride1 + max_displacement + i - padding;
+
+                if ((x1 >= 0) && (x1 < bottomwidth)) {
+                    int o_min = -neighborhood_grid_radius * stride2;
+                    int o_max = neighborhood_grid_width * stride2 - neighborhood_grid_radius * stride2;
+                    if ((o_min) < (-x1)) {
+                        o_min -= ((x1 + o_min - (stride2 - 1)) / stride2) * stride2;
+                    }
+                    if ((o_max) >= (bottomwidth + stride2 - x1)) {
+                        o_max -= ((x1 + o_max - bottomwidth) / stride2) * stride2;
+                    }
+
+                    int o = o_min;
+                    for (; o <= o_max - 4 * stride2; o += 4 * stride2) {
+                        half8 *bottom0   = (half8 *)(line0 + x1 * max_channels);
+                        half8 *bottom1_0 = (half8 *)(line1 + (x1 + o + 0 * stride2) * max_channels);
+                        half8 *bottom1_1 = (half8 *)(line1 + (x1 + o + 1 * stride2) * max_channels);
+                        half8 *bottom1_2 = (half8 *)(line1 + (x1 + o + 2 * stride2) * max_channels);
+                        half8 *bottom1_3 = (half8 *)(line1 + (x1 + o + 3 * stride2) * max_channels);
+
+                        int c = 0;
+
+                        half8 sum40 = 0;
+                        half8 sum41 = 0;
+                        half8 sum42 = 0;
+                        half8 sum43 = 0;
+
+                        for (; c <= cur_subchannels / 8 - 4; c += 4) {
+                            sum40 += bottom0[c + 0] * bottom1_0[c + 0];
+                            sum40 += bottom0[c + 1] * bottom1_0[c + 1];
+                            sum40 += bottom0[c + 2] * bottom1_0[c + 2];
+                            sum40 += bottom0[c + 3] * bottom1_0[c + 3];
+
+                            sum41 += bottom0[c + 0] * bottom1_1[c + 0];
+                            sum41 += bottom0[c + 1] * bottom1_1[c + 1];
+                            sum41 += bottom0[c + 2] * bottom1_1[c + 2];
+                            sum41 += bottom0[c + 3] * bottom1_1[c + 3];
+
+                            sum42 += bottom0[c + 0] * bottom1_2[c + 0];
+                            sum42 += bottom0[c + 1] * bottom1_2[c + 1];
+                            sum42 += bottom0[c + 2] * bottom1_2[c + 2];
+                            sum42 += bottom0[c + 3] * bottom1_2[c + 3];
+
+                            sum43 += bottom0[c + 0] * bottom1_3[c + 0];
+                            sum43 += bottom0[c + 1] * bottom1_3[c + 1];
+                            sum43 += bottom0[c + 2] * bottom1_3[c + 2];
+                            sum43 += bottom0[c + 3] * bottom1_3[c + 3];
+                        }
+
+                        for (; c < cur_subchannels / 8; c++) {
+                            sum40 += bottom0[c] * bottom1_0[c];
+                            sum41 += bottom0[c] * bottom1_1[c];
+                            sum42 += bottom0[c] * bottom1_2[c];
+                            sum43 += bottom0[c] * bottom1_3[c];
+                        }
+
+                        half sum0 = __builtin_shave_sau_sumx_f16_r(sum40);
+                        half sum1 = __builtin_shave_sau_sumx_f16_r(sum41);
+                        half sum2 = __builtin_shave_sau_sumx_f16_r(sum42);
+                        half sum3 = __builtin_shave_sau_sumx_f16_r(sum43);
+
+                        for (c = c * 8; c < cur_subchannels; c++) {
+                            sum0 += line0[x1 * max_channels + c] * line1[(x1 + o + 0 * stride2) * max_channels + c];
+                            sum1 += line0[x1 * max_channels + c] * line1[(x1 + o + 1 * stride2) * max_channels + c];
+                            sum2 += line0[x1 * max_channels + c] * line1[(x1 + o + 2 * stride2) * max_channels + c];
+                            sum3 += line0[x1 * max_channels + c] * line1[(x1 + o + 3 * stride2) * max_channels + c];
+                        }
+
+                        dline[blockIdx_x + (((o / stride2) + 0) * topwidth + neighborhood_grid_radius * topwidth)] +=
+                            sum0;
+                        dline[blockIdx_x + (((o / stride2) + 1) * topwidth + neighborhood_grid_radius * topwidth)] +=
+                            sum1;
+                        dline[blockIdx_x + (((o / stride2) + 2) * topwidth + neighborhood_grid_radius * topwidth)] +=
+                            sum2;
+                        dline[blockIdx_x + (((o / stride2) + 3) * topwidth + neighborhood_grid_radius * topwidth)] +=
+                            sum3;
+                    }
+
+                    for (; o < o_max; o += 1 * stride2) {
+                        half8 *bottom0 = (half8 *)(line0 + x1 * max_channels);
+                        half8 *bottom1 = (half8 *)(line1 + (x1 + o) * max_channels);
+
+                        int c = 0;
+
+                        half8 sum4 = 0;
+                        for (; c <= cur_subchannels / 8 - 4; c += 4) {
+                            sum4 += bottom0[c + 0] * bottom1[c + 0];
+                            sum4 += bottom0[c + 1] * bottom1[c + 1];
+                            sum4 += bottom0[c + 2] * bottom1[c + 2];
+                            sum4 += bottom0[c + 3] * bottom1[c + 3];
+                        }
+                        for (; c < cur_subchannels / 8; c++) {
+                            sum4 += bottom0[c] * bottom1[c];
+                        }
+
+                        half sum = __builtin_shave_sau_sumx_f16_r(sum4);
+
+                        for (c = c * 8; c < cur_subchannels; c++) {
+                            sum += line0[x1 * max_channels + c] * line1[(x1 + o) * max_channels + c];
+                        }
+
+                        dline[blockIdx_x + (((o + neighborhood_grid_radius * stride2) / stride2) * topwidth)] += sum;
+                    }
+                }
+            }
+        }
+    }
+}
+
+__kernel void correlate2_half(
+    __global const half *restrict bottom0,
+    __global const half *restrict bottom1,
+    __global half *restrict top,
+    int topwidth,
+    int topheight,
+    int bottomwidth,
+    int bottomheight,
+    int bottomchannels,
+    int max_displacement,
+    int padding,
+    int neighborhood_grid_radius,
+    int neighborhood_grid_width,
+    int kernel_size,
+    int stride1,
+    int stride2)
+{
+    int max_channels = (MAX_OPENCL_BUFF_SIZE / sizeof(half) - topwidth * neighborhood_grid_width) / (3 * bottomwidth);
+    if (max_channels > 64) max_channels = 64;
+    int subchannels_count = (bottomchannels + max_channels - 1) / max_channels;
+    int subchannels       = (bottomchannels + subchannels_count - 1) / subchannels_count;
+    if (subchannels < max_channels) subchannels = max_channels;
+
+    const int sumelems = kernel_size * kernel_size * bottomchannels;
+
+    __private half cmx[MAX_OPENCL_BUFF_SIZE / sizeof(half)];
+
+    __private half *line0 = cmx;
+    __private half *line1 = line0 + bottomwidth * subchannels;
+    __private half *dline = line1 + bottomwidth * subchannels;
+
+    int blockIdx_y = get_global_id(0);
+
+#if defined(USE_DMA)
+    __private half *dmabuf = dline + topwidth * neighborhood_grid_width;
+#endif
+
+    int y1 = blockIdx_y * stride1 + max_displacement;
+
+    for (int j = 0; j < kernel_size; j++) {
+        for (int bottomchannel = 0; bottomchannel < bottomchannels; bottomchannel += subchannels) {
+            // configure channel batching
+            int startchannel = bottomchannel;
+            int endchannel = startchannel + subchannels > bottomchannels ? bottomchannels : startchannel + subchannels;
+            int deltachannels = endchannel - startchannel;
+
+            // load line form blob 0 with repackaging
+            if (y1 + j - padding >= 0 && y1 + j - padding < bottomheight) {
+#if defined(USE_DMA)
+                __global const half *curr =
+                    bottom0 + startchannel * bottomheight * bottomwidth + (y1 + j - padding) * bottomwidth;
+                dmacpyLineSrcStrideStart(
+                    curr,
+                    dmabuf,
+                    bottomwidth * deltachannels * sizeof(half),
+                    bottomwidth * sizeof(half),
+                    bottomwidth * bottomheight * sizeof(half));
+
+                for (int ch = 0; ch < deltachannels; ch++) {
+                    for (int blockIdx_x = 0; blockIdx_x < bottomwidth / 8; blockIdx_x++) {
+                        half8 val = ((half8 *)(dmabuf + ch * bottomwidth))[blockIdx_x];
+                        line0[(blockIdx_x * 8 + 0) * max_channels + ch] = val[0];
+                        line0[(blockIdx_x * 8 + 1) * max_channels + ch] = val[1];
+                        line0[(blockIdx_x * 8 + 2) * max_channels + ch] = val[2];
+                        line0[(blockIdx_x * 8 + 3) * max_channels + ch] = val[3];
+
+                        line0[(blockIdx_x * 8 + 4) * max_channels + ch] = val[4];
+                        line0[(blockIdx_x * 8 + 5) * max_channels + ch] = val[5];
+                        line0[(blockIdx_x * 8 + 6) * max_channels + ch] = val[6];
+                        line0[(blockIdx_x * 8 + 7) * max_channels + ch] = val[7];
+                    }
+
+                    for (int blockIdx_x = bottomwidth / 8 * 8; blockIdx_x < bottomwidth; blockIdx_x++) {
+                        line0[(blockIdx_x)*max_channels + ch] = dmabuf[blockIdx_x + ch * bottomwidth];
+                    }
+                }
+
+                if (deltachannels < subchannels)
+                    for (int blockIdx_x = 0; blockIdx_x < bottomwidth; blockIdx_x++)
+                        memzero(
+                            line0 + blockIdx_x * max_channels + deltachannels,
+                            (subchannels - deltachannels) * sizeof(half));
+#else
+                for (int blockIdx_x = 0; blockIdx_x < bottomwidth; blockIdx_x++) {
+                    for (int ch = 0; ch < deltachannels; ch++)
+                        line0[blockIdx_x * max_channels + ch] = bottom0
+                            [(ch + startchannel) * bottomheight * bottomwidth + (y1 + j - padding) * bottomwidth
+                             + blockIdx_x];
+
+                    if (deltachannels < subchannels)
+                        memzero(
+                            line0 + blockIdx_x * max_channels + deltachannels,
+                            (subchannels - deltachannels) * sizeof(half));
+                }
+#endif
+            } else
+                memzero(line0, max_channels * bottomwidth * sizeof(half));
+
+            for (int top_channel_y = 0; top_channel_y < neighborhood_grid_width; top_channel_y++) {
+                int y2 = y1 + (top_channel_y - neighborhood_grid_radius) * stride2;
+
+                if (y2 + j - padding >= 0 && y2 + j - padding < bottomheight) {
+#if defined(USE_DMA)
+                    __global const half *curr =
+                        bottom1 + startchannel * bottomheight * bottomwidth + (y2 + j - padding) * bottomwidth;
+                    dmacpyLineSrcStrideStart(
+                        curr,
+                        dmabuf,
+                        bottomwidth * deltachannels * sizeof(half),
+                        bottomwidth * sizeof(half),
+                        bottomwidth * bottomheight * sizeof(half));
+
+                    for (int ch = 0; ch < deltachannels; ch++) {
+                        for (int blockIdx_x = 0; blockIdx_x < bottomwidth / 8; blockIdx_x++) {
+                            half8 val = ((half8 *)(dmabuf + ch * bottomwidth))[blockIdx_x];
+                            line1[(blockIdx_x * 8 + 0) * max_channels + ch] = val[0];
+                            line1[(blockIdx_x * 8 + 1) * max_channels + ch] = val[1];
+                            line1[(blockIdx_x * 8 + 2) * max_channels + ch] = val[2];
+                            line1[(blockIdx_x * 8 + 3) * max_channels + ch] = val[3];
+
+                            line1[(blockIdx_x * 8 + 4) * max_channels + ch] = val[4];
+                            line1[(blockIdx_x * 8 + 5) * max_channels + ch] = val[5];
+                            line1[(blockIdx_x * 8 + 6) * max_channels + ch] = val[6];
+                            line1[(blockIdx_x * 8 + 7) * max_channels + ch] = val[7];
+                        }
+
+                        for (int blockIdx_x = bottomwidth / 8 * 8; blockIdx_x < bottomwidth; blockIdx_x++) {
+                            line1[(blockIdx_x)*max_channels + ch] = dmabuf[blockIdx_x + ch * bottomwidth];
+                        }
+                    }
+#else
+                    for (int ch = 0; ch < deltachannels; ch++) {
+                        for (int blockIdx_x = 0; blockIdx_x < bottomwidth / 8; blockIdx_x++) {
+                            half8 val = ((
+                                __global half8
+                                    *)(bottom1 + (ch + startchannel) * bottomheight * bottomwidth + (y2 + j - padding) * bottomwidth))
+                                [blockIdx_x];
+                            line1[(blockIdx_x * 8 + 0) * max_channels + ch] = val[0];
+                            line1[(blockIdx_x * 8 + 1) * max_channels + ch] = val[1];
+                            line1[(blockIdx_x * 8 + 2) * max_channels + ch] = val[2];
+                            line1[(blockIdx_x * 8 + 3) * max_channels + ch] = val[3];
+
+                            line1[(blockIdx_x * 8 + 4) * max_channels + ch] = val[4];
+                            line1[(blockIdx_x * 8 + 5) * max_channels + ch] = val[5];
+                            line1[(blockIdx_x * 8 + 6) * max_channels + ch] = val[6];
+                            line1[(blockIdx_x * 8 + 7) * max_channels + ch] = val[7];
+                        }
+                        for (int blockIdx_x = bottomwidth / 8 * 8; blockIdx_x < bottomwidth; blockIdx_x++) {
+                            half val =
+                                (bottom1 + (ch + startchannel) * bottomheight * bottomwidth
+                                 + (y2 + j - padding) * bottomwidth)[blockIdx_x];
+                            line1[(blockIdx_x)*max_channels + ch] = val;
+                        }
+                    }
+#endif
+                    for (int blockIdx_x = 0; blockIdx_x < bottomwidth; blockIdx_x++) {
+                        if (deltachannels < subchannels)
+                            memzero(
+                                line1 + blockIdx_x * max_channels + deltachannels,
+                                (subchannels - deltachannels) * sizeof(half));
+                    }
+                } else
+                    memzero(line1, max_channels * bottomwidth * sizeof(half));
+
+                if (j == 0 && startchannel == 0) {
+                    memzero(dline, neighborhood_grid_width * topwidth * sizeof(half));
+                } else {
+#if defined(USE_DMA)
+                    dmacpyLineSrcStrideStart(
+                        top + top_channel_y * neighborhood_grid_width * topheight * topwidth + blockIdx_y * topwidth,
+                        dline,
+                        topwidth * neighborhood_grid_width * sizeof(half),
+                        topwidth * sizeof(half),
+                        topwidth * topheight * sizeof(half));
+#else
+                    for (int top_channel_x = 0; top_channel_x < neighborhood_grid_width; top_channel_x++) {
+                        for (int blockIdx_x = 0; blockIdx_x < topwidth / 8; blockIdx_x++) {
+                            half8 val = ((
+                                __global half8
+                                    *)(top + ((top_channel_y * neighborhood_grid_width + top_channel_x) * topheight * topwidth + blockIdx_y * topwidth)))
+                                [blockIdx_x];
+                            ((half8 *)(dline + top_channel_x * topwidth))[blockIdx_x] = val;
+                        }
+                        for (int blockIdx_x = (topwidth / 8) * 8; blockIdx_x < topwidth; blockIdx_x++) {
+                            dline[top_channel_x * topwidth + blockIdx_x] =
+                                top[(top_channel_y * neighborhood_grid_width + top_channel_x) * topheight * topwidth
+                                    + blockIdx_y * topwidth + blockIdx_x];
+                        }
+                    }
+#endif
+                }
+
+                if (y1 + j - padding >= 0 && y1 + j - padding < bottomheight && y2 + j - padding >= 0
+                    && y2 + j - padding < bottomheight) {
+                    crosscorrh(
+                        line0,
+                        line1,
+                        dline,
+                        topwidth,
+                        max_displacement,
+                        neighborhood_grid_radius,
+                        kernel_size,
+                        padding,
+                        bottomwidth,
+                        stride1,
+                        stride2,
+                        max_channels,
+                        subchannels);
+                }
+
+                if (j == kernel_size - 1 && endchannel == bottomchannels) {
+                    half8 scale = (half8){
+                        (half)sumelems,
+                        (half)sumelems,
+                        (half)sumelems,
+                        (half)sumelems,
+                        (half)sumelems,
+                        (half)sumelems,
+                        (half)sumelems,
+                        (half)sumelems};
+                    for (int top_channel_x = 0; top_channel_x < neighborhood_grid_width; top_channel_x++) {
+                        for (int blockIdx_x = 0; blockIdx_x < topwidth / 8; blockIdx_x++) {
+                            ((half8 *)(dline + top_channel_x * topwidth))[blockIdx_x] =
+                                ((half8 *)(dline + top_channel_x * topwidth))[blockIdx_x] / scale;
+                        }
+                        for (int blockIdx_x = (topwidth / 8) * 8; blockIdx_x < topwidth; blockIdx_x++) {
+                            dline[top_channel_x * topwidth + blockIdx_x] =
+                                dline[top_channel_x * topwidth + blockIdx_x] / (half)sumelems;
+                        }
+                    }
+                }
+
+#if defined(USE_DMA)
+                dmacpyLineDstStrideStart(
+                    dline,
+                    top + top_channel_y * neighborhood_grid_width * topheight * topwidth + blockIdx_y * topwidth,
+                    topwidth * neighborhood_grid_width * sizeof(half),
+                    topwidth * sizeof(half),
+                    topwidth * topheight * sizeof(half));
+#else
+                for (int top_channel_x = 0; top_channel_x < neighborhood_grid_width; top_channel_x++) {
+                    for (int blockIdx_x = 0; blockIdx_x < topwidth / 8; blockIdx_x++) {
+                        ((__global half8
+                              *)(top + ((top_channel_y * neighborhood_grid_width + top_channel_x) * topheight * topwidth + blockIdx_y * topwidth)))
+                            [blockIdx_x] = ((half8 *)(dline + top_channel_x * topwidth))[blockIdx_x]
+                                           + (half8){0, 0, 0, 0, 0, 0, 0, 0};
+                    }
+                    for (int blockIdx_x = (topwidth / 8) * 8; blockIdx_x < topwidth; blockIdx_x++) {
+                        top[(top_channel_y * neighborhood_grid_width + top_channel_x) * topheight * topwidth
+                            + blockIdx_y * topwidth + blockIdx_x] =
+                            dline[top_channel_x * topwidth + blockIdx_x] + (half)0;
+                    }
+                }
+#endif
+            }
+        }
+    }
+}

openvino/vpu_custom_kernels/ctc.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:292de0fbb8dc6ead6970576d1b9a26a323fc9febfceb92c3af6b84496d523def
+size 10196

openvino/vpu_custom_kernels/ctc.cl

@@ -0,0 +1,94 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__global half *find(__global const half *begin, __global const half *end, half value)
+{
+    while (begin != end) {
+        if (*begin == value) {
+            return begin;
+        }
+        ++begin;
+    }
+    return end;
+}
+
+__kernel void CTCDecoder(
+    __global half *restrict probabilities,
+    __global half *restrict sequence_indicators,
+    __global half *restrict output,
+    int width,
+    int height,
+    int channels)
+{
+    __local half local_src[88 * 1 * 77];
+    __local half local_dst[88 * 1];
+
+    event_t e1 = async_work_group_copy_2D2D(
+        local_src, // dst
+        probabilities, // src
+        width, // num_elements_per_line,
+        height * channels, // num_lines,
+        width * (height - 1), // src_line_stride,
+        width * (height - 1), // dst_line_stride,
+        0);
+
+    wait_group_events(1, &e1);
+
+    const int T = channels; // Time
+    const int B = height; // Batches
+    const int C = width; // Chars
+
+    #pragma unroll 4
+    for (int i = 0; i < B * T; i++) {
+        local_dst[i] = -1.h;
+    }
+
+    int output_index = 0;
+
+    for (int b = 0; b < B; ++b) {
+        __global const half *restrict seq_ind = sequence_indicators + b * T;
+        const int seq_len = find(seq_ind + 1, seq_ind + T, 0.h) - seq_ind;
+        const int time    = min(seq_len, T);
+
+        int prev_class_idx = -1;
+
+        #pragma unroll 4
+        for (int t = 0; t < time; ++t) {
+            __local const half *restrict probs = local_src + b * C + t * C * B;
+
+            int max_class_idx = 0;
+            half max_prob     = probs[0];
+            for (int c = 1; c < C; ++c) {
+                const half prob = probs[c];
+                if (prob > max_prob) {
+                    max_class_idx = c;
+                    max_prob      = prob;
+                }
+            }
+
+            if (max_class_idx < C - 1 && max_class_idx != prev_class_idx) {
+                local_dst[b * T + output_index] = (half)max_class_idx;
+                output_index++;
+            }
+
+            prev_class_idx = max_class_idx;
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy_2D2D(
+        output, // dst
+        local_dst, // src
+        channels, // num_elements_per_line,
+        height, // num_lines,
+        0, // src_line_stride,
+        0, // dst_line_stride,
+        0);
+
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/customLayerBindings.xml

@@ -0,0 +1,507 @@
+<CustomLayer name="ReorgYolo" type="MVCL" version="1">
+    <Kernel entry="reorg_hwc_naive">
+        <Source filename="reorg_hwc_naive.bin"/>
+        <Parameters>
+            <Tensor arg-name="src" type="input" port-index="0" format="BYXF"/>
+            <Tensor arg-name="dst" type="output" port-index="0" format="BYXF"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="C" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="stride" type="int" source="stride"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="F,1,1" local="stride*stride,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="ReorgYolo" type="MVCL" version="1">
+    <Kernel entry="reorg_chw">
+        <Source filename="reorg_chw.bin"/>
+        <Parameters>
+            <Tensor arg-name="src" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="C" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="stride" type="int" source="stride"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="Y*F/(stride*stride),stride*stride,1" local="stride,stride,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="RegionYolo" type="MVCL" version="1">
+    <Where do_softmax="1"/>
+    <Kernel entry="region_chw">
+        <Source filename="region_chw.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="classes" type="int" source="classes"/>
+            <Scalar arg-name="coords" type="int" source="coords"/>
+            <Scalar arg-name="num" type="int" source="num"/>
+            <Scalar arg-name="maskSize" type="int" source="3"/>
+            <Scalar arg-name="doSoftmax" type="int" source="do_softmax"/>
+        </Parameters>
+        <WorkSizes global="((X*Y+7)/8)*8,num,1" local="((X*Y+7)/8)*8,1,1" dim="input,0"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="RegionYolo" type="MVCL" version="1">-->
+   <Where do_softmax="0" mask="0,1,2"/>
+   <Kernel entry="region_chw">
+       <Source filename="region_chw.bin"/>
+       <Parameters>
+           <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+           <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+           <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+           <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+           <Scalar arg-name="classes" type="int" source="classes"/>
+           <Scalar arg-name="coords" type="int" source="coords"/>
+           <Scalar arg-name="num" type="int" source="num"/>
+           <Scalar arg-name="maskSize" type="int" source="3"/>
+           <Scalar arg-name="doSoftmax" type="int" source="do_softmax"/>
+       </Parameters>
+       <WorkSizes global="((X*Y+7)/8)*8,3,1" local="((X*Y+7)/8)*8,1,1" dim="input,0"/>
+   </Kernel>
+</CustomLayer>
+
+<CustomLayer name="RegionYolo" type="MVCL" version="1">
+   <Where do_softmax="1"/>
+   <Kernel entry="region_hwc">
+       <Source filename="region_hwc.bin"/>
+       <Parameters>
+           <Tensor arg-name="src" type="input" port-index="0" format="BYXF"/>
+           <Tensor arg-name="dst" type="output" port-index="0" format="BYXF"/>
+           <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+           <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+           <Scalar arg-name="classes" type="int" source="classes"/>
+           <Scalar arg-name="coords" type="int" source="coords"/>
+           <Scalar arg-name="num" type="int" source="num"/>
+           <Scalar arg-name="maskSize" type="int" source="3"/>
+           <Scalar arg-name="doSoftmax" type="int" source="do_softmax"/>
+       </Parameters>
+       <WorkSizes global="((X*Y+7)/8)*8,num,1" local="((X*Y+7)/8)*8,1,1" dim="input,0"/>
+   </Kernel>
+</CustomLayer>
+
+<CustomLayer name="RegionYolo" type="MVCL" version="1">
+   <Where do_softmax="0" mask="0,1,2"/>
+   <Kernel entry="region_hwc">
+       <Source filename="region_hwc.bin"/>
+       <Parameters>
+           <Tensor arg-name="src" type="input" port-index="0" format="BYXF"/>
+           <Tensor arg-name="dst" type="output" port-index="0" format="BYXF"/>
+           <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+           <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+           <Scalar arg-name="classes" type="int" source="classes"/>
+           <Scalar arg-name="coords" type="int" source="coords"/>
+           <Scalar arg-name="num" type="int" source="num"/>
+           <Scalar arg-name="maskSize" type="int" source="3"/>
+           <Scalar arg-name="doSoftmax" type="int" source="do_softmax"/>
+       </Parameters>
+       <WorkSizes global="((X*Y+7)/8)*8,3,1" local="((X*Y+7)/8)*8,1,1" dim="input,0"/>
+   </Kernel>
+</CustomLayer>
+
+<!-- Pixel-wise kernel binding, local work group config is per line in the input tensor  -->
+<CustomLayer name="GRN" type="MVCL" version="1">
+    <Kernel entry="grn">
+        <Source filename="grn.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="C" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="bias" type="float" source="bias"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="X,Y,1" local="X,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<!-- Two stage layer binding, first kernel computes mean and variance, the second one normalizes input tensor-->
+<CustomLayer name="MVN" type="MVCL" version="1">
+    <Kernel entry="reduction_mean" stage="0">
+        <Source filename="mvn_reduction.bin"/>
+        <Parameters>
+            <Tensor arg-name="src" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="mean" type="output_buffer" port-index="0" dim="output,0" size="Y*F*4"/>
+            <Tensor arg-name="variance" type="output_buffer" port-index="1" dim="output,0" size="Y*F*4"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="across_channels" type="int" source="across_channels"/>
+        </Parameters>
+        <WorkSizes dim="output,0" global="1,Y,F" local="1,1,1"/>
+    </Kernel>
+    <Kernel entry="mvn_scale" stage="1">
+        <Source filename="mvn_scale.bin"/>
+        <Parameters>
+            <Tensor arg-name="src" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst" type="output" port-index="0" format="BFYX"/>
+            <Tensor arg-name="mean_part" type="input_buffer" port-index="0" dim="output,0" size="Y*F*4"/>
+            <Tensor arg-name="power_mean" type="input_buffer" port-index="1" dim="output,0" size="Y*F*4"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="H1" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="across_channels" type="int" source="across_channels"/>
+            <Scalar arg-name="normalize_variance" type="int" source="normalize_variance"/>
+            <Scalar arg-name="nparts" type="int" port-index="0" source="I.Y"/>
+        </Parameters>
+        <WorkSizes dim="output,0" global="1,Y,F" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<!-- Single work group kernel for not embarrassingly-parallel use-case -->
+<CustomLayer name="CTCGreedyDecoder" type="MVCL" version="1" max-shaves="1">
+    <Kernel entry="CTCDecoder">
+        <Source filename="ctc.bin"/>
+        <Parameters>
+            <Tensor arg-name="probabilities" type="input" port-index="0" format="FYX"/>
+            <Tensor arg-name="sequence_indicators" type="input" port-index="1" format="BF"/>
+            <Tensor arg-name="output" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="width" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="height" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="channels" type="int" port-index="0" source="I.F"/>
+        </Parameters>
+        <WorkSizes dim="output,0" global="1,1,1" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="ShuffleChannel" type="MVCL" version="1">
+    <!-- artificially added where closure for testing reasons, kernel itself supports arbitrary grouping -->
+    <!-- <Where group="2"/> -->
+    <Kernel entry="ShuffleChannel">
+        <Source filename="shuffle_channels.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="C" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="G" type="int" source="group"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="F,1,1" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<!-- Reference version of generic quantize layer, should be changed to FakeQuantize-->
+<CustomLayer name="FakeQuantize" type="MVCL" version="1">
+    <Kernel entry="quantize">
+        <Source filename="fakequantize.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="input_low" type="input" port-index="1" format="ANY"/>
+            <Tensor arg-name="input_high" type="input" port-index="2" format="ANY"/>
+            <Tensor arg-name="output_low" type="input" port-index="3" format="ANY"/>
+            <Tensor arg-name="output_high" type="input" port-index="4" format="ANY"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="levels" type="int" source="levels"/>
+            <Scalar arg-name="input_low_size" type="int" port-index="1" source="I.F"/>
+            <Scalar arg-name="input_high_size" type="int" port-index="2" source="I.F"/>
+            <Scalar arg-name="output_low_size" type="int" port-index="3" source="I.F"/>
+            <Scalar arg-name="output_high_size" type="int" port-index="4" source="I.F"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="1,Y,F" local="1,Y,1"/>
+    </Kernel>
+</CustomLayer>
+
+<!-- Reference version of generic quantize layer, should be changed to FakeQuantize-->
+<CustomLayer name="FakeQuantizeBin" type="MVCL" version="1">
+    <Where levels="2"/>
+    <Kernel entry="binarization">
+        <Source filename="binarization.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="input_low_high" type="input" port-index="1" format="BFYX"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="switch_out" type="int" source="switch_out"/>
+            <Scalar arg-name="input_low_high_size" type="int" source="input_low_size"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="H" type="int" port-index="0" source="I.Y"/>
+            <Tensor arg-name="input_high" type="input" port-index="2" format="BFYX"/>
+            <Tensor arg-name="output_low" type="input" port-index="3" format="BFYX"/>
+            <Tensor arg-name="output_high" type="input" port-index="4" format="BFYX"/>
+            <Scalar arg-name="input_high_size" type="int" source="input_high_size"/>
+            <Scalar arg-name="output_low_size" type="int" source="output_low_size"/>
+            <Scalar arg-name="output_high_size" type="int" source="output_high_size"/>
+            <Data arg-name="src_local" type="local_data" dim="input,0" size="X*Y*2"/>
+            <Data arg-name="dst_local" type="local_data" dim="input,0" size="X*Y*2"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="1,1,F" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="BinaryConvolution" type="MVCL" version="1">
+    <Where kernel="3,3"/>
+    <Kernel entry="binary_convolution">
+        <Source filename="binary_convolution3x3.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Data arg-name="weights_data" type="data" source="weights" format="ANY"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="pad_value" type="float" source="pad_value"/>
+            <Scalar arg-name="IW" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="IH" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="IC" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="DW" type="int" port-index="0" source="dilations"/>
+            <Scalar arg-name="DH" type="int" port-index="1" source="dilations"/>
+            <Scalar arg-name="GC" type="int" source="group"/>
+            <Scalar arg-name="KW" type="int" port-index="0" source="kernel"/>
+            <Scalar arg-name="KH" type="int" port-index="1" source="kernel"/>
+            <Scalar arg-name="PW" type="int" port-index="0" source="pads_begin"/>
+            <Scalar arg-name="PH" type="int" port-index="1" source="pads_begin"/>
+            <Scalar arg-name="SW" type="int" port-index="0" source="strides"/>
+            <Scalar arg-name="SH" type="int" port-index="1" source="strides"/>
+            <Scalar arg-name="OW" type="int" port-index="0" source="O.X"/>
+        </Parameters>
+        <WorkSizes dim="output,0" global="Y,F,1" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="BinaryConvolution" type="MVCL" version="1">
+    <Where kernel="1,1"/>
+    <Kernel entry="binary_convolution">
+        <Source filename="binary_convolution1x1.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Data arg-name="weights_data" type="data" source="weights" format="ANY"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="pad_value" type="float" source="pad_value"/>
+            <Scalar arg-name="IW" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="IH" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="IC" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="DW" type="int" port-index="0" source="dilations"/>
+            <Scalar arg-name="DH" type="int" port-index="1" source="dilations"/>
+            <Scalar arg-name="GC" type="int" source="group"/>
+            <Scalar arg-name="KW" type="int" port-index="0" source="kernel"/>
+            <Scalar arg-name="KH" type="int" port-index="1" source="kernel"/>
+            <Scalar arg-name="PW" type="int" port-index="0" source="pads_begin"/>
+            <Scalar arg-name="PH" type="int" port-index="1" source="pads_begin"/>
+            <Scalar arg-name="SW" type="int" port-index="0" source="strides"/>
+            <Scalar arg-name="SH" type="int" port-index="1" source="strides"/>
+            <Scalar arg-name="OW" type="int" port-index="0" source="O.X"/>
+        </Parameters>
+        <WorkSizes dim="output,0" global="Y,F,1" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<!-- Reference version of generic quantize binary convolution -->
+<!-- An example of a kernel binding that uses data blob from IR -->
+<CustomLayer name="BinaryConvolution" type="MVCL" version="1">
+    <Kernel entry="binary_convolution">
+        <Source filename="binary_convolution.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Data arg-name="weights_data" type="data" source="weights" format="ANY"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="pad_value" type="float" source="pad_value"/>
+            <Scalar arg-name="IW" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="IH" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="IC" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="DW" type="int" port-index="0" source="dilations"/>
+            <Scalar arg-name="DH" type="int" port-index="1" source="dilations"/>
+            <Scalar arg-name="GC" type="int" source="group"/>
+            <Scalar arg-name="KW" type="int" port-index="0" source="kernel"/>
+            <Scalar arg-name="KH" type="int" port-index="1" source="kernel"/>
+            <Scalar arg-name="PW" type="int" port-index="0" source="pads_begin"/>
+            <Scalar arg-name="PH" type="int" port-index="1" source="pads_begin"/>
+            <Scalar arg-name="SW" type="int" port-index="0" source="strides"/>
+            <Scalar arg-name="SH" type="int" port-index="1" source="strides"/>
+        </Parameters>
+        <WorkSizes dim="output,0" global="X,Y,F" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="Resample" type="MVCL" version="1">
+    <Where antialias="0"/>
+    <Kernel entry="resample_nearest">
+        <Source filename="resample_noAA.bin"/>
+        <Parameters>
+            <Tensor arg-name="src" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="iw" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="ih" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="factor" type="float" source="factor"/>
+            <Scalar arg-name="ow" type="int" port-index="0" source="O.X"/>
+            <Scalar arg-name="oh" type="int" port-index="0" source="O.Y"/>
+            <Scalar arg-name="channels" type="int" port-index="0" source="I.F"/>
+        </Parameters>
+        <WorkSizes global="1,Y,1" local="1,1,1" dim="output,0"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="Resample" type="MVCL" version="1">
+    <Where antialias="1"/>
+    <Kernel entry="resample_with_antialias">
+        <Source filename="resample_AA.bin"/>
+        <Parameters>
+            <Tensor arg-name="src" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="iw" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="ih" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="factor" type="float" source="factor"/>
+            <Scalar arg-name="ow" type="int" port-index="0" source="O.X"/>
+            <Scalar arg-name="oh" type="int" port-index="0" source="O.Y"/>
+            <Scalar arg-name="channels" type="int" port-index="0" source="I.F"/>
+        </Parameters>
+        <WorkSizes global="1,round(Y*factor),F" local="1,1,F" dim="input,0"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="Convolution" type="MVCL" version="1">
+    <Where kernel="1,1" dilation="1,1"/>
+    <Kernel entry="Convolution1x1_NCHW">
+       <Source filename="convolution1x1_chw.bin"/>
+       <Parameters>
+           <Tensor arg-name="in" type="input" port-index="0" format="BFYX"/>
+           <Tensor arg-name="out" type="output" port-index="0" format="BFYX"/>
+           <Data arg-name="w" type="data" source="weights" format="ANY"/>
+           <Scalar arg-name="IW" type="int" port-index="0" source="I.X"/>
+           <Scalar arg-name="IH" type="int" port-index="0" source="I.Y"/>
+           <Scalar arg-name="IC" type="int" port-index="0" source="I.F"/>
+           <Scalar arg-name="OW" type="int" port-index="0" source="O.X"/>
+           <Scalar arg-name="OH" type="int" port-index="0" source="O.Y"/>
+           <Scalar arg-name="OC" type="int" port-index="0" source="O.F"/>
+
+           <Scalar arg-name="stride-x" type="int" port-index="0" source="stride"/>
+           <Scalar arg-name="stride-y" type="int" port-index="1" source="stride"/>
+           <Scalar arg-name="pad-x" type="int" port-index="0" source="pads_begin"/>
+           <Scalar arg-name="pad-y" type="int" port-index="1" source="pads_begin"/>
+           <Scalar arg-name="kernel-x" type="int" port-index="0" source="kernel"/>
+           <Scalar arg-name="kernel-y" type="int" port-index="1" source="kernel"/>
+           <Scalar arg-name="output" type="int" port-index="0" source="output"/>
+           <Scalar arg-name="group" type="int" port-index="0" source="group"/>
+       </Parameters>
+       <WorkSizes global="Y,F,B" local="1,1,1" dim="output,0"/>
+   </Kernel>
+</CustomLayer>
+
+<CustomLayer name="Convolution" type="MVCL" version="1">
+    <Where kernel="1,1" dilation="1,1"/>
+    <Kernel entry="Convolution1x1_NHWC">
+       <Source filename="convolution1x1_hwc.bin"/>
+       <Parameters>
+           <Tensor arg-name="in" type="input" port-index="0" format="BYXF"/>
+           <Tensor arg-name="out" type="output" port-index="0" format="BFYX"/>
+           <Data arg-name="w" type="data" source="weights" format="ANY"/>
+           <Scalar arg-name="IW" type="int" port-index="0" source="I.X"/>
+           <Scalar arg-name="IH" type="int" port-index="0" source="I.Y"/>
+           <Scalar arg-name="IC" type="int" port-index="0" source="I.F"/>
+           <Scalar arg-name="OW" type="int" port-index="0" source="O.X"/>
+           <Scalar arg-name="OH" type="int" port-index="0" source="O.Y"/>
+           <Scalar arg-name="OC" type="int" port-index="0" source="O.F"/>
+
+           <Scalar arg-name="stride-x" type="int" port-index="0" source="stride"/>
+           <Scalar arg-name="stride-y" type="int" port-index="1" source="stride"/>
+           <Scalar arg-name="pad-x" type="int" port-index="0" source="pads_begin"/>
+           <Scalar arg-name="pad-y" type="int" port-index="1" source="pads_begin"/>
+           <Scalar arg-name="kernel-x" type="int" port-index="0" source="kernel"/>
+           <Scalar arg-name="kernel-y" type="int" port-index="1" source="kernel"/>
+           <Scalar arg-name="output" type="int" port-index="0" source="output"/>
+           <Scalar arg-name="group" type="int" port-index="0" source="group"/>
+       </Parameters>
+       <WorkSizes global="Y,F,B" local="1,1,1" dim="output,0"/>
+   </Kernel>
+</CustomLayer>
+
+<CustomLayer name="Convolution" type="MVCL" version="1">
+    <Where kernel="3,3" dilation="1,1"/>
+    <Kernel entry="Convolution3x3">
+        <Source filename="convolution3x3.bin"/>
+        <Parameters>
+            <Tensor arg-name="in_param" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="out" type="output" port-index="0" format="BFYX"/>
+            <Data arg-name="w" type="data" source="weights" format="BFYX"/>
+            <Scalar arg-name="IW" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="IH" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="IC" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="OW" type="int" port-index="0" source="O.X"/>
+            <Scalar arg-name="OH" type="int" port-index="0" source="O.Y"/>
+            <Scalar arg-name="OC" type="int" port-index="0" source="O.F"/>
+            <Scalar arg-name="KX" type="int" port-index="0" source="kernel"/>
+            <Scalar arg-name="KY" type="int" port-index="1" source="kernel"/>
+            <Scalar arg-name="stride_x" type="int" port-index="0" source="stride"/>
+            <Scalar arg-name="stride_y" type="int" port-index="1" source="stride"/>
+            <Scalar arg-name="pad_x" type="int" port-index="0" source="pads_begin"/>
+            <Scalar arg-name="pad_y" type="int" port-index="1" source="pads_begin"/>
+            <Scalar arg-name="dilation_x" type="int" port-index="0" source="dilation"/>
+            <Scalar arg-name="dilation_y" type="int" port-index="1" source="dilation"/>
+            <Scalar arg-name="output" type="int" port-index="0" source="output"/>
+
+            <Data arg-name="in_local" type="local_data" dim="input,0" size="X*F*3*2"/>
+            <Data arg-name="out_local" type="local_data" dim="output,0" size="X*F*2"/>
+            <Data arg-name="w_local" type="local_data" dim="input,0" size="3*3*F*2"/>
+        </Parameters>
+        <WorkSizes global="Y,F,B" local="1,1,1" dim="output,0"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="ExperimentalDetectronPriorGridGenerator" type="MVCL" version="1">
+    <Kernel entry="experimental_detectron_prior_grid_generator">
+        <Source filename="detectron_prior_grid_gen.bin"/>
+        <Parameters>
+            <Tensor arg-name="input_priors" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="input_feature_map" type="input" port-index="1" format="BFYX"/>
+            <Tensor arg-name="input_rois" type="input" port-index="2" format="BFYX"/>
+            <Tensor arg-name="output" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="grid_h" type="int" port-index="1" source="I.Y"/>
+            <Scalar arg-name="grid_w" type="int" port-index="1" source="I.X"/>
+            <Scalar arg-name="stride_h" type="float" source="stride_h"/>
+            <Scalar arg-name="stride_w" type="float" source="stride_w"/>
+            <Scalar arg-name="num_priors" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="num_anchors_per_prior" type="int" port-index="0" source="I.X"/>
+        </Parameters>
+        <WorkSizes dim="input,1" global="((X+31)/32)*32,Y,1" local="32,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="Convert" type="MVCL" version="1">
+    <Kernel entry="cvtu8f16">
+        <Source filename="cvtu8f16.bin"/>
+        <Parameters>
+            <Tensor arg-name="src" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="dst" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="scale" type="float" source="scale"/>
+            <Scalar arg-name="bias" type="float" source="bias"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="X,Y,F" local="X,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="Correlate" type="MVCL" version="1">
+    <Kernel entry="correlate2_half">
+        <Source filename="correlate.bin"/>
+        <Parameters>
+            <Tensor arg-name="bottom0" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="bottom1" type="input" port-index="1" format="BFYX"/>
+            <Tensor arg-name="top" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="topwidth" type="int" source="top_width"/>
+            <Scalar arg-name="topheight" type="int" source="top_height"/>
+            <Scalar arg-name="bottomwidth" type="int" port-index="0" source="I.X"/>
+            <Scalar arg-name="bottomheight" type="int" port-index="0" source="I.Y"/>
+            <Scalar arg-name="bottomchannels" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="max_displacement" type="int" source="displacement"/>
+            <Scalar arg-name="padding" type="int" source="pad"/>
+            <Scalar arg-name="neighborhood_grid_radius" type="int" source="neighborhood_grid_radius"/>
+            <Scalar arg-name="neighborhood_grid_width" type="int" source="neighborhood_grid_width"/>
+            <Scalar arg-name="kernel_size" type="int" source="kernel_size"/>
+            <Scalar arg-name="stride1" type="int" port-index="0" source="stride"/>
+            <Scalar arg-name="stride2" type="int" port-index="1" source="stride"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="top_height,1,1" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>
+
+<CustomLayer name="SpatialTransform" type="MVCL" version="1">
+    <Kernel entry="ocl_st">
+        <Source filename="st.bin"/>
+        <Parameters>
+            <Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
+            <Tensor arg-name="theta" type="input" port-index="1" format="ANY"/>
+            <Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
+            <Scalar arg-name="C" type="int" port-index="0" source="I.F"/>
+            <Scalar arg-name="W" type="int" port-index="0" source="I.X"/>
+        </Parameters>
+        <WorkSizes dim="input,0" global="(X+511)/512,Y,1" local="1,1,1"/>
+    </Kernel>
+</CustomLayer>

openvino/vpu_custom_kernels/cvtf32f16.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:703ef56f84299e76d36b3ba5a632ae3d5e3ecd54761dcfe0006ca69ddce4bc6d
+size 2664

openvino/vpu_custom_kernels/cvtf32f16.cl

@@ -0,0 +1,17 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+__kernel void cvtf32f16(const __global float* restrict inImage,
+                              __global half*  restrict outImage,
+                                       float   scale,
+                                       float   bais)
+{
+    int idx = get_global_id(0)
+            + get_global_id(1) * get_global_size(0)
+            + get_global_id(2) * get_global_size(0) * get_global_size(1);
+
+    outImage[idx] = convert_half(inImage[idx]*scale+bais);
+}

openvino/vpu_custom_kernels/cvtu8f16.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:46c943e08f37cedac77f727f55835637d4878edcc20aaa24f16ed5888d13bd43
+size 4588

openvino/vpu_custom_kernels/cvtu8f16.cl

@@ -0,0 +1,48 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__kernel void cvtu8f16(__global const uchar *restrict src, __global half *restrict dst, float scale, float bias)
+{
+    __local uchar local_src[8 * 1024];
+    __local half local_dst[8 * 1024];
+
+    event_t e1 = async_work_group_copy_3D3D(
+        local_src, // dst
+        src + get_group_id(0) * get_local_size(0) + get_group_id(1) * get_local_size(1) * get_global_size(0)
+            + get_group_id(2) * get_local_size(2) * get_global_size(0) * get_global_size(1), // src
+        get_local_size(0), // num_elements_per_line
+        get_local_size(0) * get_local_size(1) / (get_local_size(0)), // num_lines
+        get_global_size(0) - get_local_size(0), // src_line_stride
+        0, // dst_line_stride
+        get_local_size(2), // num planes
+        get_global_size(0) * (get_global_size(1) - get_local_size(1)), // src plane stride
+        0, // dst plane stride
+        0);
+    wait_group_events(1, &e1);
+
+    size_t idx = get_local_id(0)
+               + get_local_id(1) * get_local_size(0)
+               + get_local_id(2) * get_local_size(0) * get_local_size(1);
+
+    local_dst[idx] = convert_half(local_src[idx]) * (half)scale + (half)bias;
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy_3D3D(
+        dst + get_group_id(0) * get_local_size(0) + get_group_id(1) * get_local_size(1) * get_global_size(0)
+            + get_group_id(2) * get_local_size(2) * get_global_size(0) * get_global_size(1), // dst
+        local_dst, // src
+        get_local_size(0), // num_elements_per_line
+        get_local_size(1), // num_lines
+        0, // src_line_stride
+        get_global_size(0) - get_local_size(0), // dst_line_stride
+        get_local_size(2), // num_planes
+        0, // src_plane_stride
+        get_global_size(0) * (get_global_size(1) - get_local_size(1)), // dst_plane_stride
+        0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/detectron_prior_grid_gen.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:4c80d556d23f1c959fa10c00ff1cd9c3ae10aba607b37c7a0620d903fc7cedd8
+size 6972

openvino/vpu_custom_kernels/detectron_prior_grid_gen.cl

@@ -0,0 +1,65 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__kernel void experimental_detectron_prior_grid_generator(
+    __global const half *restrict input_priors,
+    __global const half *restrict input_feature_map,
+    __global const half *restrict input_rois,
+    __global half *restrict output,
+    int grid_h,
+    int grid_w,
+    float stride_h,
+    float stride_w,
+    int num_priors,
+    int num_anchors_per_prior)
+{
+    __local half local_input_priors[8 * 1024];
+    __local half local_output[8 * 1024];
+
+    event_t e1 = async_work_group_copy(
+        local_input_priors,
+        input_priors,
+        num_anchors_per_prior * num_priors,
+        0);
+    wait_group_events(1, &e1);
+
+    int width_start = get_group_id(0) * get_local_size(0);
+    int width_end   = min(width_start + get_local_size(0), (unsigned)grid_w);
+    int width       = width_end - width_start;
+
+    int h     = get_group_id(1);
+    int w_idx = get_group_id(0) * get_local_size(0);
+    for (int w = 0; w < width; ++w) {
+        #pragma unroll 4
+        for (int p = 0; p < num_priors; ++p) {
+            local_output[(w * num_priors + p) * num_anchors_per_prior + 0] =
+                local_input_priors[4 * p + 0]
+                + convert_half(stride_w) * (convert_half(w_idx + w) + 0.5);
+            local_output[(w * num_priors + p) * num_anchors_per_prior + 1] =
+                local_input_priors[4 * p + 1] + convert_half(stride_h) * (convert_half(h) + 0.5);
+            local_output[(w * num_priors + p) * num_anchors_per_prior + 2] =
+                local_input_priors[4 * p + 2]
+                + convert_half(stride_w) * (convert_half(w_idx + w) + 0.5);
+            local_output[(w * num_priors + p) * num_anchors_per_prior + 3] =
+                local_input_priors[4 * p + 3] + convert_half(stride_h) * (convert_half(h) + 0.5);
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy_2D2D(
+        output + get_group_id(0) * get_local_size(0) * num_anchors_per_prior * num_priors
+            + get_group_id(1) * get_local_size(1) * grid_w * num_anchors_per_prior
+                  * num_priors, // dst
+        local_output, // src
+        width * num_anchors_per_prior * num_priors, // num_elements_per_line
+        1, // num_lines
+        (grid_w - width) * num_anchors_per_prior * num_priors, // src_line_stride
+        (grid_w - width) * num_anchors_per_prior * num_priors, // dst_line_stride
+        0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/fakequantize.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:d17659bbf12a172849085003a055bfb4b91d3bb5bdc7f820395820eaa90b46ef
+size 15688

openvino/vpu_custom_kernels/fakequantize.cl

@@ -0,0 +1,111 @@
+// Copyright (C) 2018-2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+//
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
+
+__kernel void quantize(
+    __global const half *restrict src_data,
+    __global const half *restrict input_low,
+    __global const half *restrict input_high,
+    __global const half *restrict output_low,
+    __global const half *restrict output_high,
+    __global half *restrict dst_data,
+    int levels,
+    int input_low_size,
+    int input_high_size,
+    int output_low_size,
+    int output_high_size,
+    int W,
+    int H)
+{
+    __local half local_src[15 * 1024];
+    __local half local_dst[15 * 1024];
+
+    event_t e1 = async_work_group_copy(local_src, src_data + get_group_id(2) * W * H, W * H, 0);
+    wait_group_events(1, &e1);
+
+    int c = get_group_id(2);
+
+    half h_ilow  = (input_low_size == 1 ? input_low[0] : input_low[c]);
+    half h_ihigh = (input_high_size == 1 ? input_high[0] : input_high[c]);
+    half h_olow  = (output_low_size == 1 ? output_low[0] : output_low[c]);
+    half h_ohigh = (output_high_size == 1 ? output_high[0] : output_high[c]);
+
+    half const1 = (half)(
+        !(h_ihigh - h_ilow) ? 0.0f : convert_float(levels - 1) / (convert_float(h_ihigh) - convert_float(h_ilow)));
+    half const2 =
+        (half)(!(levels - 1) ? 0.0f : (convert_float(h_ohigh) - convert_float(h_olow)) / convert_float(levels - 1));
+
+    __local const half *restrict src = local_src + W * get_local_id(1);
+    __local half       *restrict dst = local_dst + W * get_local_id(1);
+
+    for (int w = 0; w < W / 8; w++) {
+        half8 val = *((__local half8 *)src + w);
+        half8 aux = (val - (half8)h_ilow) * (half8)const1 + (half8)0.5h;
+
+        aux = (half8){
+            (half)(short)(aux.s0),
+            (half)(short)(aux.s1),
+            (half)(short)(aux.s2),
+            (half)(short)(aux.s3),
+            (half)(short)(aux.s4),
+            (half)(short)(aux.s5),
+            (half)(short)(aux.s6),
+            (half)(short)(aux.s7)};
+
+        aux = aux * (half8)const2 + (half8)h_olow;
+
+        short8 a;
+        short8 b;
+        a.s0 = (val.s0 <= h_ilow);
+        a.s1 = (val.s1 <= h_ilow);
+        a.s2 = (val.s2 <= h_ilow);
+        a.s3 = (val.s3 <= h_ilow);
+        a.s4 = (val.s4 <= h_ilow);
+        a.s5 = (val.s5 <= h_ilow);
+        a.s6 = (val.s6 <= h_ilow);
+        a.s7 = (val.s7 <= h_ilow);
+
+        b.s0 = (val.s0 > h_ihigh);
+        b.s1 = (val.s1 > h_ihigh);
+        b.s2 = (val.s2 > h_ihigh);
+        b.s3 = (val.s3 > h_ihigh);
+        b.s4 = (val.s4 > h_ihigh);
+        b.s5 = (val.s5 > h_ihigh);
+        b.s6 = (val.s6 > h_ihigh);
+        b.s7 = (val.s7 > h_ihigh);
+
+        a = ~(a - (short8)1);
+        b = ~(b - (short8)1);
+
+        short8 c1 = (~a & b);
+        short8 c2 = (~a & ~b);
+
+        short8 res = (a & as_short8((half8)h_olow)) | (c1 & as_short8((half8)h_ohigh)) | (c2 & as_short8(aux));
+
+        *((__local half8 *)dst + w) = as_half8(res);
+    }
+
+    for (int w = W & (~0x7); w < W; w++) {
+        half val = src[w];
+        short a  = val <= h_ilow;
+        a        = ~(a - 1);
+        short b  = val > h_ihigh;
+        b        = ~(b - 1);
+
+        short c1 = (~a & b);
+        short c2 = (~a & ~b);
+
+        short res = (a & as_short(h_olow)) | (c1 & as_short(h_ohigh))
+                    | (c2 & as_short(((half)(round((val - h_ilow) * const1) * const2) + h_olow)));
+
+        dst[w] = as_half(res);
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    event_t e2 = async_work_group_copy(dst_data + get_group_id(2) * W * H, local_dst, W * H, 0);
+    wait_group_events(1, &e2);
+}

openvino/vpu_custom_kernels/grn.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:6e3dbe5173ca93f39fecaf29f820e1704bcb485affc1a09554e4c86f8de46214
+size 7972