face match detector

Dockerfile

@@ -1,19 +0,0 @@
-FROM ubuntu:20.04
-RUN ln -snf /usr/share/zoneinfo/$CONTAINER_TIMEZONE /etc/localtime && echo $CONTAINER_TIMEZONE > /etc/timezone
-RUN apt-get update -y
-RUN apt-get install -y python3 python3-pip python3-opencv
-RUN apt-get install -y libcurl4-openssl-dev libssl-dev
-RUN mkdir -p /home/FaceOnLive_v6
-RUN mkdir -p /home/FaceOnLive_v6/facewrapper
-WORKDIR /home/FaceOnLive_v6
-COPY ./facewrapper ./facewrapper
-COPY ./facewrapper/libs/libimutils.so /usr/lib
-COPY ./gradio ./gradio
-COPY ./openvino /usr/lib
-COPY ./app.py ./app.py
-COPY ./run.sh .
-COPY ./requirements.txt ./requirements.txt
-RUN pip3 install -r requirements.txt
-RUN chmod a+x run.sh
-CMD ["./run.sh"]
-EXPOSE 8000

README.md

@@ -1,11 +1,12 @@
 ---
-title: Face Recognition SDK
-emoji: 🤨🙂
+title: Face Matching Tool
+emoji: 🤡
 colorFrom: yellow
-colorTo: pink
-sdk: docker
+colorTo: red
+sdk: gradio
+sdk_version: 4.36.1
+app_file: app.py
 pinned: false
-license: mit
 short_description: FaceOnLive On-Premise Solution
 ---
 

app.py

@@ -1,217 +1,85 @@
-import sys
-sys.path.append('.')
-
-from flask import Flask, request, jsonify
-from time import gmtime, strftime
-import os
-import base64
-import json
 import cv2
-import numpy as np
+import face_recognition
+import gradio as gr
+import datetime as DT
+import pytz
 
-from facewrapper.facewrapper import ttv_version
-from facewrapper.facewrapper import ttv_get_hwid
-from facewrapper.facewrapper import ttv_init
-from facewrapper.facewrapper import ttv_init_offline
-from facewrapper.facewrapper import ttv_extract_feature
-from facewrapper.facewrapper import ttv_compare_feature
+ipAddress = None
 
-app = Flask(__name__) 
 
-app.config['SITE'] = "http://0.0.0.0:8000/"
-app.config['DEBUG'] = False
+def __nowInIST():
+    return DT.datetime.now(pytz.timezone("Asia/Kolkata"))
 
-licenseKey = os.environ.get("LICENSE_KEY")
-licensePath = "license.txt"
-modelFolder = os.path.abspath(os.path.dirname(__file__)) + '/facewrapper/dict'
 
-version = ttv_version()
-print("version: ", version.decode('utf-8'))
+def __attachIp(request: gr.Request):
+    global ipAddress
+    x_forwarded_for = request.headers.get('x-forwarded-for')
+    if x_forwarded_for:
+        ipAddress = x_forwarded_for
 
-ret = ttv_init(modelFolder.encode('utf-8'), licenseKey.encode('utf-8'))
-if ret != 0:
-    print(f"online init failed: {ret}");
 
-    hwid = ttv_get_hwid()
-    print("hwid: ", hwid.decode('utf-8'))
+def pprint(log: str):
+    now = __nowInIST()
+    now = now.strftime("%Y-%m-%d %H:%M:%S")
+    print(f"[{now}] [{ipAddress}] {log}")
 
