SparseTest0.ipynb

{
 "cells": [
  {
   "cell_type": "code",
   "execution_count": 1,
   "id": "4309590e",
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "DEVICE:GPU\n"
     ]
    }
   ],
   "source": [
    "from tinygrad.densetensor import DenseTensor\n",
    "from tinygrad.sparsetensor import SparseTensor\n",
    "import numpy as np\n",
    "\n",
    "%load_ext autoreload\n",
    "%autoreload 2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "id": "4cb532fa",
   "metadata": {},
   "outputs": [],
   "source": [
    "x_init = np.random.randn(1,3).astype(np.float32)\n",
    "x2_init = np.random.randn(3).astype(np.float32)\n",
    "U_init = np.random.randn(3,3).astype(np.float32)\n",
    "V_init = np.random.randn(3,3).astype(np.float32)\n",
    "W_init = np.random.randn(3,3).astype(np.float32)\n",
    "m_init = np.random.randn(1,3).astype(np.float32)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "id": "3fb8a745",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(array([1.0457091], dtype=float32),\n",
       " <DenseTensor <GPUBuffer with shape (1, 3)> with grad <GPUBuffer with shape (1, 3)>>)"
      ]
     },
     "execution_count": 3,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "x = DenseTensor(x_init)\n",
    "W = DenseTensor(W_init)\n",
    "m = DenseTensor(m_init)\n",
    "out = x.dot(W).relu()\n",
    "out = out.logsoftmax()\n",
    "out = out.mul(m).add(m).sum()\n",
    "out.backward()\n",
    "\n",
    "out.cpu().data, x"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "id": "46b0ba5a",
   "metadata": {
    "scrolled": true
   },
   "outputs": [
    {
     "ename": "IndexError",
     "evalue": "tuple index out of range",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mIndexError\u001b[0m                                Traceback (most recent call last)",
      "\u001b[0;32m/tmp/ipykernel_12222/814277375.py\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[1;32m      1\u001b[0m \u001b[0mx2\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mDenseTensor\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mx2_init\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;31m#.gpu()\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m      2\u001b[0m \u001b[0mW\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mSparseTensor\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mW_init\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m----> 3\u001b[0;31m \u001b[0mout\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mW\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mdot\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mx2\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mrelu\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0msum\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m      4\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m      5\u001b[0m \u001b[0mout\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mbackward\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/mnt/HD2/ML/tinygrad/tinygrad/sparsetensor.py\u001b[0m in \u001b[0;36mdot\u001b[0;34m(self, w)\u001b[0m\n\u001b[1;32m    460\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    461\u001b[0m   \u001b[0;32mdef\u001b[0m \u001b[0mdot\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mself\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mw\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 462\u001b[0;31m     \u001b[0;32mreturn\u001b[0m \u001b[0mself\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mmatmul\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mw\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    463\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    464\u001b[0m   \u001b[0;32mdef\u001b[0m \u001b[0mmean\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mself\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0maxis\u001b[0m\u001b[0;34m=\u001b[0m\u001b[0;32mNone\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/mnt/HD2/ML/tinygrad/tinygrad/sparsetensor.py\u001b[0m in \u001b[0;36mdispatch\u001b[0;34m(*x, **kwargs)\u001b[0m\n\u001b[1;32m    577\u001b[0m     \u001b[0mf\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mcl_ctx\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mcl_queue\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mane\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mdevice\u001b[0m \u001b[0;34m=\u001b[0m \u001b[0mcl_ctx\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mcl_queue\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mane\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mtt\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mdevice\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    578\u001b[0m     \u001b[0;31m# print('APPLY:', *x, f)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 579\u001b[0;31m     \u001b[0;32mreturn\u001b[0m \u001b[0mf\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mapply\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mf\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m*\u001b[0m\u001b[0mx\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m**\u001b[0m\u001b[0mkwargs\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m\u001b[1;32m    580\u001b[0m   \u001b[0msetattr\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mSparseTensor\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mname\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mdispatch\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    581\u001b[0m   \u001b[0;32mif\u001b[0m \u001b[0mname\u001b[0m \u001b[0;32min\u001b[0m \u001b[0;34m[\u001b[0m\u001b[0;34m'add'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'sub'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'mul'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'pow'\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0;34m'matmul'\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/mnt/HD2/ML/tinygrad/tinygrad/sparsetensor.py\u001b[0m in \u001b[0;36mapply\u001b[0;34m(self, *x, **kwargs)\u001b[0m\n\u001b[1;32m    561\u001b[0m       \u001b[0msetattr\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mctx\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mk\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mv\u001b[0m\u001b[0;34m)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    562\u001b[0m     \u001b[0;32mwith\u001b[0m \u001b[0mProfileOp\u001b[0m\u001b[0;34m(\u001b[0m\u001b[0mctx\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0m__class__\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0m__name__\u001b[0m\u001b[0;34m,\u001b[0m \u001b[0mx\u001b[0m\u001b[0;34m)\u001b[0m \u001b[0;32mas\u001b[0m \u001b[0mpo\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 563\u001b[0;31m       po.output = ret = DenseTensor(self.forward(ctx, *[t for t in x], **kwargs),\n\u001b[0m\u001b[1;32m    564\u001b[0m                    device=ctx.device, requires_grad=any([t.requires_grad for t in x]))\n\u001b[1;32m    565\u001b[0m     \u001b[0;32mif\u001b[0m \u001b[0mret\u001b[0m\u001b[0;34m.\u001b[0m\u001b[0mrequires_grad\u001b[0m\u001b[0;34m:\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;32m/mnt/HD2/ML/tinygrad/tinygrad/ops_gpusparse.py\u001b[0m in \u001b[0;36mforward\u001b[0;34m(ctx, input, weight)\u001b[0m\n\u001b[1;32m    355\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n\u001b[1;32m    356\u001b[0m     \u001b[0;31m# (isize,msize) x (msize,osize) = (isize,osize)\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0;32m--> 357\u001b[0;31m     matmul(ctx.cl_queue, [weight.shape[1], weight.shape[0]], None,\n\u001b[0m\u001b[1;32m    358\u001b[0m       input.data.cl, input.idxs.cl, input.nnzs.cl, np.uint32(input.ellw), weight.data.cl, ret.cl)\n\u001b[1;32m    359\u001b[0m \u001b[0;34m\u001b[0m\u001b[0m\n",
      "\u001b[0;31mIndexError\u001b[0m: tuple index out of range"
     ]
    }
   ],
   "source": [
    "x2 = DenseTensor(x2_init)#.gpu()\n",
    "W = SparseTensor(W_init)\n",
    "out = W.dot(x2).relu().sum()\n",
    "\n",
    "out.backward()\n",
    "\n",
    "out.cpu().data, x"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "id": "c3d0659d",
   "metadata": {},
   "outputs": [],
   "source": [
    "import numpy as np\n",
    "import pyopencl as cl\n",
    "\n",
    "mf = cl.mem_flags"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "id": "a0cc978b",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "((8, 8), (4, 8))"
      ]
     },
     "execution_count": 6,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "dim1 = 8\n",
    "dim2 = 8\n",
    "dim3 = 8\n",
    "bs = 4\n",
    "\n",
    "ctx = cl.create_some_context()\n",
    "queue = cl.CommandQueue(ctx,\n",
    "        properties=cl.command_queue_properties.PROFILING_ENABLE)\n",
    "\n",
    "sparsity = 0.2\n",
    "\n",
    "a = np.zeros((dim1,dim2))\n",
    "b = np.random.rand(4,dim3).astype(np.float32)\n",
    "\n",
    "a.shape, b.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 7,
   "id": "7ecf7a2b",
   "metadata": {},
   "outputs": [],
   "source": [
    "def fill_sparse(mat, sparsity=0.1):\n",
    "    indices = np.array(range(mat.shape[1]))\n",
    "    nrows = int(mat.shape[1]*sparsity)\n",
    "    for row in range(mat.shape[0]):\n",
    "        lim = nrows #+ int(np.random.random()*3)\n",
    "        mat[row][np.random.permutation(indices)[:lim]] = np.random.random(lim)\n",
    "    return mat\n",
    "\n",
    "a = fill_sparse(a, 0.9)\n",
    "#b = fill_sparse(b, sparsity)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 8,
   "id": "b32f0e9a",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[0.59460324, 0.86305024, 0.63867457, 0.        , 0.51579547,\n",
       "        0.78790121, 0.93278368, 0.02474353],\n",
       "       [0.56595709, 0.53133109, 0.        , 0.54216719, 0.6889426 ,\n",
       "        0.32183103, 0.64553166, 0.53694062],\n",
       "       [0.80983952, 0.        , 0.73776245, 0.76804488, 0.79568409,\n",
       "        0.9053589 , 0.55798573, 0.11461218],\n",
       "       [0.09699463, 0.6963742 , 0.30319763, 0.51583918, 0.83718424,\n",
       "        0.        , 0.05103909, 0.27944975],\n",
       "       [0.75281157, 0.48989915, 0.        , 0.82777856, 0.68470816,\n",
       "        0.95266816, 0.10713863, 0.63318255],\n",
       "       [0.17814208, 0.84142313, 0.        , 0.9636037 , 0.76600269,\n",
       "        0.79484073, 0.64250607, 0.16992201],\n",
       "       [0.97628359, 0.        , 0.74801394, 0.25351422, 0.59218695,\n",
       "        0.73024714, 0.80543023, 0.57306463],\n",
       "       [0.89489469, 0.31771464, 0.51842438, 0.06335784, 0.80116333,\n",
       "        0.60369206, 0.95093155, 0.        ]])"
      ]
     },
     "execution_count": 8,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "a"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 9,
   "id": "d286d740",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[0.3313637 , 0.5479767 , 0.32370082, 0.70595866, 0.85048527,\n",
       "        0.40968382, 0.17567654, 0.5410416 ],\n",
       "       [0.28101033, 0.59984696, 0.04329206, 0.92281836, 0.73979604,\n",
       "        0.48108804, 0.04234902, 0.6342557 ],\n",
       "       [0.5997033 , 0.73747677, 0.65285623, 0.28667638, 0.8824747 ,\n",
       "        0.0911755 , 0.9040162 , 0.5719434 ],\n",
       "       [0.41850978, 0.21085866, 0.11353323, 0.0624781 , 0.7506935 ,\n",
       "        0.16010118, 0.3271686 , 0.6290892 ]], dtype=float32)"
      ]
     },
     "execution_count": 9,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "b"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 10,
   "id": "45c14776",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([-0.25337243, -0.7956384 ,  0.8874337 ], dtype=float32)"
      ]
     },
     "execution_count": 10,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "x2_init.T"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 11,
   "id": "6b2b88a2",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(8, 4)"
      ]
     },
     "execution_count": 11,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "mult = a.dot(b.T)\n",
    "mult.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 12,
   "id": "d33049e0",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(8, 4)"
      ]
     },
     "execution_count": 12,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "mult.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 13,
   "id": "67872651",
   "metadata": {},
   "outputs": [],
   "source": [
    "def to_data(mat):\n",
    "    all_rows = []\n",
    "    all_idxs = []\n",
    "    all_nnzs = []\n",
    "    for row in range(mat.shape[0]):\n",
    "        rowdata = []\n",
    "        colidxs = []\n",
    "        all_nnzs.append(0)\n",
    "        for col in range(mat.shape[1]):\n",
    "            val = mat[row][col]\n",
    "            if val != 0:\n",
    "                rowdata.append(val)\n",
    "                colidxs.append(col)\n",
    "                all_nnzs[-1] += 1\n",
    "        all_rows.append(rowdata)\n",
    "        all_idxs.append(colidxs)\n",
    "    \n",
    "    ellwidth = int(np.sqrt(np.max(all_nnzs))+1)**2\n",
    "    #all_rows = np.array(all_rows)#.astype(np.float32).flatten()           \n",
    "    for row in range(mat.shape[0]):\n",
    "        #print(row, all_rows)\n",
    "        all_rows[row] = np.array(all_rows[row])\n",
    "        all_rows[row].resize(ellwidth)\n",
    "        all_idxs[row] = np.array(all_idxs[row])\n",
    "        all_idxs[row].resize(ellwidth)\n",
    "        print(all_idxs[row])\n",
    "    all_rows = np.array(all_rows).astype(np.float32).flatten()\n",
    "    print(all_idxs)\n",
    "    all_idxs = np.array(all_idxs).astype(np.uint32).flatten()\n",
    "    all_nnzs = np.array(all_nnzs).astype(np.uint32)\n",
    "    \n",
    "    return all_rows, all_idxs, all_nnzs, ellwidth"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 14,
   "id": "5c9c6134",
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "[0 1 2 4 5 6 7 0 0]\n",
      "[0 1 3 4 5 6 7 0 0]\n",
      "[0 2 3 4 5 6 7 0 0]\n",
      "[0 1 2 3 4 6 7 0 0]\n",
      "[0 1 3 4 5 6 7 0 0]\n",
      "[0 1 3 4 5 6 7 0 0]\n",
      "[0 2 3 4 5 6 7 0 0]\n",
      "[0 1 2 3 4 5 6 0 0]\n",
      "[array([0, 1, 2, 4, 5, 6, 7, 0, 0]), array([0, 1, 3, 4, 5, 6, 7, 0, 0]), array([0, 2, 3, 4, 5, 6, 7, 0, 0]), array([0, 1, 2, 3, 4, 6, 7, 0, 0]), array([0, 1, 3, 4, 5, 6, 7, 0, 0]), array([0, 1, 3, 4, 5, 6, 7, 0, 0]), array([0, 2, 3, 4, 5, 6, 7, 0, 0]), array([0, 1, 2, 3, 4, 5, 6, 0, 0])]\n"
     ]
    },
    {
     "data": {
      "text/plain": [
       "(array([0.59460324, 0.8630502 , 0.63867456, 0.51579547, 0.7879012 ,\n",
       "        0.93278366, 0.02474353, 0.        , 0.        , 0.56595707,\n",
       "        0.53133106, 0.5421672 , 0.6889426 , 0.32183105, 0.64553165,\n",
       "        0.53694063, 0.        , 0.        , 0.80983955, 0.73776245,\n",
       "        0.7680449 , 0.7956841 , 0.9053589 , 0.5579857 , 0.11461218,\n",
       "        0.        , 0.        , 0.09699463, 0.6963742 , 0.30319762,\n",
       "        0.51583916, 0.83718425, 0.05103909, 0.27944976, 0.        ,\n",
       "        0.        , 0.75281155, 0.48989916, 0.8277786 , 0.6847082 ,\n",
       "        0.9526682 , 0.10713863, 0.6331825 , 0.        , 0.        ,\n",
       "        0.17814207, 0.84142315, 0.9636037 , 0.7660027 , 0.79484075,\n",
       "        0.64250606, 0.16992201, 0.        , 0.        , 0.9762836 ,\n",
       "        0.748014  , 0.25351423, 0.5921869 , 0.73024714, 0.80543023,\n",
       "        0.5730646 , 0.        , 0.        , 0.8948947 , 0.31771463,\n",
       "        0.5184244 , 0.06335784, 0.8011633 , 0.60369205, 0.95093155,\n",
       "        0.        , 0.        ], dtype=float32),\n",
       " array([0, 1, 2, 4, 5, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4,\n",
       "        5, 6, 7, 0, 0, 0, 1, 2, 3, 4, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0,\n",
       "        0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4, 5, 6, 7, 0, 0, 0, 1, 2,\n",
       "        3, 4, 5, 6, 0, 0], dtype=uint32),\n",
       " array([7, 7, 7, 7, 7, 7, 7, 7], dtype=uint32),\n",
       " 9)"
      ]
     },
     "execution_count": 14,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "adata, acols, annz, ellwa = to_data(a)\n",
    "adata, acols, annz, ellwa"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 15,
   "id": "d8e29e3c",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(8, 8)"
      ]
     },
     "execution_count": 15,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "a.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 16,
   "id": "233fb7fe",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "((72,), (72,), (8,), 9)"
      ]
     },
     "execution_count": 16,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "adata.shape, acols.shape, annz.shape, ellwa"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 17,
   "id": "fbb29574",
   "metadata": {},
   "outputs": [],
   "source": [
    "#acols = acols.astype(np.uint32)\n",
    "#annz = annz.astype(np.uint32)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 18,
   "id": "008c1a26",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(array([0.59460324, 0.8630502 , 0.63867456, 0.51579547, 0.7879012 ,\n",
       "        0.93278366, 0.02474353, 0.        , 0.        , 0.56595707,\n",
       "        0.53133106, 0.5421672 , 0.6889426 , 0.32183105, 0.64553165,\n",
       "        0.53694063, 0.        , 0.        , 0.80983955, 0.73776245,\n",
       "        0.7680449 , 0.7956841 , 0.9053589 , 0.5579857 , 0.11461218,\n",
       "        0.        , 0.        , 0.09699463, 0.6963742 , 0.30319762,\n",
       "        0.51583916, 0.83718425, 0.05103909, 0.27944976, 0.        ,\n",
       "        0.        , 0.75281155, 0.48989916, 0.8277786 , 0.6847082 ,\n",
       "        0.9526682 , 0.10713863, 0.6331825 , 0.        , 0.        ,\n",
       "        0.17814207, 0.84142315, 0.9636037 , 0.7660027 , 0.79484075,\n",
       "        0.64250606, 0.16992201, 0.        , 0.        , 0.9762836 ,\n",
       "        0.748014  , 0.25351423, 0.5921869 , 0.73024714, 0.80543023,\n",
       "        0.5730646 , 0.        , 0.        , 0.8948947 , 0.31771463,\n",
       "        0.5184244 , 0.06335784, 0.8011633 , 0.60369205, 0.95093155,\n",
       "        0.        , 0.        ], dtype=float32),\n",
       " array([0, 1, 2, 4, 5, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4,\n",
       "        5, 6, 7, 0, 0, 0, 1, 2, 3, 4, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0,\n",
       "        0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4, 5, 6, 7, 0, 0, 0, 1, 2,\n",
       "        3, 4, 5, 6, 0, 0], dtype=uint32),\n",
       " array([7, 7, 7, 7, 7, 7, 7, 7], dtype=uint32),\n",
       " array([[0.3313637 , 0.5479767 , 0.32370082, 0.70595866, 0.85048527,\n",
       "         0.40968382, 0.17567654, 0.5410416 ],\n",
       "        [0.28101033, 0.59984696, 0.04329206, 0.92281836, 0.73979604,\n",
       "         0.48108804, 0.04234902, 0.6342557 ],\n",
       "        [0.5997033 , 0.73747677, 0.65285623, 0.28667638, 0.8824747 ,\n",
       "         0.0911755 , 0.9040162 , 0.5719434 ],\n",
       "        [0.41850978, 0.21085866, 0.11353323, 0.0624781 , 0.7506935 ,\n",
       "         0.16010118, 0.3271686 , 0.6290892 ]], dtype=float32))"
      ]
     },
     "execution_count": 18,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "adata, acols, annz, b"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 19,
   "id": "6e741bf7",
   "metadata": {},
   "outputs": [],
   "source": [
    "adata_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=adata)\n",
    "acols_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=acols)\n",
    "annzs_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=annz)\n",
    "b_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=b.T)\n",
    "\n",
    "prg = cl.Program(ctx, \"\"\"\n",
    "// Every global_id_0 works on a row\n",
    "    __kernel void matmul(__global  float* matData,     // INPUT MATRIX DATA\n",
    "                            __global  uint*  colIdx,\n",
    "                            __global  uint*  rowNnz,\n",
    "                            uint   ellwidth,\n",
    "                            __global  float* vector_x,    // INPUT\n",
    "                            __global  float* vector_y    // OUTPUT\n",
    "                            ) { // LOCAL SHARED BUFFER\n",
    "      uint gid = get_global_id(0);\n",
    "      uint nrows = get_global_size(0);\n",
    "      uint gid2 = get_global_id(1);\n",
    "\n",
    "      uint nnz    = rowNnz[gid];\n",
    "      uint baseidx = gid2*nrows;\n",
    "      float sum = 0;\n",
    "      for (uint i = 0; i < nnz; i++) {\n",
    "        uint index   = (gid * ellwidth) + i;\n",
    "        uint col     = colIdx[index];\n",
    "        float aval  = matData[index];\n",
    "        float xval  = vector_x[baseidx+col];\n",
    "        //printf(\"aval, xval: %.2f,%.2f: (%i,%i) \\\\n\", aval, xval, col, index);\n",
    "        sum  += aval * xval;\n",
    "      }\n",
    "      printf(\"SUM/NNZ: %.2f %i \\\\n\", sum, nnz);\n",
    "      vector_y[baseidx+gid] = sum;\n",
    "    }\"\"\").build()"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 20,
   "id": "b4c85bcb",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "((8, 8), (4, 8))"
      ]
     },
     "execution_count": 20,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "a.shape, b.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 21,
   "id": "c740a4fe",
   "metadata": {},
   "outputs": [],
   "source": [
    "res = np.zeros(a.shape[0]).astype(np.float32)\n",
    "#res"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 22,
   "id": "07aa4d30",
   "metadata": {},
   "outputs": [],
   "source": [
    "rows = a.shape[0]"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 23,
   "id": "98da857f",
   "metadata": {},
   "outputs": [],
   "source": [
    "mult = mult.astype(np.float32)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 24,
   "id": "cc7365fc",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "<pyopencl._cl.NannyEvent at 0x7fab40115310>"
      ]
     },
     "execution_count": 24,
     "metadata": {},
     "output_type": "execute_result"
    },
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "SUM/NNZ: 2.79 7 \n",
      "SUM/NNZ: 2.41 7 \n",
      "SUM/NNZ: 2.54 7 \n",
      "SUM/NNZ: 1.86 7 \n",
      "SUM/NNZ: 2.20 7 \n",
      "SUM/NNZ: 2.43 7 \n",
      "SUM/NNZ: 2.79 7 \n",
      "SUM/NNZ: 2.75 7 \n",
      "SUM/NNZ: 1.34 7 \n",
      "SUM/NNZ: 1.50 7 \n",
      "SUM/NNZ: 1.47 7 \n",
      "SUM/NNZ: 1.08 7 \n",
      "SUM/NNZ: 1.57 7 \n",
      "SUM/NNZ: 1.33 7 \n",
      "SUM/NNZ: 1.69 7 \n",
      "SUM/NNZ: 1.51 7 \n",
      "SUM/NNZ: 1.82 7 \n",
      "SUM/NNZ: 1.98 7 \n",
      "SUM/NNZ: 2.26 7 \n",
      "SUM/NNZ: 1.75 7 \n",
      "SUM/NNZ: 2.44 7 \n",
      "SUM/NNZ: 2.38 7 \n",
      "SUM/NNZ: 2.00 7 \n",
      "SUM/NNZ: 1.78 7 \n",
      "SUM/NNZ: 1.53 7 \n",
      "SUM/NNZ: 2.01 7 \n",
      "SUM/NNZ: 2.09 7 \n",
      "SUM/NNZ: 1.73 7 \n",
      "SUM/NNZ: 2.64 7 \n",
      "SUM/NNZ: 2.53 7 \n",
      "SUM/NNZ: 1.73 7 \n",
      "SUM/NNZ: 1.45 7 \n"
     ]
    }
   ],
   "source": [
    "res_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes)\n",
    "knl = prg.matmul  # Use this Kernel object for repeated calls\n",
    "knl(queue, [rows,bs], None, adata_buf, acols_buf, annzs_buf, np.uint32(ellwa), b_buf, res_buf)\n",
    "\n",
    "res_np = np.empty_like(b.T)\n",
    "cl.enqueue_copy(queue, res_np, res_buf)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 25,
   "id": "33d8a584",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "<pyopencl._cl.Buffer at 0x7fab40110d10>"
      ]
     },
     "execution_count": 25,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "res_buf"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 26,
   "id": "99faec15",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[1.8154231, 1.5282668, 2.7944448, 1.3374313],\n",
       "       [1.9831388, 2.0104804, 2.4146652, 1.5004594],\n",
       "       [2.2570374, 2.0888033, 2.5421941, 1.467593 ],\n",
       "       [1.7482169, 1.7328761, 1.8623147, 1.0750473],\n",
       "       [2.436312 , 2.6403   , 2.2001545, 1.5699865],\n",
       "       [2.3822901, 2.5280738, 2.4300745, 1.3315717],\n",
       "       [1.998972 , 1.7276658, 2.7915556, 1.6948323],\n",
       "       [1.7789352, 1.4463638, 2.749304 , 1.5135268]], dtype=float32)"
      ]
     },
     "execution_count": 26,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "res_np"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 27,
   "id": "4da1c066",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[1.8154231, 1.5282668, 2.7944448, 1.3374313],\n",
       "       [1.9831388, 2.0104804, 2.4146655, 1.5004594],\n",
       "       [2.2570372, 2.0888033, 2.5421941, 1.4675928],\n",
       "       [1.748217 , 1.732876 , 1.8623147, 1.0750474],\n",
       "       [2.436312 , 2.6403   , 2.2001548, 1.5699865],\n",
       "       [2.38229  , 2.5280738, 2.4300745, 1.3315718],\n",
       "       [1.9989719, 1.7276659, 2.791556 , 1.6948323],\n",
       "       [1.7789351, 1.4463638, 2.749304 , 1.5135268]], dtype=float32)"
      ]
     },
     "execution_count": 27,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "mult"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 28,
   "id": "c197fea2",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(8, 4)"
      ]
     },
     "execution_count": 28,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "res_np.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 29,
   "id": "b20b9da3",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(8, 4)"
      ]
     },
     "execution_count": 29,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "mult.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 30,
   "id": "2a6322f5",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "-2.3841858e-07"
      ]
     },
     "execution_count": 30,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "(res_np-mult).sum()"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "b3ad7225",
   "metadata": {},
   "source": [
    "## Weight update kernel"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 31,
   "id": "a5138cee",
   "metadata": {},
   "outputs": [],
   "source": [
    "bs = 4"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 32,
   "id": "76e44e2d",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "((4, 8), (4, 8), 2)"
      ]
     },
     "execution_count": 32,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "dim = 8\n",
    "topk = 2\n",
    "\n",
    "x = np.random.rand(bs,dim).astype(np.float32)\n",
    "y = np.random.rand(bs,dim).astype(np.float32)\n",
    "x.shape,y.shape, topk"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 33,
   "id": "75611f3a",
   "metadata": {},
   "outputs": [],
   "source": [
    "x_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=x)\n",
    "y_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=y)\n",
    "x_cp_buf = cl.Buffer(ctx, mf.WRITE_ONLY, bs*topk*topk*4)\n",
    "x_idx_buf = cl.Buffer(ctx, mf.WRITE_ONLY, bs*topk*4)\n",
    "y_idx_buf = cl.Buffer(ctx, mf.WRITE_ONLY, bs*topk*4)\n",
    "\n",
    "prg = cl.Program(ctx, \"\"\"\n",
    "    // sorts x and y in ascending order and returns sorted indices\n",
    "    __kernel void genwupdate2(__global  float* x,     // INPUT MATRIX DATA\n",
    "                             __global  float* y,    // INPUT\n",
    "                             __global  float* xout,    // INPUT\n",
    "                             uint topk,\n",
    "                             __global  uint* xoutidx,    // INPUT\n",
    "                             __global  uint* youtidx    // INPUT\n",
    "                            ) { // LOCAL SHARED BUFFER\n",
    "      uint gid = get_global_id(0);\n",
    "      uint n = get_global_size(0);\n",
    "      uint bs = get_global_size(1);\n",
    "      uint gid2 = get_global_id(1);\n",
    "\n",
    "      uint idx = n*gid2+gid;\n",
    "\n",
    "      float valx = x[idx];\n",
    "      float valy = y[idx];\n",
    "      uint posx = 0;\n",
    "      uint posy = 0;\n",
    "      for (uint i = 0; i < n; i++) {\n",
    "        uint idx2 = n*gid2+i;\n",
    "        float tempval = x[idx2];\n",
    "        float tempval2 = y[idx2];\n",
    "        bool larger = tempval > valx;\n",
    "        bool larger2 = tempval2 > valy;\n",
    "\n",
    "        barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "        posx += (larger)?1:0;\n",
    "        posy += (larger2)?1:0;\n",
    "        barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "      }\n",
    "      barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "      //printf(\"posx:%i\", posx);\n",
    "      if (posx < topk) {\n",
    "        xoutidx[posx+topk*gid2] = gid;\n",
    "      }\n",
    "      if (posy < topk) {\n",
    "        youtidx[posy+topk*gid2] = gid;\n",
    "      }\n",
    "      barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "      if (gid < topk) {\n",
    "        for (uint j=0; j<topk; j++) {\n",
    "          float res = x[xoutidx[gid+topk*gid2]+gid2*n] * y[youtidx[j+topk*gid2]+gid2*n];\n",
    "          //printf(\"\\\\nJ:%i  gid:%i\", j, gid);\n",
    "          //printf(\"\\\\nRES:%.2f - %i - %i -  %.2f - %.2f\",res, xoutidx[gid+topk*gid2], youtidx[j+topk*gid2], x[xoutidx[gid+topk*gid2]+gid2*n], y[youtidx[j+topk*gid2]+gid2*n]);\n",
    "          barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "          xout[gid2*topk*topk+j*topk+gid] = res;\n",
    "          barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "          \n",
    "        }\n",
    "      }\n",
    "      barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "    }\"\"\").build()"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 34,
   "id": "88864470",
   "metadata": {},
   "outputs": [],
   "source": [
    "x_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=x)\n",
    "y_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=y)\n",
    "x_cp_buf = cl.Buffer(ctx, mf.WRITE_ONLY, bs*topk*topk*4)\n",
    "x_idx_buf = cl.Buffer(ctx, mf.WRITE_ONLY, bs*topk*4)\n",
    "y_idx_buf = cl.Buffer(ctx, mf.WRITE_ONLY, bs*topk*4)\n",
    "\n",
    "prg = cl.Program(ctx, \"\"\"\n",
    "    // sorts x and y in ascending order and returns sorted indices\n",
    "    __kernel void genwupdate2(__global  float* x,     // INPUT MATRIX DATA\n",
    "                             __global  float* y,    // INPUT\n",
    "                             __global  float* xout,    // INPUT\n",
    "                             uint topk,\n",
    "                             uint bs,\n",
    "                             __global  uint* xoutidx,    // INPUT\n",
    "                             __global  uint* youtidx    // INPUT\n",
    "                            ) { // LOCAL SHARED BUFFER\n",
    "      uint gid = get_global_id(0);\n",
    "      uint n = get_global_size(0);\n",
    "      //uint bs = get_global_size(1);\n",
    "      //uint gid2 = get_global_id(1);\n",
    "\n",
    "      for (uint gid2=0; gid2<bs; gid2++){\n",
    "        uint idx = n*gid2+gid;\n",
    "\n",
    "        float valx = x[idx];\n",
    "        float valy = y[idx];\n",
    "        uint posx = 0;\n",
    "        uint posy = 0;\n",
    "        for (uint i = 0; i < n; i++) {\n",
    "          uint idx2 = n*gid2+i;\n",
    "          float tempval = x[idx2];\n",
    "          float tempval2 = y[idx2];\n",
    "          bool larger = tempval > valx;\n",
    "          bool larger2 = tempval2 > valy;\n",
    "\n",
    "          posx += (larger)?1:0;\n",
    "          posy += (larger2)?1:0;\n",
    "        }\n",
    "        //printf(\"posx:%i\", posx);\n",
    "        if (posx < topk) {\n",
    "        xoutidx[posx+topk*gid2] = gid;\n",
    "        }\n",
    "        if (posy < topk) {\n",
    "          youtidx[posy+topk*gid2] = gid;\n",
    "        }\n",
    "      }\n",
    "      for (uint gid2=0; gid2<bs; gid2++){\n",
    "        if (gid < topk) {\n",
    "          for (uint j=0; j<topk; j++) {\n",
    "            float res = x[xoutidx[gid+topk*gid2]+gid2*n] * y[youtidx[j+topk*gid2]+gid2*n];\n",
    "            //printf(\"\\\\nJ:%i  gid:%i\", j, gid);\n",
    "            //printf(\"\\\\nRES:%.2f - %i - %i -  %.2f - %.2f\",res, xoutidx[gid+topk*gid2], youtidx[j+topk*gid2], x[xoutidx[gid+topk*gid2]+gid2*n], y[youtidx[j+topk*gid2]+gid2*n]);\n",
    "            //barrier(CLK_GLOBAL_MEM_FENCE);\n",
    "            xout[gid2*topk*topk+j*topk+gid] = res;\n",
    "          }\n",
    "        }\n",
    "      }\n",
    "    }\"\"\").build()"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 35,
   "id": "94682965",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "<pyopencl._cl.NannyEvent at 0x7fab40141b80>"
      ]
     },
     "execution_count": 35,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "knl = prg.genwupdate2  # Use this Kernel object for repeated calls\n",
    "evt = knl(queue, [dim], None, x_buf, y_buf, x_cp_buf, np.uint32(topk), np.uint32(bs), x_idx_buf, y_idx_buf)\n",
    "\n",
    "#evt.wait()\n",
    "resx = np.zeros(bs*topk*topk).astype(np.float32)\n",
    "resxidx = np.zeros(bs*topk).astype(np.uint32)\n",
    "resyidx = np.zeros(bs*topk).astype(np.uint32)\n",
    "\n",
    "cl.enqueue_copy(queue, resx, x_cp_buf)\n",
    "cl.enqueue_copy(queue, resxidx, x_idx_buf)\n",
    "cl.enqueue_copy(queue, resyidx, y_idx_buf)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 36,
   "id": "88725ae2",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[0.48934165, 0.71112716, 0.29748935, 0.782211  , 0.39926562,\n",
       "        0.8514658 , 0.2073185 , 0.52108914],\n",
       "       [0.9927594 , 0.22187403, 0.3556742 , 0.21139157, 0.73280114,\n",
       "        0.6351652 , 0.12583221, 0.74821895],\n",
       "       [0.1520464 , 0.4024638 , 0.17809196, 0.4235243 , 0.36142918,\n",
       "        0.41935408, 0.10778701, 0.6391944 ],\n",
       "       [0.6635596 , 0.01897575, 0.52861106, 0.47245625, 0.67950535,\n",
       "        0.50114566, 0.81095266, 0.5377907 ]], dtype=float32)"
      ]
     },
     "execution_count": 36,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "x"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 37,
   "id": "4c150d0f",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[0.13566868, 0.12836814, 0.82512575, 0.19204263, 0.21516852,\n",
       "        0.6004559 , 0.34215394, 0.85616165],\n",
       "       [0.51213527, 0.02328744, 0.24240972, 0.83984864, 0.3979477 ,\n",
       "        0.37296107, 0.11364152, 0.265691  ],\n",
       "       [0.5349935 , 0.83051544, 0.21390152, 0.34057328, 0.6541128 ,\n",
       "        0.11135104, 0.32611912, 0.9343506 ],\n",
       "       [0.7771125 , 0.07796126, 0.36196333, 0.9460012 , 0.9146572 ,\n",
       "        0.9022052 , 0.85295373, 0.03266399]], dtype=float32)"
      ]
     },
     "execution_count": 37,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "y"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 38,
   "id": "990b8af7",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "((4, 8), (4, 8))"
      ]
     },
     "execution_count": 38,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "x.shape, y.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 39,
   "id": "1c18b7bb",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([0.7289924 , 0.6696991 , 0.7025664 , 0.64542246, 0.83376765,\n",
       "       0.62839067, 0.5084271 , 0.38318932, 0.5972317 , 0.39572018,\n",
       "       0.53086084, 0.35174346, 0.76716214, 0.64281285, 0.7417437 ,\n",
       "       0.62151444], dtype=float32)"
      ]
     },
     "execution_count": 39,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resx"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 40,
   "id": "58d4154c",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[[0.7289924 , 0.6696991 ],\n",
       "        [0.7025664 , 0.64542246]],\n",
       "\n",
       "       [[0.83376765, 0.62839067],\n",
       "        [0.5084271 , 0.38318932]],\n",
       "\n",
       "       [[0.5972317 , 0.39572018],\n",
       "        [0.53086084, 0.35174346]],\n",
       "\n",
       "       [[0.76716214, 0.64281285],\n",
       "        [0.7417437 , 0.62151444]]], dtype=float32)"
      ]
     },
     "execution_count": 40,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resx.reshape(bs,topk,topk)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 41,
   "id": "5cca2876",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([5, 3, 0, 7, 7, 3, 6, 4], dtype=uint32)"
      ]
     },
     "execution_count": 41,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resxidx"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 94,
   "id": "80deeac2",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([7, 2, 3, 0, 7, 1, 3, 4], dtype=uint32)"
      ]
     },
     "execution_count": 94,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resyidx"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 42,
   "id": "f366fda3",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([7, 2, 3, 0, 7, 1, 3, 4], dtype=uint32)"
      ]
     },
     "execution_count": 42,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resyidx// Every global_id_0 works on a row\n",
    "    __kernel void addvals(__global  float* matData,     // INPUT MATRIX DATA\n",
    "                         __global  uint*  colIdx,\n",
    "                         __global  uint*  rowNnz,\n",
    "                         float lr,\n",
    "                         uint   ellwidth,\n",
    "                         __global  float* updatevals,    // INPUT\n",
    "                         __global  uint* updatexidx,\n",
    "                         __global  uint* updateyidx\n",
    "                         ) { // LOCAL SHARED BUFFER\n",
    "      uint gid = get_global_id(0);\n",
    "      uint gid2 = get_global_id(1);\n",
    "      uint topk = get_global_size(0);\n",
    "      uint bs = get_global_size(1);\n",
    "      uint baseupdateidx = topk*topk*gid2;\n",
    "      uint baseidxidx = topk*gid2;\n",
    "      uint row = updateyidx[baseidxidx+gid];\n",
    "\n",
    "      for (uint i=0; i<topk; i++) {\n",
    "        float val = updatevals[baseupdateidx+gid*topk+i];\n",
    "        uint col = updatexidx[baseidxidx+i];\n",
    "        for (uint i=0; i<rowNnz[row]; i++) {\n",
    "          uint idx = row*ellwidth+i;\n",
    "          if (colIdx[idx] >= col) {\n",
    "            //printf(\"\\\\nFOUND:%i/%i  - idx:%i\", colIdx[idx], col, idx);\n",
    "            if (colIdx[idx] == col) {\n",
    "              matData[idx] += val*lr;\n",
    "              printf(\"\\\\nUPDATE[%i,%i]: %f\", row,col, val);\n",
    "              break;\n",
    "            } else {\n",
    "              // insert new column\n",
    "              printf(\"\\\\nINSERT[%i,%i]: %.2f\", row,col, val);\n",
    "              for (uint j=rowNnz[row]+1; j>i; j--) {\n",
    "                uint idx2 = row*ellwidth+j;\n",
    "                matData[idx2] = matData[idx2-1];\n",
    "                colIdx[idx2] = colIdx[idx2-1];\n",
    "              }\n",
    "              matData[idx] = val*lr;\n",
    "              colIdx[idx] = col;\n",
    "              rowNnz[row] += 1;\n",
    "              break;\n",
    "            }\n",
    "          }\n",
    "        }\n",
    "        if (rowNnz[row] >= ellwidth) {\n",
    "          break;\n",
    "        }\n",
    "      }\n",
    "    }\"\"\""
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 43,
   "id": "3020c7d7",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "(8, 8)"
      ]
     },
     "execution_count": 43,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "idx = 0\n",
    "xy0 = x[idx].reshape(dim,1)*y[idx]\n",
    "xy0.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 44,
   "id": "fc865ccb",
   "metadata": {},
   "outputs": [
    {
     "ename": "IndexError",
     "evalue": "index 382 is out of bounds for axis 0 with size 8",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mIndexError\u001b[0m                                Traceback (most recent call last)",
      "\u001b[0;32m/tmp/ipykernel_12222/4282766181.py\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0mxy0\u001b[0m\u001b[0;34m[\u001b[0m\u001b[0;36m382\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m[\u001b[0m\u001b[0;36m430\u001b[0m\u001b[0;34m]\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;31mIndexError\u001b[0m: index 382 is out of bounds for axis 0 with size 8"
     ]
    }
   ],
   "source": [
    "xy0[382][430]"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "d3a0961d",
   "metadata": {},
   "outputs": [],
   "source": [
    "asdf"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "6644b755",
   "metadata": {},
   "source": [
    "### update vals"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 95,
   "id": "73b60325",
   "metadata": {},
   "outputs": [],
   "source": [
    "adata_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=adata)\n",
    "acols_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=acols)\n",
    "annzs_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=annz)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 96,
   "id": "7cb03af3",
   "metadata": {},
   "outputs": [],
   "source": [
    "prg = cl.Program(ctx, \"\"\"\n",
    "// Every global_id_0 works on a row\n",
    "    __kernel void addvals(__global  float* matData,     // INPUT MATRIX DATA\n",
    "                         __global  uint*  colIdx,\n",
    "                         __global  uint*  rowNnz,\n",
    "                         float lr,\n",
    "                         uint   ellwidth,\n",
    "                         __global  float* updatevals,    // INPUT\n",
    "                         __global  uint* updatexidx,\n",
    "                         __global  uint* updateyidx\n",
    "                         ) { // LOCAL SHARED BUFFER\n",
    "      uint gid = get_global_id(0);\n",
    "      uint gid2 = get_global_id(1);\n",
    "      uint topk = get_global_size(0);\n",
    "      uint bs = get_global_size(1);\n",
    "      uint baseupdateidx = topk*topk*gid2;\n",
    "      uint baseidxidx = topk*gid2;\n",
    "      uint col = updateyidx[baseidxidx+gid];\n",
    "\n",
    "      for (uint i=0; i<topk; i++) {\n",
    "        float val = updatevals[baseupdateidx+gid*topk+i];\n",
    "        uint row = updatexidx[baseidxidx+i];\n",
    "        for (uint i=0; i<rowNnz[row]; i++) {\n",
    "          uint idx = row*ellwidth+i;\n",
    "          if (colIdx[idx] >= col) {\n",
    "            //printf(\"\\\\nFOUND:%i/%i  - idx:%i\", colIdx[idx], col, idx);\n",
    "            if (colIdx[idx] == col) {\n",
    "              matData[idx] += -val*lr;\n",
    "              printf(\"\\\\nUPDATE[%i,%i]: %f\", row,col, val);\n",
    "              break;\n",
    "            } else {\n",
    "              // insert new column\n",
    "              printf(\"\\\\nINSERT[%i,%i]: %.2f\", row,col, val);\n",
    "              for (uint j=rowNnz[row]+1; j>i; j--) {\n",
    "                uint idx2 = row*ellwidth+j;\n",
    "                matData[idx2] = matData[idx2-1];\n",
    "                colIdx[idx2] = colIdx[idx2-1];\n",
    "              }\n",
    "              matData[idx] = -val*lr;\n",
    "              colIdx[idx] = col;\n",
    "              rowNnz[row] += 1;\n",
    "              break;\n",
    "            }\n",
    "          }\n",
    "        }\n",
    "        if (rowNnz[row] >= ellwidth) {\n",
    "          break;\n",
    "        }\n",
    "      }\n",
    "    }\"\"\").build()"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 97,
   "id": "0aab86e1",
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "\n",
      "UPDATE[6,3]: 0.767162\n",
      "UPDATE[7,1]: 0.530861\n",
      "UPDATE[0,0]: 0.508427\n",
      "INSERT[5,2]: 0.70\n",
      "UPDATE[6,4]: 0.741744\n",
      "INSERT[0,3]: 0.83\n",
      "UPDATE[3,1]: 0.351743\n",
      "UPDATE[3,2]: 0.645422\n",
      "UPDATE[4,3]: 0.642813\n",
      "UPDATE[3,7]: 0.395720\n",
      "UPDATE[7,0]: 0.383189\n",
      "UPDATE[3,7]: 0.669699\n",
      "UPDATE[4,4]: 0.621514\n",
      "UPDATE[7,3]: 0.628391"
     ]
    },
    {
     "data": {
      "text/plain": [
       "<pyopencl._cl.NannyEvent at 0x7faafd3df7c0>"
      ]
     },
     "execution_count": 97,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "knl = prg.addvals  # Use this Kernel object for repeated calls\n",
    "knl(queue, [topk,bs], None, adata_buf, acols_buf, annzs_buf, np.float32(1), np.uint32(ellwa), x_cp_buf, x_idx_buf, y_idx_buf)\n",
    "\n",
    "resa = np.empty_like(adata)\n",
    "resaidx = np.zeros(acols.shape).astype(np.uint32)\n",
    "resannz = np.zeros(annz.shape).astype(np.uint32)\n",
    "\n",
    "cl.enqueue_copy(queue, resa, adata_buf)\n",
    "cl.enqueue_copy(queue, resaidx, acols_buf)\n",
    "cl.enqueue_copy(queue, resannz, annzs_buf)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 98,
   "id": "922fdfb7",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "9"
      ]
     },
     "execution_count": 98,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "ellwa"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 99,
   "id": "9da55fa5",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[0.59460324, 0.8630502 , 0.63867456, 0.51579547, 0.7879012 ,\n",
       "        0.93278366, 0.02474353, 0.        , 0.        ],\n",
       "       [0.56595707, 0.53133106, 0.5421672 , 0.6889426 , 0.32183105,\n",
       "        0.64553165, 0.53694063, 0.        , 0.        ],\n",
       "       [0.80983955, 0.73776245, 0.7680449 , 0.7956841 , 0.9053589 ,\n",
       "        0.5579857 , 0.11461218, 0.        , 0.        ],\n",
       "       [0.09699463, 0.6963742 , 0.30319762, 0.51583916, 0.83718425,\n",
       "        0.05103909, 0.27944976, 0.        , 0.        ],\n",
       "       [0.75281155, 0.48989916, 0.8277786 , 0.6847082 , 0.9526682 ,\n",
       "        0.10713863, 0.6331825 , 0.        , 0.        ],\n",
       "       [0.17814207, 0.84142315, 0.9636037 , 0.7660027 , 0.79484075,\n",
       "        0.64250606, 0.16992201, 0.        , 0.        ],\n",
       "       [0.9762836 , 0.748014  , 0.25351423, 0.5921869 , 0.73024714,\n",
       "        0.80543023, 0.5730646 , 0.        , 0.        ],\n",
       "       [0.8948947 , 0.31771463, 0.5184244 , 0.06335784, 0.8011633 ,\n",
       "        0.60369205, 0.95093155, 0.        , 0.        ]], dtype=float32)"
      ]
     },
     "execution_count": 99,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "adata2 = adata.reshape(-1, ellwa)\n",
    "adata2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 100,
   "id": "466c61fb",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[ 0.08617616,  0.8630502 ,  0.63867456, -0.83376765,  0.51579547,\n",
       "         0.7879012 ,  0.93278366,  0.02474353,  0.        ],\n",
       "       [ 0.56595707,  0.53133106,  0.5421672 ,  0.6889426 ,  0.32183105,\n",
       "         0.64553165,  0.53694063,  0.        ,  0.        ],\n",
       "       [ 0.80983955,  0.73776245,  0.7680449 ,  0.7956841 ,  0.9053589 ,\n",
       "         0.5579857 ,  0.11461218,  0.        ,  0.        ],\n",
       "       [ 0.09699463,  0.34463072, -0.34222484,  0.51583916,  0.83718425,\n",
       "         0.05103909, -0.7859695 ,  0.        ,  0.        ],\n",
       "       [ 0.75281155,  0.48989916,  0.18496573,  0.06319374,  0.9526682 ,\n",
       "         0.10713863,  0.6331825 ,  0.        ,  0.        ],\n",
       "       [ 0.17814207,  0.84142315, -0.7025664 ,  0.9636037 ,  0.7660027 ,\n",
       "         0.79484075,  0.64250606,  0.16992201,  0.        ],\n",
       "       [ 0.9762836 ,  0.748014  , -0.5136479 , -0.14955676,  0.73024714,\n",
       "         0.80543023,  0.5730646 ,  0.        ,  0.        ],\n",
       "       [ 0.5117054 , -0.21314621,  0.5184244 , -0.56503284,  0.8011633 ,\n",
       "         0.60369205,  0.95093155,  0.        ,  0.        ]],\n",
       "      dtype=float32)"
      ]
     },
     "execution_count": 100,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resa = resa.reshape(-1, ellwa)\n",
    "resa"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 101,
   "id": "30f36719",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[-0.5084271 ,  0.        ,  0.        , -1.3495631 , -0.27210575,\n",
       "        -0.14488244,  0.9080401 ,  0.02474353,  0.        ],\n",
       "       [ 0.        ,  0.        ,  0.        ,  0.        ,  0.        ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        ,  0.        ,  0.        ,  0.        ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ],\n",
       "       [ 0.        , -0.35174346, -0.64542246,  0.        ,  0.        ,\n",
       "         0.        , -1.0654192 ,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        , -0.64281285, -0.62151444,  0.        ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        , -1.6661701 ,  0.19760096, -0.02883804,\n",
       "         0.15233469,  0.47258407,  0.16992201,  0.        ],\n",
       "       [ 0.        ,  0.        , -0.76716214, -0.7417437 ,  0.        ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ],\n",
       "       [-0.38318932, -0.53086084,  0.        , -0.62839067,  0.        ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ]],\n",
       "      dtype=float32)"
      ]
     },
     "execution_count": 101,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resa - adata2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 102,
   "id": "9dd6baf4",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([0, 1, 2, 4, 5, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4,\n",
       "       5, 6, 7, 0, 0, 0, 1, 2, 3, 4, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0,\n",
       "       0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4, 5, 6, 7, 0, 0, 0, 1, 2,\n",
       "       3, 4, 5, 6, 0, 0], dtype=uint32)"
      ]
     },
     "execution_count": 102,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "acols"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 103,
   "id": "5c526051",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([0, 1, 2, 3, 4, 5, 6, 7, 0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4,\n",
       "       5, 6, 7, 0, 0, 0, 1, 2, 3, 4, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0,\n",
       "       0, 0, 1, 2, 3, 4, 5, 6, 7, 0, 0, 2, 3, 4, 5, 6, 7, 0, 0, 0, 1, 2,\n",
       "       3, 4, 5, 6, 0, 0], dtype=uint32)"
      ]
     },
     "execution_count": 103,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resaidx"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 104,
   "id": "f635a14d",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([8, 7, 7, 7, 7, 8, 7, 7], dtype=uint32)"
      ]
     },
     "execution_count": 104,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resannz"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 105,
   "id": "d99713b6",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([7, 7, 7, 7, 7, 7, 7, 7], dtype=uint32)"
      ]
     },
     "execution_count": 105,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "annz"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "7c703466",
   "metadata": {},
   "source": [
    "### update vals2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 131,
   "id": "ed7e786b",
   "metadata": {},
   "outputs": [],
   "source": [
    "adata_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=adata)\n",
    "acols_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=acols)\n",
    "annzs_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=annz)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 132,
   "id": "a1e8e861",
   "metadata": {},
   "outputs": [],
   "source": [
    "prg = cl.Program(ctx, \"\"\"\n",
    "// Every global_id_0 works on a row\n",
    "    __kernel void addvals(__global  float* matData,     // INPUT MATRIX DATA\n",
    "                         __global  uint*  colIdx,\n",
    "                         __global  uint*  rowNnz,\n",
    "                         float lr,\n",
    "                         uint   ellwidth,\n",
    "                         __global  float* updatevals,    // INPUT\n",
    "                         __global  uint* updatexidx,\n",
    "                         __global  uint* updateyidx\n",
    "                         ) { // LOCAL SHARED BUFFER\n",
    "      uint gid = get_global_id(0);\n",
    "      uint gid2 = get_global_id(1);\n",
    "      uint topk = get_global_size(0);\n",
    "      uint bs = get_global_size(1);\n",
    "      uint baseupdateidx = topk*topk*gid2;\n",
    "      uint baseidxidx = topk*gid2;\n",
    "      uint row = updateyidx[baseidxidx+gid];\n",
    "\n",
    "      for (uint i=0; i<topk; i++) {\n",
    "        float val = updatevals[baseupdateidx+gid*topk+i];\n",
    "        uint col = updatexidx[baseidxidx+i];\n",
    "        for (uint i=0; i<rowNnz[row]; i++) {\n",
    "          uint idx = row*ellwidth+i;\n",
    "          if (colIdx[idx] >= col) {\n",
    "            //printf(\"\\\\nFOUND:%i/%i  - idx:%i\", colIdx[idx], col, idx);\n",
    "            if (colIdx[idx] == col) {\n",
    "              matData[idx] += -val*lr;\n",
    "              printf(\"\\\\nUPDATE[%i,%i]: %f\", row,col, val);\n",
    "              break;\n",
    "            } else {\n",
    "              // insert new column\n",
    "              printf(\"\\\\nINSERT[%i,%i]: %.2f\", row,col, val);\n",
    "              for (uint j=rowNnz[row]+1; j>i; j--) {\n",
    "                uint idx2 = row*ellwidth+j;\n",
    "                matData[idx2] = matData[idx2-1];\n",
    "                colIdx[idx2] = colIdx[idx2-1];\n",
    "              }\n",
    "              matData[idx] = -val*lr;\n",
    "              colIdx[idx] = col;\n",
    "              rowNnz[row] += 1;\n",
    "              break;\n",
    "            }\n",
    "          }\n",
    "        }\n",
    "        if (rowNnz[row] >= ellwidth) {\n",
    "          break;\n",
    "        }\n",
    "      }\n",
    "    }\"\"\").build()"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 133,
   "id": "0053ed70",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "<pyopencl._cl.NannyEvent at 0x7faafd3df680>"
      ]
     },
     "execution_count": 133,
     "metadata": {},
     "output_type": "execute_result"
    },
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "\n",
      "UPDATE[2,5]: 0.702566\n",
      "UPDATE[1,7]: 0.530861\n",
      "UPDATE[3,0]: 0.833768\n",
      "UPDATE[0,0]: 0.508427\n",
      "UPDATE[3,6]: 0.767162\n",
      "UPDATE[4,6]: 0.741744\n",
      "UPDATE[7,5]: 0.728992\n",
      "UPDATE[1,3]: 0.351743\n",
      "UPDATE[3,7]: 0.628391\n",
      "UPDATE[0,7]: 0.383189\n",
      "UPDATE[4,4]: 0.621514\n",
      "UPDATE[2,3]: 0.645422\n",
      "UPDATE[7,3]: 0.395720\n",
      "UPDATE[3,4]: 0.642813\n",
      "UPDATE[7,3]: 0.669699"
     ]
    }
   ],
   "source": [
    "knl = prg.addvals  # Use this Kernel object for repeated calls\n",
    "knl(queue, [topk,bs], None, adata_buf, acols_buf, annzs_buf, np.float32(1), np.uint32(ellwa), x_cp_buf, x_idx_buf, y_idx_buf)\n",
    "\n",
    "resa = np.empty_like(adata)\n",
    "resaidx = np.zeros(acols.shape).astype(np.uint32)\n",
    "resannz = np.zeros(annz.shape).astype(np.uint32)\n",
    "\n",
    "cl.enqueue_copy(queue, resa, adata_buf)\n",
    "cl.enqueue_copy(queue, resaidx, acols_buf)\n",
    "cl.enqueue_copy(queue, resannz, annzs_buf)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 134,
   "id": "8f781e72",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "9"
      ]
     },
     "execution_count": 134,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "ellwa"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 135,
   "id": "64a4fdb0",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[0.59460324, 0.8630502 , 0.63867456, 0.51579547, 0.7879012 ,\n",
       "        0.93278366, 0.02474353, 0.        , 0.        ],\n",
       "       [0.56595707, 0.53133106, 0.5421672 , 0.6889426 , 0.32183105,\n",
       "        0.64553165, 0.53694063, 0.        , 0.        ],\n",
       "       [0.80983955, 0.73776245, 0.7680449 , 0.7956841 , 0.9053589 ,\n",
       "        0.5579857 , 0.11461218, 0.        , 0.        ],\n",
       "       [0.09699463, 0.6963742 , 0.30319762, 0.51583916, 0.83718425,\n",
       "        0.05103909, 0.27944976, 0.        , 0.        ],\n",
       "       [0.75281155, 0.48989916, 0.8277786 , 0.6847082 , 0.9526682 ,\n",
       "        0.10713863, 0.6331825 , 0.        , 0.        ],\n",
       "       [0.17814207, 0.84142315, 0.9636037 , 0.7660027 , 0.79484075,\n",
       "        0.64250606, 0.16992201, 0.        , 0.        ],\n",
       "       [0.9762836 , 0.748014  , 0.25351423, 0.5921869 , 0.73024714,\n",
       "        0.80543023, 0.5730646 , 0.        , 0.        ],\n",
       "       [0.8948947 , 0.31771463, 0.5184244 , 0.06335784, 0.8011633 ,\n",
       "        0.60369205, 0.95093155, 0.        , 0.        ]], dtype=float32)"
      ]
     },
     "execution_count": 135,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "adata2 = adata.reshape(-1, ellwa)\n",
    "adata2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 136,
   "id": "3b8f6062",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[ 0.08617616,  0.8630502 ,  0.63867456,  0.51579547,  0.7879012 ,\n",
       "         0.93278366, -0.3584458 ,  0.        ,  0.        ],\n",
       "       [ 0.56595707,  0.53133106,  0.19042373,  0.6889426 ,  0.32183105,\n",
       "         0.64553165,  0.00607979,  0.        ,  0.        ],\n",
       "       [ 0.80983955,  0.73776245,  0.12262243,  0.7956841 ,  0.20279253,\n",
       "         0.5579857 ,  0.11461218,  0.        ,  0.        ],\n",
       "       [-0.736773  ,  0.6963742 ,  0.30319762,  0.51583916,  0.1943714 ,\n",
       "        -0.71612304, -0.3489409 ,  0.        ,  0.        ],\n",
       "       [ 0.75281155,  0.48989916,  0.8277786 ,  0.06319374,  0.9526682 ,\n",
       "        -0.63460505,  0.6331825 ,  0.        ,  0.        ],\n",
       "       [ 0.17814207,  0.84142315,  0.9636037 ,  0.7660027 ,  0.79484075,\n",
       "         0.64250606,  0.16992201,  0.        ,  0.        ],\n",
       "       [ 0.9762836 ,  0.748014  ,  0.25351423,  0.5921869 ,  0.73024714,\n",
       "         0.80543023,  0.5730646 ,  0.        ,  0.        ],\n",
       "       [ 0.8948947 ,  0.31771463,  0.5184244 , -1.0020614 ,  0.8011633 ,\n",
       "        -0.12530035,  0.95093155,  0.        ,  0.        ]],\n",
       "      dtype=float32)"
      ]
     },
     "execution_count": 136,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resa = resa.reshape(-1, ellwa)\n",
    "resa"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 137,
   "id": "eb51f891",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([[-0.5084271 ,  0.        ,  0.        ,  0.        ,  0.        ,\n",
       "         0.        , -0.38318932,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        , -0.35174346,  0.        ,  0.        ,\n",
       "         0.        , -0.53086084,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        , -0.64542246,  0.        , -0.7025664 ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ],\n",
       "       [-0.83376765,  0.        ,  0.        ,  0.        , -0.64281285,\n",
       "        -0.76716214, -0.62839067,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        ,  0.        , -0.62151444,  0.        ,\n",
       "        -0.7417437 ,  0.        ,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        ,  0.        ,  0.        ,  0.        ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        ,  0.        ,  0.        ,  0.        ,\n",
       "         0.        ,  0.        ,  0.        ,  0.        ],\n",
       "       [ 0.        ,  0.        ,  0.        , -1.0654192 ,  0.        ,\n",
       "        -0.7289924 ,  0.        ,  0.        ,  0.        ]],\n",
       "      dtype=float32)"
      ]
     },
     "execution_count": 137,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resa - adata2"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 138,
   "id": "f7e150fe",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([0, 1, 2, 4, 5, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4,\n",
       "       5, 6, 7, 0, 0, 0, 1, 2, 3, 4, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0,\n",
       "       0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4, 5, 6, 7, 0, 0, 0, 1, 2,\n",
       "       3, 4, 5, 6, 0, 0], dtype=uint32)"
      ]
     },
     "execution_count": 138,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "acols"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 139,
   "id": "8f1f18f1",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([0, 1, 2, 4, 5, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4,\n",
       "       5, 6, 7, 0, 0, 0, 1, 2, 3, 4, 6, 7, 0, 0, 0, 1, 3, 4, 5, 6, 7, 0,\n",
       "       0, 0, 1, 3, 4, 5, 6, 7, 0, 0, 0, 2, 3, 4, 5, 6, 7, 0, 0, 0, 1, 2,\n",
       "       3, 4, 5, 6, 0, 0], dtype=uint32)"
      ]
     },
     "execution_count": 139,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resaidx"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 140,
   "id": "bd326e73",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([7, 7, 7, 7, 7, 7, 7, 7], dtype=uint32)"
      ]
     },
     "execution_count": 140,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "resannz"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 141,
   "id": "357413c1",
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/plain": [
       "array([7, 7, 7, 7, 7, 7, 7, 7], dtype=uint32)"
      ]
     },
     "execution_count": 141,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "annz"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 130,
   "id": "e7fcc7fe",
   "metadata": {},
   "outputs": [
    {
     "ename": "NameError",
     "evalue": "name 'asdf' is not defined",
     "output_type": "error",
     "traceback": [
      "\u001b[0;31m---------------------------------------------------------------------------\u001b[0m",
      "\u001b[0;31mNameError\u001b[0m                                 Traceback (most recent call last)",
      "\u001b[0;32m/tmp/ipykernel_12222/2827374946.py\u001b[0m in \u001b[0;36m<module>\u001b[0;34m\u001b[0m\n\u001b[0;32m----> 1\u001b[0;31m \u001b[0masdf\u001b[0m\u001b[0;34m\u001b[0m\u001b[0;34m\u001b[0m\u001b[0m\n\u001b[0m",
      "\u001b[0;31mNameError\u001b[0m: name 'asdf' is not defined"
     ]
    }
   ],
   "source": [
    "asdf"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "7f618bfd",
   "metadata": {},
   "source": [
    "# OTHER"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "28fd6b45",
   "metadata": {},
   "source": [
    "import numpy as np\n",
    "import pyopencl as cl\n",
    "\n",
    "mf = cl.mem_flags"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "ebed13d1",
   "metadata": {},
   "source": [
    "dim = 16\n",
    "topk = 4\n",
    "\n",
    "x = np.random.rand(dim).astype(np.float32)\n",
    "y = np.random.rand(dim).astype(np.float32)\n",
    "x.shape,y.shape"
   ]
  },
  {
   "cell_type": "markdown",
   "id": "444a4a7f",
   "metadata": {},
   "source": [
    "dim1 = 4\n",
    "dim2 = 8\n",
    "dim3 = 1\n",
    "\n",
    "ctx = cl.create_some_context()\n",
    "queue = cl.CommandQueue(ctx,\n",
    "        properties=cl.command_queue_properties.PROFILING_ENABLE)\n",
    "\n",
    "sparsity = 0.2\n",
    "\n",
    "a = np.zeros((dim1,dim2))\n",
    "b = np.random.rand(dim2,dim3).flatten().astype(np.float32)\n",
    "\n",
    "a.shape, b.shape"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "e2d4594d",
   "metadata": {},
   "outputs": [],
   "source": [
    "x_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=x)\n",
    "y_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=y)\n",
    "val_out_buf = cl.Buffer(ctx, mf.READ_WRITE, 4*topk*topk)\n",
    "x_idx_buf = cl.Buffer(ctx, mf.READ_WRITE, topk*4)\n",
    "y_idx_buf = cl.Buffer(ctx, mf.READ_WRITE, topk*4)\n",
    "\n",
    "prg = cl.Program(ctx, \"\"\"\n",
    "// Every global_id_0 works on a row\n",
    "__kernel void genwupdate2(__global  float* x,     // INPUT MATRIX DATA\n",
    "                         __global  float* y,    // INPUT\n",
    "                         __global  float* xout,    // INPUT\n",
    "                         uint topk,\n",
    "                         __global  uint* xoutidx,    // INPUT\n",
    "                         __global  uint* youtidx    // INPUT\n",
    "                        ) { // LOCAL SHARED BUFFER\n",
    "  uint gid = get_global_id(0);\n",
    "  uint n = get_global_size(0);\n",
    "  \n",
    "  xout[gid] = x[gid];\n",
    "  xoutidx[gid] = gid;\n",
    "  youtidx[gid] = gid;\n",
    "  \n",
    "  float valx = x[gid];\n",
    "  float valy = y[gid];\n",
    "  uint posx = 0;\n",
    "  uint posy = 0;\n",
    "  for (uint i = 0; i < n; i++) {\n",
    "    float tempval = x[i];\n",
    "    float tempval2 = y[i];\n",
    "    bool larger = tempval > valx;\n",
    "    bool larger2 = tempval2 > valy;\n",
    "      \n",
    "    posx += (larger)?1:0;\n",
    "    posy += (larger2)?1:0;\n",
    "  }\n",
    "  //printf(\"posx:%i\", posx);\n",
    "  if (posx < topk) {\n",
    "    xoutidx[posx] = gid;\n",
    "  }\n",
    "  if (posy < topk) {\n",
    "    youtidx[posy] = gid;\n",
    "  }\n",
    "  if (gid < topk) {\n",
    "    uint i = gid;\n",
    "    for (uint j=0; j<topk; j++) {\n",
    "      xout[gid*topk+j] = x[xoutidx[gid]] * y[youtidx[j]];\n",
    "    }\n",
    "  }\n",
    "}\"\"\").build()"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "f45b3b69",
   "metadata": {},
   "outputs": [],
   "source": [
    "knl = prg.genwupdate2  # Use this Kernel object for repeated calls\n",
    "event = knl(queue, [dim,], None, x_buf, y_buf, val_out_buf, np.uint32(topk), x_idx_buf, y_idx_buf)\n",
    "\n",
    "#event.wait()\n",
    "val_out = np.zeros(topk*topk).astype(np.float32)\n",
    "resxidx = np.zeros(topk).astype(np.uint32)\n",
    "resyidx = np.zeros(topk).astype(np.uint32)\n",
    "\n",
    "cl.enqueue_copy(queue, val_out, val_out_buf)\n",
    "cl.enqueue_copy(queue, resxidx, x_idx_buf, wait_for=[event])\n",
    "cl.enqueue_copy(queue, resyidx, y_idx_buf)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "d883394f",
   "metadata": {},
   "outputs": [],
   "source": [
    "val_out"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "2335cf69",
   "metadata": {},
   "outputs": [],
   "source": [
    "resxidx"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "5e714f2a",
   "metadata": {},
   "outputs": [],
   "source": [
    "resyidx"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "02a5c5dd",
   "metadata": {},
   "outputs": [],
   "source": [
    "from __future__ import division\n",
    "\n",
    "KERNEL_CODE = \"\"\"\n",
    "// Thread block size\n",
    "#define BLOCK_SIZE %(block_size)d\n",
    "// Matrix dimensions\n",
    "// (chosen as multiples of the thread block size for simplicity)\n",
    "#define WA %(w_a)d // Matrix A width\n",
    "#define HA %(h_a)d // Matrix A height\n",
    "#define WB %(w_b)d // Matrix B width\n",
    "#define HB WA  // Matrix B height\n",
    "#define WC WB  // Matrix C width\n",
    "#define HC HA  // Matrix C height\n",
    "/*\n",
    " * Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.\n",
    " *\n",
    " * NVIDIA Corporation and its licensors retain all intellectual property and\n",
    " * proprietary rights in and to this software and related documentation.\n",
    " * Any use, reproduction, disclosure, or distribution of this software\n",
    " * and related documentation without an express license agreement from\n",
    " * NVIDIA Corporation is strictly prohibited.\n",
    " *\n",
    " * Please refer to the applicable NVIDIA end user license agreement (EULA)\n",
    " * associated with this source code for terms and conditions that govern\n",
    " * your use of this NVIDIA software.\n",
    " *\n",
    " */\n",
    "/* Matrix multiplication: C = A * B.\n",
    " * Device code.\n",
    " */\n",
    "#define AS(j, i) As[i + j * BLOCK_SIZE]\n",
    "#define BS(j, i) Bs[i + j * BLOCK_SIZE]\n",
    "////////////////////////////////////////////////////////////////////////////////\n",
    "//! Matrix multiplication on the device: C = A * B\n",
    "//! WA is A's width and WB is B's width\n",
    "////////////////////////////////////////////////////////////////////////////////\n",
    "__kernel __attribute__((reqd_work_group_size(16,16,1))) \n",
    "void\n",
    "matrixMul( __global float* C, __global float* A, __global float* B)\n",
    "{\n",
    "    __local float As[BLOCK_SIZE*BLOCK_SIZE];\n",
    "    __local float Bs[BLOCK_SIZE*BLOCK_SIZE];\n",
    "    // Block index\n",
    "    int bx = get_group_id(0);\n",
    "    int by = get_group_id(1);\n",
    "    // Thread index\n",
    "    int tx = get_local_id(0);\n",
    "    int ty = get_local_id(1);\n",
    "    // Index of the first sub-matrix of A processed by the block\n",
    "    int aBegin = WA * BLOCK_SIZE * by;\n",
    "    // Index of the last sub-matrix of A processed by the block\n",
    "    int aEnd   = aBegin + WA - 1;\n",
    "    // Step size used to iterate through the sub-matrices of A\n",
    "    int aStep  = BLOCK_SIZE;\n",
    "    // Index of the first sub-matrix of B processed by the block\n",
    "    int bBegin = BLOCK_SIZE * bx;\n",
    "    // Step size used to iterate through the sub-matrices of B\n",
    "    int bStep  = BLOCK_SIZE * WB;\n",
    "    // Csub is used to store the element of the block sub-matrix\n",
    "    // that is computed by the thread\n",
    "    float Csub = 0.0f;\n",
    "    // Loop over all the sub-matrices of A and B\n",
    "    // required to compute the block sub-matrix\n",
    "    for (int a = aBegin, b = bBegin;\n",
    "             a <= aEnd;\n",
    "             a += aStep, b += bStep) {\n",
    "        // Load the matrices from device memory\n",
    "        // to shared memory; each thread loads\n",
    "        // one element of each matrix\n",
    "        AS(ty, tx) = A[a + WA * ty + tx];\n",
    "        BS(ty, tx) = B[b + WB * ty + tx];\n",
    "        // Synchronize to make sure the matrices are loaded\n",
    "        barrier(CLK_LOCAL_MEM_FENCE);\n",
    "        // Multiply the two matrices together;\n",
    "        // each thread computes one element\n",
    "        // of the block sub-matrix\n",
    "        for (int k = 0; k < BLOCK_SIZE; ++k)\n",
    "            Csub += AS(ty, k) * BS(k, tx);\n",
    "        // Synchronize to make sure that the preceding\n",
    "        // computation is done before loading two new\n",
    "        // sub-matrices of A and B in the next iteration\n",
    "        barrier(CLK_LOCAL_MEM_FENCE);\n",
    "    }\n",
    "    // Write the block sub-matrix to device memory;\n",
    "    // each thread writes one element\n",
    "    C[get_global_id(1) * get_global_size(0) + get_global_id(0)] = Csub;\n",
    "}\n",
    "\"\"\"\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "c1eb509e",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "77b4a04e",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "a0521855",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "9a61206e",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "16d754cb",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "c91d5425",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "d9f4a6f0",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "9f708713",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "e59475d1",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "3dd6c3a3",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "7e83b426",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "af6af470",
   "metadata": {},
   "outputs": [],
   "source": []
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "id": "e2b3a039",
   "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.8.12"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 5
}