diff --git a/OpenCL to CUDA.ipynb b/OpenCL to CUDA.ipynb index 9cc0eb7..ff68603 100644 --- a/OpenCL to CUDA.ipynb +++ b/OpenCL to CUDA.ipynb @@ -106,26 +106,10 @@ "scrolled": false }, "outputs": [ - { - "name": "stderr", - "output_type": "stream", - "text": [ - "c:\\Users\\anbro\\Documents\\projects\\ShallowWaterGPU\\SWESimulators\\Common.py:28: UserWarning: The CUDA compiler succeeded, but said the following:\n", - "kernel.cu\r\n", - "c:\\users\\anbro\\documents\\projects\\shallowwatergpu\\swesimulators\\common.cu(837): warning: variable \"rv_1\" was declared but never referenced\r\n", - "\r\n", - "c:\\users\\anbro\\documents\\projects\\shallowwatergpu\\swesimulators\\common.cu(839): warning: variable \"rh_2\" was declared but never referenced\r\n", - "\r\n", - "c:\\users\\anbro\\documents\\projects\\shallowwatergpu\\swesimulators\\common.cu(843): warning: variable \"rv_3\" was declared but never referenced\r\n", - "\r\n", - "\n", - " kernel = cuda_compiler.SourceModule(kernel_string, include_dirs=[module_path])\n" - ] - }, { "data": { "text/plain": [ - "" + "" ] }, "execution_count": 4, @@ -174,10 +158,60 @@ }, { "cell_type": "code", - "execution_count": null, - "metadata": {}, - "outputs": [], - "source": [] + "execution_count": 10, + "metadata": { + "scrolled": false + }, + "outputs": [ + { + "data": { + "text/plain": [ + "" + ] + }, + "execution_count": 10, + "metadata": {}, + "output_type": "execute_result" + }, + { + "data": { + "image/png": "iVBORw0KGgoAAAANSUhEUgAAAZIAAADxCAYAAADhlTG6AAAABHNCSVQICAgIfAhkiAAAAAlwSFlzAAALEgAACxIB0t1+/AAAADl0RVh0U29mdHdhcmUAbWF0cGxvdGxpYiB2ZXJzaW9uIDIuMi4yLCBodHRwOi8vbWF0cGxvdGxpYi5vcmcvhp/UCwAAIABJREFUeJzt3X/wXXV95/HnixCCBJEfAUSh6m5Z3VQxrSm04+gIDBgYJ9jBWrCjiQuT4ixLa7cdoWuxQ7c7sm5Lf+jWBkyRqohmF40VCYG1q27FEmzkp0CkWEKQEALIDwnJ9/veP8653xxu7vd7P997zs099/N9PWbO5N5zP+fH/c775H0/5/M5n48iAjMzs0HtN+oTMDOz8eZEYmZmtTiRmJlZLU4kZmZWixOJmZnV4kRiZma1OJGYmVktTiRmZlaLE4mZmdWy/6hPwKyXd568MJ7YMZFU9vY7dq6PiGVDPiWzRqTG9jjFtROJtdITOyb4p/U/l1R23jEPLBry6Zg1JjW2xymunUislQKYZHLUp2HWuBxj24nEWikIdkXarS2zcZJjbDuRWGvl9qvNrCO32HYisVYKgglPcWAZyjG2nUistSbJ62Iz68gttp1IrJUCmMjsYjODPGPbicRaK7dfbWYducW2E4m1UgC7MruPbAZ5xrYTibVSENlV/80gz9h2IrF2CpjI61ozK2QY204k1krF079m+ckxtp1IrKXEBBr1SZgNQX6x7URirVQ0SOZ1sZlBnrHtRGKtVPS1z+tiM4M8Y9uJxFprMrNfbWYducW2E4m1Uo6/2swgz9h2IrFWCsSEZ4K2DOUY23l9G8vKZChp6UfSMkn3Sdos6eIen6+U9LikTeVyfrn+5Mq6TZJekPTuIXxVm2OaiOs2cY3EWikQL8a82vuRNA/4FHAasAW4TdK6iLinq+h1EXHhS84h4pvAknI/hwObgZtqn5TNaU3Fdps4kVgrFQ9tNVJhPhHYHBEPAkj6InAW0J1I+nkP8I2IeL6Jk7K5q8HYbo28vo1lZaJ8cKvfAiyStLGyrKrs5tXAw5X3W8p13c6WdIektZKO6/H5OcC1jX05m9MS43pstLJGcoAWxIEsHPVp2BC9wHO8GDunvVoixEQk/87ZHhFLp/ms1zG6Rzr6GnBtROyUdAHwWeCUqR1IxwBvAtannlAvjuu54Rme3B4RR073+SxjeyzUSiSSlgF/AcwDroqIj3d9vgC4BngL8ATwGxHxUL/9HshCTtKpdU7NWu57cUvfMpPN/CrbAlRrGMcCW6sFIuKJytsrgcu7Yvse4PqI2NUpNEhsO67nhptj7Y/7lWkotltj4LRYacQ8A1gMnCtpcVex84AnI+LngSuAywc9ns0tRYPk/klLH7cBx0t6naQDKG5RrasWKGscHcuBe3lpbJ8K/GPXfh3bNpDU2B4ndepXU42YEfEi0GnErDqL4jYBwFrgVEl5pWIbik6DZMoy434idgMXUtyWuhf4UkTcLekyScvLYhdJulvSD4CLgD9nTwP9q4Dd5b9Vjm0bSGpsj5M6aa9XI+ZJ05WJiN2SngaOALbXOK7NERMN9aWPiBuAG7rWXVp5fQlwSee9pPewJ24fknQhjm1rUFOx3RZ1EklKI2ZKmaJg0dNmFcCBHFTjtCwHI376t7HYdlxbtxyfbK+TSPo2YlbKbJG0P/AKYEevnUXEamA1wCE6PLP5w2wQk6Pr2dJYbDuurZcRxvZQ1Pk2fRsxy/crytfvAf5PRGaz3ttQFAPb7Ze0DIFj24YmNbbHycBnm9iI+RngCEmbgd8F9hrnyKyXQOyKeUlL48d2bNsQpcZ2ihrjyC2R9N2yk8kdkn6jss3Vkv6lss2SfudRq49ZQiPmC8Cv1zmGzU0RjPShLce2DUtTsV1nHDngeeADEfGApFcBt0taHxFPlZ//fkSsTT2X8eqsbHOIsntoy6zQWGwPPI5cRNxfeb1V0jbgSOCp6bea3njdiLM5Iyh+taUsZuMkNbYTNDKOnKQTgQOAH1VW/0m5zRXlKA4z8lVorTXCxnazoUqM65kGI4X0ceReGxEnADez5yHaYgfFqA5/B3wwIibL1ZcAbwB+GTgc+Ei/7+NbW9ZKwfhN7mOWYhaxPdNgpDDgOHKdN5IOAb4OfDQibq1s82j5cqekvwV+r9+JOpFYKwWwa8zGGzJL0WBsT3VTBx6h6Kb+vmoBScdUEkNnHDnKbu3XA9dExJd7bVMO+fNu4K5+J+Ir1Vpq/OZkMEvTTGyXQ/N0uqnPA9Z0uqkDGyNiHcU4csspxovbAawsN38v8HaKLuyddSsjYhPweUlHUtw62wRc0O9cnEislYL8nv41g2Zje7bjyFXWfw743DT7PKXX+pk4kVhruUZiucottp1IrJUi5BqJZSnH2HYisVYqGiSbH/7EbNRyjG0nEmup/Oa1NivkF9tOJNZKRYNkXveRzSDP2HYisdbyU+uWq9xi24nEWslPtluucoztvNKiZWWS/ZKWfgads6H87Ock3STpXkn3SHpto1/S5qQm4rpNXCOxVoqAXZMjn7MB4BrgTyJig6SDgckeZcySNRXbbeJEYq1UVP8budgGnrNB0mJg/4jYABARzzZxQja3NRjbrTHwt5F0nKRvllX+uyX9do8y75D0dOWWwaW99mXWy0Q5JlG/hZmH2x5kzoYTJX2TYgyjN0v6gaR/lvSJsobj2LZaEuN6bNSpkewG/nNEfF/SyymmatzQ45bBtyPiXTWOY3PQLLtIzjTcduqcDddGxE5JFwB/BlwE/BuKudkPoRg59Y8oBr37TLmdY9tmLcfuvwPXSCLi0Yj4fvn6GYrhiXv90jMbQFH9T1n6SJqzISJ2lm+vBN5UxvYW4J+BO4BXAl8BfqmBL2dzWlpsj5NGzrbsyfKLwPd6fPyr5a2Bb0j6hRn2sapza2IXO6crZnPIZDm3db+lj6k5G8o5GM4B1lULlLPEdUzN2VBuexTwForYPoWXtq30je2xi2up3UsmGojrVqnd2F72ZPlfwO9ExE+7Pv4+8JqIeFbSmRS/6I7vtZ+IWA2sBjhEh3fferA5pujZUn88oppzNryM4hrZDfw/4HaKGgskxrbj2ro1FdttUiuRSJpPkUQ+HxH/u/vzamKJiBsk/U9JiyJie53jWv6afGhrkDkbKrH96Yj4sx77dGzbQHJ8IHHgRFJOw/gZ4N5eF1pZ5pXAYxERkk6kuJX2RK+yZt1GVb13bNuwjdutq37q1EjeCrwfuFPSpnLdHwA/BxARnwbeA3xI0m7gZ8A5EeHqvfU14p4tjm0bmhx7bQ2cSCLiO/TuWlkt80ngk4Mew+a2UfVccWzbsI1br6x+/GS7tVKE2J3ZxWYGeca2E4m1Vm7Vf7OO3GLbicRaKcf7yGaQZ2w7kVhr5XaxmXXkFttOJNZKOfa1H5n90h9+2+/ABUnltCCtXKrYmfbU/+QLiaMDTE7UOJvhyjG2nUistXLra2/WkVtsO5FYK0XA7swm/zGDPGPbicRaK7fqv1lHbrHtRGKtlON9ZDPIM7bzql9ZViKUtJiNm6biWtIySfdJ2izp4h6fr5T0eGUmz/PL9Uskfbec3fYOSb9R2eZ1kr4n6QFJ15XTL8zIicRaq6H5SMxap4m4Lqd9/hRwBrAYOFfS4h5Fr4uIJeVyVbnueeADEfELwDLgzyUdWn52OXBFRBwPPAmc1+9cnEislSKK+8gpi9k4SY3tBCcCmyPiwYh4EfgicFbaOcT9EfFA+XorsA04shz5+hRgbVn0s8C7++3PbSTWUmIis54tZoXGYvvVwMOV91uAk3qUO1vS24H7gQ9HRHUbymkQDgB+BBwBPBURuyv77DuFuq9Uay23kViuEuN6UWea5nJZ1bWbXsHfPZXB14DXRsQJwM0UNYw9Oyimmf474IMRMZm4z724RjJkmt+3nWpWYteLje6vrXIcj6hxiU+szzvk4PR9HnNUUrFdi2axzwTztz+bVG7eo9uSyk38NG1/wD5/Cn4Wsb09IpbO8PkW4LjK+2OBrS85VkR1srUrKdo/AJB0CPB14KMRcWvnmMChkvYvayV77bMX10isnaK4l5yymI2V5uL6NuD4spfVAcA5wLpqgbLG0bEcuLdcfwBwPXBNRHx56tSKydm+STFxG8AK4Kv9TqR2IpH0kKQ7y65lG3t8Lkl/WXZPu0PSL9U9ps0NTfXaGqSLZCWuQ9Lz5fp1lW0c1zawJuK6rDFcCKynSBBfioi7JV0maXlZ7KKyi+8PgIuAleX69wJvB1ZW4n5J+dlHgN+VtJmizeQz/c6lqVtbJ0fE9mk+OwM4vlxOAv6a3g1CZlOioQbJShfJ0yhuBdwmaV1E3NNV9LqIuLCy3UeBk4GHIqLXfRzHtQ2kqdgGiIgbgBu61l1aeX0JcEmP7T4HfG6afT5I0SMs2b64tXUWRfUpyvtwh3ZVt8x6aujW1sBdJPtwXNvAcrtl20QiCeAmSbf36FUAvbuo9e1OZjaLXlsz9W5Jjb+zy1tUayUdRxnXwMLyNtetkqr96R3XNrDceiM2cWvrrRGxVdJRwAZJP4yIb1U+T+pOVl78qwAO5KAGTsvGWfGrLPlimql3S2oXyWsjYqekCyi6SHbi+k0UtwCuoHj6986I+FHifh3XtpdZxvZYqF0jKZ+KJCK2UfQC6L631reLWrn96ohYGhFL59PspDk2nhp6sj2pi2REdGZMuhJ4SyWu76SI6+OAfwB+MXW/5faOa9tLbiM21EokkhZKennnNXA6cFdXsXXAB8peLr8CPB0Rj9Y5rs0NDbWRDNJF8j5JL5d0mKTDKOL6X4G3Ap1Gese1DSy3NpK6t7aOBq4vhmdhf+ALEXFjeXuAiPg0RY+CM4HNFAOFfbDmMW0OCMRkAz1bImK3pE4XyXnAmk4XSWBjRKyj6CK5HNgN7AD+APgOcBBFreNJ4A+B71J0mbwHx7UNqKnYbpNaiaTsJvbmHus/XXkdwH+sc5w2Sp2zevKX3tDocff7/g+TyqXOgd1mTf0oG7CL5F5x3bX9yONa8xMv36OPTN7nE7+8KKncjjcm7zLJ4XcdmFTuiO+lRcV+qXO7A5Mv7Pv53ceswtGXh0ixdsqwQdIMyDK2nUisvXL72WbWkVlsO5FYa+X2q82sI7fYdiKxVgpgcjKvi80M8oxtJxJrpwAy+9VmBmQZ204k1lrj1pfeLFVuse1EYu2V2cVmNiWz2HYisZYav4HrzNLkF9tOJNZemf1qM5uSWWw7kQxo3ivT5rZ+9CPNzrF+7G+nHXf3jx/uX6jNAiKzni1N2y9xdIWdr35F8j4fP2kyqdyH33Fj8j5TXLFwWVK5g7ekfZcFiXO7A/DCC+llm5BhbDuRWIvldbGZ7ZFXbDuRWHtlVv03m5JZbDuRWHtldrGZTckstp1IrJ0yfGjLDMgytp1IrLVye2jLrCO32HYisfbKrGeL2ZTMYnvgabokvV7SpsryU0m/01XmHZKerpS5dLr9mXVTpC2NH9exbUM2irgepoFrJBFxH7AEQNI84BHg+h5Fvx0R7xr0ODZHBSNrkHRs21CNMLaHpamJg08FfhQRP25ofzbnqWiQTFn67UlaJuk+SZslXdzj85WSHq/ULs6vfHwq8BDwj5I+2djXszmsmbhuk6baSM4Brp3ms1+V9ANgK/B7EXF3Q8ccqclXLEwq94MTp/uzDOaMV5zT6P5arYFfbWWN4lPAacAW4DZJ6yLinq6i10XEhT12cQ7wM+D/9vhsOLGtxP9E5s1LKrZ7YVo5gIOOeTap3H86rNnfjH+TeNzdC9OuuwWJfxsg/e/dZAu5ayQvJekAYDnw5R4ffx94TUS8Gfgr4Csz7GeVpI2SNu5iZ93TshxMJi4zOxHYHBEPRsSLwBeBs1IOX8b2rwFPATd1fZwU245r66l+XAP1atuSbpT0lKS/79rmakn/UtlmSb/zaOLW1hnA9yPise4PIuKnEfFs+foGYL6kRb12EhGrI2JpRCydT9oYQpaxTl/7+rcAXg1UBx7bUq7rdrakOyStlXRcue7M8kwu2uv0EmPbcW17SY3tPiq17TOAxcC5khb3KHpdRCwpl6sq6z8BvH+a3f9+ZZtN/c6liURyLtPc1pL0SqmoN0o6sTzeEw0c0+aAWfTaWtT51V8uq6q76bHr7hsLXwNeGxEnADcDny3X/yFwc0TsNQKmY9vqaKjX1sC1bYCIuAV4ZqAv0KVWG4mkgyjuPf9WZd0FABHxaeA9wIck7aa4z3xORG6P4tjQpEfK9ohYOs1nW4DjKu+PpWjT2HOYiGoCuBK4vIztNwJHSXoIOBhYKOnnI2IZjm2ro5lI6VXbPqlHubMlvR24H/hwrx9GPfxJ2aX9FuDiiJjxvmytRBIRzwNHdK37dOX1JwH3dLFRug04XtLrKLrxngO8r1pA0jER8Wj5djlwbxnbCyplVgJLOw3yjm3bBxZJ2lh5vzoiVlfep9a2r42IneWP/M8Cp/Q57iXAT4ADgNXAR4DLZtrAT7ZbazXxUFZE7JZ0IbAemAesiYi7JV0GbIyIdcBFkpYDu4EdwMr6RzabXmJsz1TThgFr2/0OWvlRtVPS3wK/128bJxJrp6CxYSTKxvAbutZdWnl9CcWvsJn2cTVwdSMnZHNbc7E9UG27304725RtgO8G7uq3jROJtZdbHCxXLahtS/o28AbgYElbgPMiYj3weUlHUtw62wRc0O9cnEistcZtvCGzVE3Fdp3adkS8bZr1/dpQ9uJEYu3lRGK5yiy2nUgGtN/TzyWVe/M/ndvocY99Om0oicQHY9sts4stWWov4omJpGL7P5dWDuD5Rw9OKvdXT74mqdx+SovE536SNvTJ0anfJfFvA4xmcpDMYtuJxFppHIfSNkuRY2w7kVh7ZTb5j9mUzGLbicRaK7dfbWYducW2E4m1V2YXm9mUzGLbicTaKcP7yGZAlrHtRGLtldnFZjYls9h2IrHWSuw5ajZ2covtpuZsNzOzOco1EmuvzKr/ZlMyi20nkgFN/GRbUrmjLz98JMcdexk2SDYtXnwxqdyCR55O3ueR3zsyqdxfPndm8j6Tjnt3WrkFj+xIKpf6txmJDGPbicTaK7OLzWxKZrGd1EYiaY2kbZLuqqw7XNIGSQ+U/x42zbYryjIPSFrR1InbHBCJy4Ac1zYyQ4zrUUhtbL8aWNa17mLglog4nnJe3+6NJB0OfIxiHuETgY9Nd2GaVYmiZ0vKUsPVOK5tH0uN7XGSlEgi4lsUk6JUnUUx/y/lv+/usek7gQ0RsSMingQ2sPeFa7a32DO4Xb9l4EM4rm0UhhzXo1CnjeTozhSO5bSMR/Uo82rg4cr7LeW6vUhaBawCOJCDapyWZWM0F5Pj2oZvzBJFP8N+jqTXEJc9/4QRsToilkbE0vksGPJp2VhoqI1E0jJJ90naLGmvW1XAoZIel7QJeLmk88vtXlO+3yTpbkmdKUcd11bPHG0j6eUxScdAMVk80Ktf6hbguMr7Y4GtNY5pc0gTt7YkzQM+BZwBLAbOlbS4R9HrImIJsBn4erkugAfL9ScBF0t6FY5rqym3W1t1Esk6oNNbZQXw1R5l1gOnSzqsbIw8vVxn1l8zNZITgc0R8WBEvAh8kaIdZDrVuH4f8JXy9QL2XC+Oa6tnLtZIJF0LfBd4vaQtks4DPg6cJukB4LTyPZKWSroKICJ2AH8M3FYul5XrzGYWjfXamrY9oxLXrwQ+JOlh4PXAuypxfY2kO4BHgIcjYqvj2mpJjO1xktTYHhHTTTx+ao+yG4HzK+/XAGsGOrsWi507k8rtt/HeZo+7q8VP7DYt/VfZIkkbK+9XR8Tq8vW07RmduJZ0BPBsROws20HeW3b/7TihvKX1FUlHR8RjbYjryRd3JZWb99jjyfs84ra0mfsO+XHa3O6p5m97Nq1g4ndJ/duMzJjVOPrxk+3WWrO4T7w9IpZO81nf9oyIeKLy9krg8u6dRMRWSXcDbwPWJp+ZWQ/j1gbSj0f/tfZqpo3kNuB4Sa+TdABwDkU7yJROp5HScuDecv2xkl5Wvj4MeCtwX63vZAbZtZG4RmLt1NDFFBG7JV1I0Rg+D1gTEXdLugzYGBHrgIskLQd2UzyguLLc/N8DfyopKG6R/Y+IuLP+WdmcNoaJoh8nEmsl0Vz1PyJuAG7oWndp5fUlwCU9ttsAnNDMWZgVmozttnAisdbK7WIz68gttp1IrL0yu9jMpmQW225st/ZqaIgUs9ZpKK77Df8jaWVn+J9yOb/y2Y2SnpL0913bvE7S98opEq4rO6nMyInE2ilxeJTcbhHYHNBQXM92+J9yuaqy/hPA+3uUvxy4onyW6kngvH7n4kRi7eUaieWqmbie7fA/Lz2FiFuAZ6rrJAk4hT3PSk03lcJLuI1kyObUk+gNG7dhIva5yYmkYhM/TXxqHNgv8Ynw+T9pdiTj1JEiJl9IK5f6txmVhmK71/A/J/Uod7aktwP3Ax+OiId7lOk4AngqInZX9tlzioQq10istXxry3KVGNeLJG2sLKu6d9Nj191XxNeA10bECcDN7Jm0bdpTS9jnXlwjsXbybSvLVXpszzT0DzQ0/E/3MSnm59m/rJUkTZHgGom1l9tILFfNxPXAw/9Me1oRAXwTeE+5agW9pwh5CddIrJVyfPrXDJqL7ZrD/yDp28AbgIMlbQHOi4j1wEeAL0r6r8A/A5/pdy5OJNZamnQmsTw1FduDDv9Tfva2adY/SNEjLJkTibWTb1tZrjKM7b5tJJLWSNom6a7Kuk9I+qGkOyRdL+nQabZ9SNKd5ROVG3uVMZvOsHttObZtVHLrjZjS2H41sKxr3QbgjWWXsvuZpupUOrl8onKm3gdmext+Y/vVOLZtFDLrRNI3kUTEtygaaarrbqo8sHIrRRcxs0YNu0bi2LZRya1G0kQbyX8ArpvmswBuKicG+pvKPNpm/Y3+YsojtmfxlPfkz36WVjC1XNNi9EHRiEy+RketRCLpv1B0K/v8NEXeWs51fRSwQdIPy1+Bvfa1ClgFcCAH1Tkty0GMdoiUpmLbcW17GXFsD8PADyRKWgG8C/jN8iGWvUTE1vLfbcD1zNClLCJWR8TSiFg6n2bH8bHx0+lrP4pbAE3GtuPauqXG9jgZKJFIWkbx0MryiHh+mjILJb288xo4HbirV1mzniLSlgY5tm2f2MdxPWx9b21JuhZ4B8UAYluAj1H0ZFlAUaUHuDUiLpD0KuCqiDgTOBq4vvx8f+ALEXHjUL6FZampX2VlcvgLiqd/r4qIj5frO7F9pKQJ4BGKuH2OIrZfBrwCeILih+TOiPhFHNtW07jVOPrpm0gi4tweq3s+Ml9W988sXz8IvLnW2dnc1VAXyMrkP6dRDHJ3m6R1EXFPJ7YlrQSWRsSFXdv+O4rhhx4ofyTdLulQx7bVMobde/vxk+3WWg01SE5N/gMgqTP5zz39NoyI+yuvt0raBhwJPNXImdmc5cZ2s31Ek2lLH70m/+k1Uc/Z5dPsayUd1/2hpBOBA4AfDfh1zKY0ENet4kRi7RTMprF9pgmAak/+Uw7F/XfAByNizC5xa53U2B4jvrVlrTWLBsmZJgCqNfmPpEOArwMfjYhbk8/IbAZzrrHdbGSaudimJv+h6JV1DvC+agFJx0TEo+Xbqcl/ysmCrgeuiYgvN3I2bTdmv4THVmZ/ZicSa6WWTP7zXuDtwBFlzy6AlRGxqf6Z2VyV46RtTiTWThEjn/wnIj4HfK6RkzDraDC228KJxNorr2vNbI/MYtuJxFort+q/WUduse1EYu0UQGbVfzMgy9h2IrH2yutaM9sjs9h2IrHWyq36b9aRW2w7kVhr5dazxawjt9h2IrF2ynCEVDMgy9h2IrFWKh7ayuxqMyPP2HYisfby8IiWq8xiu+/ov5LWSNom6a7Kuj+S9IikTeVy5jTbLpN0n6TNki5u8sQtf4pIWgbev2PbRmSYcT0KKcPIXw0s67H+iohYUi43dH9YmZnuDGAxcK6kxXVO1uaQmMUyuKtxbNu+Nvy43uf6JpKI+BbFQHazNTUzXUS8CHRmpjNLUIxHlLIMfATHto1Ec3Hdr2YsaaWkxys17PMrn62Q9EC5rKis/4dyn51tjup3HnUmtrqwnFFujaTDenyeOjOdWW/pE1s1zbFtw9VAXM+iZnxdpYZ9Vbnt4cDHgJMofhh9rCvWf7OyzbZ+5zJoIvlr4N8CS4BHgT/tUSZlZro9haVVnRnudrFzwNOybERjU+3OVqOx7bi2vSTGdoI6NeN3AhsiYkdEPAlsoPdt3iQDJZKIeCwiJsppR6+k+ELd+s5M17XP1RGxNCKWzmfBIKdluRlBjaTp2HZcW0/1p5CG9Jrx2WUNe62kTtz22/Zvy9tafyip1w+nlxgokZRzWHf8GnBXj2JTM9OVM82dA6wb5Hg2R42gUdKxbftEWlxv7/wIKZfVXXtJqRl/DXhtRJwA3Ax8NmHb34yINwFvK5f39/s6Kd1/rwW+C7xe0hZJ5wH/XdKdku4ATgY+XJZ9laQboJiZDujMTHcv8KWIuLvf8cw6NDmZtAy8f8e2jUhDcd23ZhwRT0RE557qlcBb+m0bEY+U/z4DfIHetfKX6PtAYkSc22P1Z6YpuxU4s/J+r5npzJIEQ39oy7FtI9FcbE/VjIFHKGrG76sWkHRMRDxavl1O8cMHih9B/63SwH46cImk/YFDI2K7pPnAuyhqMjPyk+3WSmL8HsoyS9FUbEfEbkmdmvE8YE1E3C3pMmBjRKwDLpK0HNhN0dV9ZbntDkl/TJGMAC4r1y0E1pdJZB5FErmy37k4kVh7NZRIJC0D/oLiwrgqIj7e9flK4BMUv+oAPlnpJnkj8CvAdyLiXY2ckFlDsd2rZhwRl1ZeXwJcMs22a4A1XeueY8/tr2ROJNZeDVxslb72p1HcF75N0rqIuKer6HURcWGPXXwCOAj4rdonY9aRWW27zgOJZsPTuY+cssys1lPoEXEL8Mwsz95seqmxPUacSKy1Guq1VaevvdlQDLM34ig4kVhLJT6M2P/BrTp97c2GYGRD/wyN20isnYLZXEzbI2LpNJ8l9bWvvL0SuDz1wGazNrvYHguukVjOxVI3AAAEqUlEQVR7NdNG0vcp9K6n2at97c2GI7M2EtdIrLVG3dceQNK3gTcAB0vaApwXEetrn5jNabk9I+VEYu3Vjr72b2vkJMyqnEiG7xme3H5zrP1x1+pFwPZRnE/DcvkeUO+7vGbGTyNgYszq931kHtfg79Ix52K7lYkkIo7sXidp4wwNqmMjl+8B++C7ZParLee4Bn+XWckstluZSMyA7C42symZxbYTibVTADXmYzdrrQxje5wSSfekLuMql+8BQ/0uAZHXfeRpOB7aybE9C2OTSHrMDjaWcvkeMOTvEmTXINmL46GdHNuzMzaJxOagzO4jm03JLLadSKy9MrvYzKZkFtutHyJF0jJJ90naLOniUZ9PHZIeKucD3yRp46jPZzYkrZG0TdJdlXWHS9og6YHy38Nm2sfszGrQxrHk2G6H1sb2GGl1IqlMSnQGsBg4V9Li0Z5VbSdHxJIx7G9/NbCsa93FwC0RcTxwS/m+GQFMTqYtY8ix3SpX08bYHiOtTiTUnJTImhMR36IYh6rqLPYMuf5Z4N0NHzS7X24Vju2WaG1sj5G2J5LUSYnGRQA3Sbq9a86McXV0RDwKUP57VHO7LoeRSFnGk2O73UYf22Ok7Y3tKZMSjZO3RsRWSUcBGyT9sPw1ZN0CIrO+9l0c23NVhrHd9hpJ30mJxklEbC3/3QZcT3F7Y5w91pnLo/x3W6N7n4y0ZTw5tttt9LE9RtqeSPpOSjQuJC2U9PLOa+B04K6Zt2q9dcCK8vUK4KuN7j3vNhLHdruNPrbHSKtvbU03KdGIT2tQRwPXS4Li7/6FiLhxtKeUTtK1wDso5kffAnwM+DjwJUnnAf8K/HpjB4wYu54rs+HYbg/Hdn2tTiTQe1KicRQRDwJvHvV5DCoizp3mo1OHeNCh7boNHNvt4Niur/WJxOaqICYmRn0SZkOQX2w7kVg7ZTjUthmQZWy3vbHd5rKYTFv66DcUiaSVkh4vh/fYJOn8ymcrymEyHpC0ontbs4E0ENcwnNiW9JZyuJvNkv5SZePXTFwjsVYKIBr41VYZiuQ0ii63t0laFxH3dBW9LiIu7Nr2cIqG16XlKd1ebvtk7ROzOWsMYvuvgVXArRRteMuAb8x0Lq6RWDtFNFUjqTMUyTuBDRGxo7zANrD3mExms5Ma2/01HtvlMzOHRMR3IyKAa0gYHsaJxForJiaSlj5ShyI5W9IdktZK6jwomNswJtYSDcQ1DCe2X12+7rfPl/CtLWulZ3hy/c2xdlFi8QO7hi5fXZnhLmUokq8B10bETkkXUAzSd0ritmazMovYnimuYTixPVDMO5FYK0VEU7eQ+g5FEhFPVN5eCVxe2fYdXdv+Q0PnZXNUy2N7S/l62n324ltblru+Q5F0xlQqLQfuLV+vB06XdFg5sdHp5TqzNmg8tsuRjp+R9Ctlb60PkDA8jGsklrXphiKRdBmwMSLWARdJWg7sppiXYmW57Q5Jf0xxwQJcFhHd81aYjcQQY/tDFJN9vYyit9aMPbYAFJk9qm9mZvuWb22ZmVktTiRmZlaLE4mZmdXiRGJmZrU4kZiZWS1OJGZmVosTiZmZ1eJEYmZmtfx/ycJNLd4sbvMAAAAASUVORK5CYII=\n", + "text/plain": [ + "
" + ] + }, + "metadata": {}, + "output_type": "display_data" + } + ], + "source": [ + "from SWESimulators import FORCE\n", + "importlib.reload(FORCE)\n", + "\n", + "nx = 10\n", + "ny = 15\n", + "num_ghost_cells = 1\n", + "dt = 0.01\n", + "g = 9.81\n", + "\n", + "h0, hu0, hv0, dx, dy, nx, ny = gen_test_data(nx, ny, num_ghost_cells)\n", + "plt.figure()\n", + "plt.subplot(121)\n", + "plt.imshow(h0)\n", + "plt.colorbar()\n", + "\n", + "sim = FORCE.FORCE(h0, hu0, hv0, \\\n", + " nx, ny, \\\n", + " dx, dy, dt, \\\n", + " g)\n", + "\n", + "t = sim.step(0.02)\n", + "h1, hu1, hv1 = sim.download()\n", + "\n", + "plt.subplot(122)\n", + "plt.imshow(h1)\n", + "plt.colorbar()" + ] }, { "cell_type": "code", diff --git a/SWESimulators/FORCE.py b/SWESimulators/FORCE.py index 926977b..a935ac4 100644 --- a/SWESimulators/FORCE.py +++ b/SWESimulators/FORCE.py @@ -22,7 +22,11 @@ along with this program. If not, see . #Import packages we need import numpy as np -import pyopencl as cl #OpenCL in Python + +import pycuda.compiler as cuda_compiler +import pycuda.gpuarray +import pycuda.driver as cuda + from SWESimulators import Common @@ -53,24 +57,26 @@ class FORCE: g: Gravitational accelleration (9.81 m/s^2) """ def __init__(self, \ - cl_ctx, \ h0, hu0, hv0, \ nx, ny, \ dx, dy, dt, \ g, \ block_width=16, block_height=16): - self.cl_ctx = cl_ctx - - #Create an OpenCL command queue - self.cl_queue = cl.CommandQueue(self.cl_ctx) + #Create a CUDA stream + self.stream = cuda.Stream() #Get kernels - self.kernel = Common.get_kernel(self.cl_ctx, "FORCE_kernel.opencl", block_width, block_height) + self.force_module = Common.get_kernel("FORCE_kernel.cu", block_width, block_height) + self.force_kernel = self.force_module.get_function("FORCEKernel") + self.force_kernel.prepare("iiffffPiPiPiPiPiPi") #Create data by uploading to device ghost_cells_x = 1 ghost_cells_y = 1 - self.cl_data = Common.SWEDataArkawaA(self.cl_ctx, nx, ny, ghost_cells_x, ghost_cells_y, h0, hu0, hv0) + self.data = Common.SWEDataArakawaA(nx, ny, \ + ghost_cells_x, ghost_cells_y, \ + h0, hu0, hv0, \ + stream=self.stream) #Save input parameters #Notice that we need to specify them in the correct dataformat for the @@ -86,7 +92,7 @@ class FORCE: self.t = np.float32(0.0) #Compute kernel launch parameters - self.local_size = (block_width, block_height) + self.local_size = (block_width, block_height, 1) self.global_size = ( \ int(np.ceil(self.nx / float(self.local_size[0])) * self.local_size[0]), \ int(np.ceil(self.ny / float(self.local_size[1])) * self.local_size[1]) \ @@ -109,20 +115,20 @@ class FORCE: if (local_dt <= 0.0): break - self.kernel.swe_2D(self.cl_queue, self.global_size, self.local_size, \ + self.force_kernel.prepared_async_call(self.global_size, self.local_size, self.stream, \ self.nx, self.ny, \ self.dx, self.dy, local_dt, \ self.g, \ - self.cl_data.h0.data, self.cl_data.h0.pitch, \ - self.cl_data.hu0.data, self.cl_data.hu0.pitch, \ - self.cl_data.hv0.data, self.cl_data.hv0.pitch, \ - self.cl_data.h1.data, self.cl_data.h1.pitch, \ - self.cl_data.hu1.data, self.cl_data.hu1.pitch, \ - self.cl_data.hv1.data, self.cl_data.hv1.pitch) + self.data.h0.data.gpudata, self.data.h0.pitch, \ + self.data.hu0.data.gpudata, self.data.hu0.pitch, \ + self.data.hv0.data.gpudata, self.data.hv0.pitch, \ + self.data.h1.data.gpudata, self.data.h1.pitch, \ + self.data.hu1.data.gpudata, self.data.hu1.pitch, \ + self.data.hv1.data.gpudata, self.data.hv1.pitch) self.t += local_dt - self.cl_data.swap() + self.data.swap() return self.t @@ -131,5 +137,5 @@ class FORCE: def download(self): - return self.cl_data.download(self.cl_queue) + return self.data.download(self.stream) diff --git a/SWESimulators/FORCE_kernel.opencl b/SWESimulators/FORCE_kernel.cu similarity index 70% rename from SWESimulators/FORCE_kernel.opencl rename to SWESimulators/FORCE_kernel.cu index 4dfc8ba..1c5af82 100644 --- a/SWESimulators/FORCE_kernel.opencl +++ b/SWESimulators/FORCE_kernel.cu @@ -19,14 +19,15 @@ along with this program. If not, see . */ -#include "common.opencl" +#include "common.cu" /** * Computes the flux along the x axis for all faces */ -void computeFluxF(__local float Q[3][block_height+2][block_width+2], - __local float F[3][block_height+1][block_width+1], +__device__ +void computeFluxF(float Q[3][block_height+2][block_width+2], + float F[3][block_height+1][block_width+1], const float g_, const float dx_, const float dt_) { //Index of thread within block @@ -40,12 +41,12 @@ void computeFluxF(__local float Q[3][block_height+2][block_width+2], const int k = i; // Q at interface from the right and left - const float3 Qp = (float3)(Q[0][l][k+1], - Q[1][l][k+1], - Q[2][l][k+1]); - const float3 Qm = (float3)(Q[0][l][k], - Q[1][l][k], - Q[2][l][k]); + const float3 Qp = make_float3(Q[0][l][k+1], + Q[1][l][k+1], + Q[2][l][k+1]); + const float3 Qm = make_float3(Q[0][l][k], + Q[1][l][k], + Q[2][l][k]); // Computed flux const float3 flux = FORCE_1D_flux(Qm, Qp, g_, dx_, dt_); @@ -54,15 +55,16 @@ void computeFluxF(__local float Q[3][block_height+2][block_width+2], F[2][j][i] = flux.z; } } - barrier(CLK_LOCAL_MEM_FENCE); + __syncthreads(); } /** * Computes the flux along the y axis for all faces */ -void computeFluxG(__local float Q[3][block_height+2][block_width+2], - __local float G[3][block_height+1][block_width+1], +__device__ +void computeFluxG(float Q[3][block_height+2][block_width+2], + float G[3][block_height+1][block_width+1], const float g_, const float dy_, const float dt_) { //Index of thread within block const int tx = get_local_id(0); @@ -76,12 +78,12 @@ void computeFluxG(__local float Q[3][block_height+2][block_width+2], // Q at interface from the right and left // Note that we swap hu and hv - const float3 Qp = (float3)(Q[0][l+1][k], - Q[2][l+1][k], - Q[1][l+1][k]); - const float3 Qm = (float3)(Q[0][l][k], - Q[2][l][k], - Q[1][l][k]); + const float3 Qp = make_float3(Q[0][l+1][k], + Q[2][l+1][k], + Q[1][l+1][k]); + const float3 Qm = make_float3(Q[0][l][k], + Q[2][l][k], + Q[1][l][k]); // Computed flux // Note that we swap back @@ -91,24 +93,24 @@ void computeFluxG(__local float Q[3][block_height+2][block_width+2], G[2][j][i] = flux.y; } } - barrier(CLK_LOCAL_MEM_FENCE); + __syncthreads(); } -__kernel void swe_2D( +__global__ void FORCEKernel( int nx_, int ny_, float dx_, float dy_, float dt_, float g_, //Input h^n - __global float* h0_ptr_, int h0_pitch_, - __global float* hu0_ptr_, int hu0_pitch_, - __global float* hv0_ptr_, int hv0_pitch_, + float* h0_ptr_, int h0_pitch_, + float* hu0_ptr_, int hu0_pitch_, + float* hv0_ptr_, int hv0_pitch_, //Output h^{n+1} - __global float* h1_ptr_, int h1_pitch_, - __global float* hu1_ptr_, int hu1_pitch_, - __global float* hv1_ptr_, int hv1_pitch_) { + float* h1_ptr_, int h1_pitch_, + float* hu1_ptr_, int hu1_pitch_, + float* hv1_ptr_, int hv1_pitch_) { //Index of thread within block const int tx = get_local_id(0); @@ -122,8 +124,8 @@ __kernel void swe_2D( const int ti = get_global_id(0) + 1; //Skip global ghost cells, i.e., +1 const int tj = get_global_id(1) + 1; - __local float Q[3][block_height+2][block_width+2]; - __local float F[3][block_height+1][block_width+1]; + __shared__ float Q[3][block_height+2][block_width+2]; + __shared__ float F[3][block_height+1][block_width+1]; //Read into shared memory @@ -131,7 +133,7 @@ __kernel void swe_2D( hu0_ptr_, hu0_pitch_, hv0_ptr_, hv0_pitch_, Q, nx_, ny_); - barrier(CLK_LOCAL_MEM_FENCE); + __syncthreads(); //Save our input variables @@ -142,23 +144,21 @@ __kernel void swe_2D( //Set boundary conditions noFlowBoundary1(Q, nx_, ny_); - barrier(CLK_LOCAL_MEM_FENCE); + __syncthreads(); //Compute flux along x, and evolve computeFluxF(Q, F, g_, dx_, dt_); - barrier(CLK_LOCAL_MEM_FENCE); evolveF1(Q, F, nx_, ny_, dx_, dt_); - barrier(CLK_LOCAL_MEM_FENCE); + __syncthreads(); //Set boundary conditions noFlowBoundary1(Q, nx_, ny_); - barrier(CLK_LOCAL_MEM_FENCE); + __syncthreads(); //Compute flux along y, and evolve computeFluxG(Q, F, g_, dy_, dt_); - barrier(CLK_LOCAL_MEM_FENCE); evolveG1(Q, F, nx_, ny_, dy_, dt_); - barrier(CLK_LOCAL_MEM_FENCE); + __syncthreads(); //Write to main memory writeBlock1(h1_ptr_, h1_pitch_,