-    ret = ttv_init_offline(modelFolder.encode('utf-8'), licensePath.encode('utf-8'))
-    if ret != 0:
-        print(f"offline init failed: {ret}")
-        exit(-1)
-    else:
-        print(f"offline init ok")
-
-else:
-    print(f"online init ok")
-
-@app.route('/api/compare_face', methods=['POST'])
-def compare_face():
-    file1 = request.files['image1']
-    image1 = cv2.imdecode(np.fromstring(file1.read(), np.uint8), cv2.IMREAD_COLOR)
-    if image1 is None:
-        result = "image1: is null!"
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-
-    file2 = request.files['image2']
-    image2 = cv2.imdecode(np.fromstring(file2.read(), np.uint8), cv2.IMREAD_COLOR)
-    if image2 is None:
-        result = "image2: is null!"
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-    
-    faceRect1 = np.zeros([4], dtype=np.int32)
-    feature1 = np.zeros([2048], dtype=np.uint8)
-    featureSize1 = np.zeros([1], dtype=np.int32)
-
-    ret = ttv_extract_feature(image1, image1.shape[1], image1.shape[0], faceRect1, feature1, featureSize1)
-    if ret <= 0:
-        if ret == -1:
-            result = "license error!"
-        elif ret == -2:
-            result = "init error!"
-        elif ret == 0:
-            result = "image1: no face detected!"
-
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-
-    faceRect2 = np.zeros([4], dtype=np.int32)
-    feature2 = np.zeros([2048], dtype=np.uint8)
-    featureSize2 = np.zeros([1], dtype=np.int32)
-
-    ret = ttv_extract_feature(image2, image2.shape[1], image2.shape[0], faceRect2, feature2, featureSize2)
-    if ret <= 0:
-        if ret == -1:
-            result = "license error!"
-        elif ret == -2:
-            result = "init error!"
-        elif ret == 0:
-            result = "image2: no face detected!"
-
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-
-    similarity = ttv_compare_feature(feature1, feature2)
-    if similarity > 0.7:
-        result = "same"
-    else:
-        result = "different"
-  
-    status = "ok"
-    response = jsonify(
-    {
-        "status": status, 
-        "data": {
-            "result": result, 
-            "similarity": float(similarity), 
-            "face1": {"x1": int(faceRect1[0]), "y1": int(faceRect1[1]), "x2": int(faceRect1[2]), "y2" : int(faceRect1[3])}, 
-            "face2": {"x1": int(faceRect2[0]), "y1": int(faceRect2[1]), "x2": int(faceRect2[2]), "y2" : int(faceRect2[3])}, 
-            }
-    })
-
-    response.status_code = 200
-    response.headers["Content-Type"] = "application/json; charset=utf-8"
-    return response
-
-
-@app.route('/api/compare_face_base64', methods=['POST'])
-def coompare_face_base64():
-    content = request.get_json()
-    imageBase641 = content['image1']
-    image1 = cv2.imdecode(np.frombuffer(base64.b64decode(imageBase641), dtype=np.uint8), cv2.IMREAD_COLOR)
-
-    if image1 is None:
-        result = "image1: is null!"
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-
-    imageBase642 = content['image2']
-    image2 = cv2.imdecode(np.frombuffer(base64.b64decode(imageBase642), dtype=np.uint8), cv2.IMREAD_COLOR)
-
-    if image2 is None:
-        result = "image2: is null!"
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-
-    faceRect1 = np.zeros([4], dtype=np.int32)
-    feature1 = np.zeros([2048], dtype=np.uint8)
-    featureSize1 = np.zeros([1], dtype=np.int32)
-
-    ret = ttv_extract_feature(image1, image1.shape[1], image1.shape[0], faceRect1, feature1, featureSize1)
-    if ret <= 0:
-        if ret == -1:
-            result = "license error!"
-        elif ret == -2:
-            result = "init error!"
-        elif ret == 0:
-            result = "image1: no face detected!"
-
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-
-    faceRect2 = np.zeros([4], dtype=np.int32)
-    feature2 = np.zeros([2048], dtype=np.uint8)
-    featureSize2 = np.zeros([1], dtype=np.int32)
-
-    ret = ttv_extract_feature(image2, image2.shape[1], image2.shape[0], faceRect2, feature2, featureSize2)
-    if ret <= 0:
-        if ret == -1:
-            result = "license error!"
-        elif ret == -2:
-            result = "init error!"
-        elif ret == 0:
-            result = "image2: no face detected!"
-
-        status = "ok"
-        response = jsonify({"status": status, "data": {"result": result}})
-        response.status_code = 200
-        response.headers["Content-Type"] = "application/json; charset=utf-8"
-        return response
-
-    similarity = ttv_compare_feature(feature1, feature2)
-    if similarity > 0.7:
-        result = "same"
+
+def __findFaceEncodings(imagePath):
+    image = cv2.imread(imagePath)
+    faceEncodings = face_recognition.face_encodings(image)
+    return faceEncodings[0] if len(faceEncodings) > 0 else None
+
+
+def predictMatch(firstImage, secondImage):
+    image1Encoding = __findFaceEncodings(firstImage)
+    image2Encoding = __findFaceEncodings(secondImage)
+
+    pprint("Starting the job...")
+    distance = face_recognition.face_distance([image1Encoding], image2Encoding)
+    distancePercent = round(distance[0] * 100)
+    matchPercent = 100 - distancePercent
+    pprint(f"Job finished. Match    : {matchPercent}%")
+
+    isSame = matchPercent > 50
+
+    amplificationFactor = 1.5
+    adjustment = (amplificationFactor - 0.5) * 100
+
+    outputTexts = []
+    if isSame:
+        outputTexts.append("The images are of the same person ✅")
+        matchScore = round((matchPercent + adjustment) / (100 + adjustment) * 100)
+        outputTexts.append(f"\nMatch Score: {matchScore}%")
     else:
