2_Kernels_and_convolution.ipynb

{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# Lesson 2. Computational kernels and convolution operation"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "# Import the required modules\n",
    "import cv2\n",
    "from IPython.display import clear_output\n",
    "import time\n",
    "import PIL.Image\n",
    "from io import BytesIO\n",
    "import IPython.display\n",
    "import numpy as np\n",
    "from numba import cuda"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#Use 'jpeg' instead of 'png' (~5 times faster)\n",
    "def array_to_image(a, fmt='jpeg'):\n",
    "    #Create binary stream object\n",
    "    f = BytesIO()\n",
    "    \n",
    "    #Convert array to binary stream object\n",
    "    PIL.Image.fromarray(a).save(f, fmt)\n",
    "    \n",
    "    return IPython.display.Image(data=f.getvalue())"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "def get_frame(cam):\n",
    "    # Capture frame-by-frame\n",
    "    ret, frame = cam.read()\n",
    "    \n",
    "    #flip image for natural viewing\n",
    "    frame = cv2.flip(frame, 1)\n",
    "    \n",
    "    return frame"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit\n",
    "def blur_kernel(x, out):\n",
    "    blur_size = 3\n",
    "    i,j = cuda.grid(2)\n",
    "    \n",
    "    if i < x.shape[0] and j < x.shape[1]:\n",
    "        for c in range(0, 3):\n",
    "            out[i,j,c] = 0\n",
    "        count = 0\n",
    "        for ii in range(-blur_size,blur_size+1):\n",
    "            for jj in range(-blur_size,blur_size+1):\n",
    "                if (i+ii)>-blur_size and (i+ii)<x.shape[0]-blur_size+1 and (j+jj)>-blur_size and (j+jj)<x.shape[1]-blur_size+1:\n",
    "                    for c in range(0, 3):\n",
    "                        out[i,j,c] += x[i+ii,j+jj,c]\n",
    "                    count = count + 1\n",
    "        for c in range(0, 3):\n",
    "            out[i,j,c] = out[i,j,c]/count\n",
    "    return"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit\n",
    "def bw_kernel(x, out):\n",
    "    i,j = cuda.grid(2)\n",
    "    if i < x.shape[0] and j < x.shape[1]:\n",
    "        out[i,j] = 0.2126*x[i,j,0] + 0.7152*x[i,j,1] + 0.0722*x[i,j,2]\n",
    "    return"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit\n",
    "def conv_kernel(x, kernel, out):\n",
    "    kernel_size = (kernel.shape[0]-1)//2\n",
    "    i,j = cuda.grid(2)\n",
    "    \n",
    "    if i < x.shape[0] and j < x.shape[1]:\n",
    "        out[i,j] = 0\n",
    "\n",
    "        for ii in range(-kernel_size,kernel_size+1):\n",
    "            for jj in range(-kernel_size,kernel_size+1):\n",
    "                if (i+ii)>-kernel_size and (i+ii)<x.shape[0]-kernel_size+1 and (j+jj)>-kernel_size and (j+jj)<x.shape[1]-kernel_size+1:\n",
    "                    out[i,j] += x[i+ii,j+jj] * kernel[ii+kernel_size,jj+kernel_size]\n",
    "    return"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "@cuda.jit\n",
    "def max_pool(x, out):\n",
    "    i,j = cuda.grid(2)\n",
    "    \n",
    "    if i < out.shape[0] and j < out.shape[1]:\n",
    "            \n",
    "        max_brightness = 0\n",
    "\n",
    "        for ii in range(0,6):\n",
    "            for jj in range(0,6):\n",
    "                if (x[5*i+ii,5*j+jj] > max_brightness):\n",
    "                    max_brightness = x[5*i+ii,5*j+jj]\n",
    "        \n",
    "        out[i,j] = max_brightness\n",
    "    return  "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#Try different convolution kernels\n",
    "gauss_blur = np.array([[1,2,1],[2,4,2],[1,2,1]])/16.0\n",
    "edge_detector_1 = np.array([[1,0,-1],[0,0,0],[-1,0,1]])\n",
    "edge_detector_2 = np.array([[0,1,0],[1,-4,1],[0,1,0]])\n",
    "edge_detector_3 = np.array([[-1,-1,-1],[-1,8,-1],[-1,-1,-1]])\n",
    "sharpen = np.array([[0,-1,0],[-1,5,-1],[0,-1,0]])"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "scrolled": true
   },
   "outputs": [],
   "source": [
    "d_conv_kernel_1 = cuda.to_device(edge_detector_1)\n",
    "d_conv_kernel_2 = cuda.to_device(edge_detector_2)\n",
    "d_conv_kernel_3 = cuda.to_device(edge_detector_3)\n",
    "cam = cv2.VideoCapture(0)\n",
    "frame = get_frame(cam)\n",
    "d_out = cuda.device_array(frame.shape[0:2])\n",
    "d_out_1 = cuda.device_array(frame.shape[0:2])\n",
    "d_out_2 = cuda.device_array(frame.shape[0:2])\n",
    "d_out_3 = cuda.device_array(frame.shape[0:2])\n",
    "d_out_pool_1 = cuda.device_array( tuple(x//5 for x in frame.shape[0:2]) )\n",
    "d_out_pool_2 = cuda.device_array( tuple(x//5 for x in frame.shape[0:2]) )\n",
    "d_out_pool_3 = cuda.device_array( tuple(x//5 for x in frame.shape[0:2]) )\n",
    "\n",
    "\n",
    "d = IPython.display.display(\"\", display_id=1)\n",
    "d2 = IPython.display.display(\"\", display_id=2)\n",
    "while True:\n",
    "    try:\n",
    "        t1 = time.time()\n",
    "        \n",
    "        # Capture frame-by-frame\n",
    "        frame = get_frame(cam)\n",
    "        \n",
    "        # Convert the image from OpenCV BGR format to matplotlib RGB format\n",
    "        # to display the image\n",
    "        frame = cv2.cvtColor(frame, cv2.COLOR_BGR2RGB)\n",
    "        frame = frame.astype(np.uint32)\n",
    "        \n",
    "        # Send to GPU\n",
    "        d_frame = cuda.to_device(frame)\n",
    "        \n",
    "        #Run kernel\n",
    "        griddim = 30, 40\n",
    "        blockdim = 16, 16\n",
    "        \n",
    "        bw_kernel[griddim, blockdim](d_frame, d_out)\n",
    "        conv_kernel[griddim, blockdim](d_out, d_conv_kernel_1, d_out_1)\n",
    "        conv_kernel[griddim, blockdim](d_out, d_conv_kernel_2, d_out_2)\n",
    "        conv_kernel[griddim, blockdim](d_out, d_conv_kernel_3, d_out_3)\n",
    "        \n",
    "        #Apply max_pooling layer\n",
    "        max_pool[griddim, blockdim](d_out_1, d_out_pool_1)\n",
    "        max_pool[griddim, blockdim](d_out_2, d_out_pool_2)\n",
    "        max_pool[griddim, blockdim](d_out_3, d_out_pool_3)\n",
    "        \n",
    "        #Send back to CPU\n",
    "        frame_GPU_1 = d_out_pool_1.copy_to_host().astype(np.uint8)\n",
    "        frame_GPU_2 = d_out_pool_2.copy_to_host().astype(np.uint8)\n",
    "        frame_GPU_3 = d_out_pool_3.copy_to_host().astype(np.uint8)\n",
    "        \n",
    "        #Concatenate three resulting layers\n",
    "        frame_GPU = np.concatenate((frame_GPU_1, frame_GPU_2, frame_GPU_3), axis=1)\n",
    "    \n",
    "        im = array_to_image(frame_GPU)\n",
    "        \n",
    "        d.update(im)\n",
    "        \n",
    "        t2 = time.time()\n",
    "        \n",
    "        s = f\"\"\"{int(1/(t2-t1))} FPS\"\"\"\n",
    "        d2.update( IPython.display.HTML(s) )\n",
    "    except KeyboardInterrupt:\n",
    "        print()\n",
    "        cam.release()\n",
    "        IPython.display.clear_output()\n",
    "        print (\"Stream stopped\")\n",
    "        break"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": []
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.7.1"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}