-        result = "different"
-  
-    status = "ok"
-    response = jsonify(
-    {
-        "status": status, 
-        "data": {
-            "result": result, 
-            "similarity": float(similarity), 
-            "face1": {"x1": int(faceRect1[0]), "y1": int(faceRect1[1]), "x2": int(faceRect1[2]), "y2" : int(faceRect1[3])}, 
-            "face2": {"x1": int(faceRect2[0]), "y1": int(faceRect2[1]), "x2": int(faceRect2[2]), "y2" : int(faceRect2[3])}, 
-            }
-    })
-    response.status_code = 200
-    response.headers["Content-Type"] = "application/json; charset=utf-8"
-    return response
-
-if __name__ == '__main__':
-    port = int(os.environ.get("PORT", 8000))
-    app.run(host='0.0.0.0', port=port)
+        outputTexts.append("The images are not of the same person ❌")
+
+    outputText = "\n".join(outputTexts)
+    pprint(f"{outputText=}")
+
+    return outputText
+
+
+with gr.Row(elem_classes=["main-container"]):
+    with gr.Row(elem_classes=["img-container"]):
+        firstImage = gr.Image(type='filepath', height=250, elem_classes=["image"])
+        secondImage = gr.Image(type='filepath', height=250, elem_classes=["image"])
+
+    with gr.Row(elem_classes=["output-container"]):
+        result = gr.Textbox(label="Result", elem_classes=["output"])
+
+
+with gr.Interface(
+    fn=predictMatch,
+    inputs=[
+        firstImage,
+        secondImage,
+    ],
+    outputs=[
+        result
+    ],
+    title="Face Match Detector",
+    allow_flagging="never",
+) as demo:
+    demo.load(__attachIp, None, None)
+
+demo.launch(debug=True)

facewrapper/dict/data1.bin

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

facewrapper/dict/data2.bin

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

facewrapper/dict/data3.bin

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

facewrapper/dict/detect.bin

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

facewrapper/facewrapper.py

@@ -1,32 +0,0 @@
-import ctypes, ctypes.util
-from ctypes import *
-from numpy.ctypeslib import ndpointer
-import sys
-import os
-
-lib_path = os.path.abspath(os.path.dirname(__file__)) + '/libs/libttvfaceengine6.so'
-liveness_engine = cdll.LoadLibrary(lib_path)
-
-ttv_version = liveness_engine.ttv_version
-ttv_version.argtypes = []
-ttv_version.restype = ctypes.c_char_p
-
-ttv_get_hwid = liveness_engine.ttv_get_hwid
-ttv_get_hwid.argtypes = []
-ttv_get_hwid.restype = ctypes.c_char_p
-
-ttv_init = liveness_engine.ttv_init
-ttv_init.argtypes = [ctypes.c_char_p, ctypes.c_char_p]
-ttv_init.restype = ctypes.c_int32
-
-ttv_init_offline = liveness_engine.ttv_init_offline
-ttv_init_offline.argtypes = [ctypes.c_char_p, ctypes.c_char_p]
-ttv_init_offline.restype = ctypes.c_int32
-
-ttv_extract_feature = liveness_engine.ttv_extract_feature
-ttv_extract_feature.argtypes = [ndpointer(ctypes.c_ubyte, flags='C_CONTIGUOUS'), ctypes.c_int32, ctypes.c_int32, ndpointer(ctypes.c_int32, flags='C_CONTIGUOUS'), ndpointer(ctypes.c_ubyte, flags='C_CONTIGUOUS'), ndpointer(ctypes.c_int32, flags='C_CONTIGUOUS')]
-ttv_extract_feature.restype = ctypes.c_int
-
-ttv_compare_feature = liveness_engine.ttv_compare_feature
-ttv_compare_feature.argtypes = [ndpointer(ctypes.c_ubyte, flags='C_CONTIGUOUS'), ndpointer(ctypes.c_ubyte, flags='C_CONTIGUOUS')]
-ttv_compare_feature.restype = ctypes.c_double

facewrapper/libs/libttvfaceengine6.so

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

gradio/demo.py

@@ -1,115 +0,0 @@
-import gradio as gr
-import requests
-import json
-from PIL import Image
-
-def compare_face(frame1, frame2):
-    url = "http://127.0.0.1:8000/api/compare_face"
-    files = {'image1': open(frame1, 'rb'), 'image2': open(frame2, 'rb')}
-
-    r = requests.post(url=url, files=files)
-    faces = None
-
-    try:
-        image1 = Image.open(frame1)
-        image2 = Image.open(frame2)
-
-        face1 = None
-        face2 = None
-        data = r.json().get('data')
-        if data.get('face1') is not None:
-            face = data.get('face1')
-            x1 = face.get('x1')
-            y1 = face.get('y1')
-            x2 = face.get('x2')
-            y2 = face.get('y2')
-            if x1 < 0:
-                x1 = 0
-            if y1 < 0:
-                y1 = 0
-            if x2 >= image1.width:
-                x2 = image1.width - 1
-            if y2 >= image1.height:
-                y2 = image1.height - 1
-
-            face1 = image1.crop((x1, y1, x2, y2))
-            face_image_ratio = face1.width / float(face1.height)
-            resized_w = int(face_image_ratio * 150)
-            resized_h = 150
-
-            face1 = face1.resize((int(resized_w), int(resized_h)))
-
-        if data.get('face2') is not None:
-            face = data.get('face2')
-            x1 = face.get('x1')
-            y1 = face.get('y1')
-            x2 = face.get('x2')
-            y2 = face.get('y2')
-
-            if x1 < 0:
-                x1 = 0
-            if y1 < 0:
-                y1 = 0
-            if x2 >= image2.width:
-                x2 = image2.width - 1
-            if y2 >= image2.height:
-                y2 = image2.height - 1
-
-            face2 = image2.crop((x1, y1, x2, y2))
-            face_image_ratio = face2.width / float(face2.height)
-            resized_w = int(face_image_ratio * 150)
-            resized_h = 150
-
-            face2 = face2.resize((int(resized_w), int(resized_h)))
-        
-        if face1 is not None and face2 is not None:
-            new_image = Image.new('RGB',(face1.width + face2.width + 10, 150), (80,80,80))
-
-            new_image.paste(face1,(0,0))
-            new_image.paste(face2,(face1.width + 10, 0))
-            faces = new_image.copy()
-        elif face1 is not None and face2 is None:
-            new_image = Image.new('RGB',(face1.width + face1.width + 10, 150), (80,80,80))
-
-            new_image.paste(face1,(0,0))
-            faces = new_image.copy()
-        elif face1 is None and face2 is not None:
-            new_image = Image.new('RGB',(face2.width + face2.width + 10, 150), (80,80,80))
-
-            new_image.paste(face2,(face2.width + 10, 0))
-            faces = new_image.copy()
-    except:
-        pass
-
-    return [r.json(), faces]
-
-with gr.Blocks() as demo:
-    gr.Markdown(
-        """
-    # Face Recognition
-    Get your own Face Recognition Server by duplicating this space.<br/>
-        Or run on your own machine using docker.<br/>
-    ```docker run -it -p 7860:7860 --platform=linux/amd64 \
-	-e LICENSE_KEY="YOUR_VALUE_HERE" \
-	registry.hf.space/faceonlive-face-recognition-sdk:latest ```<br/><br/>
-    Contact us at https://faceonlive.com for issues and support.<br/>
-    """
-    )
-    with gr.Row():
-        with gr.Column():
-            compare_face_input1 = gr.Image(type='filepath', height=480)
-            gr.Examples(['gradio/examples/1.jpg', 'gradio/examples/2.jpg'], 
-                        inputs=compare_face_input1)
-            compare_face_button = gr.Button("Compare Face")
-        with gr.Column():
-            compare_face_input2 = gr.Image(type='filepath', height=480)
-            gr.Examples(['gradio/examples/3.jpg', 'gradio/examples/4.jpg'], 
-                        inputs=compare_face_input2)
-        with gr.Column():
-            compare_face_output = gr.Image(type="pil", height=300)
-            compare_result_output = gr.JSON(label='Result')
-
-    compare_face_button.click(compare_face, inputs=[compare_face_input1, compare_face_input2], outputs=[compare_result_output, compare_face_output])
-    gr.HTML('<a href="https://visitorbadge.io/status?path=https%3A%2F%2Fhuggingface.co%2Fspaces%2FFaceOnLive%2FFace-Recognition-SDK"><img src="https://api.visitorbadge.io/api/combined?path=https%3A%2F%2Fhuggingface.co%2Fspaces%2FFaceOnLive%2FFace-Recognition-SDK&labelColor=%23ff8a65&countColor=%2337d67a&style=flat&labelStyle=upper" /></a>')
-
-demo.launch(server_name="0.0.0.0", server_port=7860)

openvino/libgna.so

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

openvino/libgna.so.2

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

openvino/libgna.so.3.0.0.1455

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

openvino/libopenvino.so

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

openvino/libopenvino_gapi_preproc.so

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

openvino/libopenvino_intel_cpu_plugin.so

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

openvino/libopenvino_intel_gna_plugin.so

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

openvino/libopenvino_intel_hddl_plugin.so

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

openvino/libopenvino_intel_myriad_plugin.so

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

openvino/libopenvino_onnx_frontend.so

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

openvino/libopenvino_tensorflow_fe.so

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

openvino/pcie-ma2x8x.mvcmd

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

openvino/plugins.xml

@@ -1,27 +0,0 @@
-<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

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

openvino/vpu_custom_kernels/binarization.bin

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

openvino/vpu_custom_kernels/binarization.cl

@@ -1,67 +0,0 @@
-// 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

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

openvino/vpu_custom_kernels/binary_convolution.cl

@@ -1,95 +0,0 @@
-// 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

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

openvino/vpu_custom_kernels/binary_convolution1x1.cl

@@ -1,117 +0,0 @@
-// 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

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

openvino/vpu_custom_kernels/binary_convolution3x3.cl

@@ -1,278 +0,0 @@
-// 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

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

openvino/vpu_custom_kernels/convolution1x1_chw.cl

@@ -1,114 +0,0 @@
-// 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

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

openvino/vpu_custom_kernels/convolution1x1_hwc.cl

@@ -1,126 +0,0 @@
-// 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);
-}