{ "cells": [ { "cell_type": "code", "execution_count": 1, "metadata": { "nbsphinx": "hidden" }, "outputs": [], "source": [ "import psutil\n", "\n", "import pystencils as ps\n", "import pystencils.plot as psplot\n", "import numpy as np\n", "import sympy as sp\n", "\n", "import shutil" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "# Demo: Finite differences - 2D wave equation\n", "\n", "In this tutorial we show how to use the finite difference module of *pystencils* to solve a 2D wave equations. The time derivative is discretized by a simple forward Euler method." ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "$$ \\frac{\\partial^2 u}{\\partial t^2} = \\mbox{div} \\left( q(x,y) \\nabla u \\right)$$" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "We begin by creating three *numpy* arrays for the current, the previous and the next timestep. Actually we will see later that two fields are enough, but let's keep it simple. From these *numpy* arrays we create *pystencils* fields to formulate our update rule." ] }, { "cell_type": "code", "execution_count": 2, "metadata": {}, "outputs": [], "source": [ "size = (60, 70) # domain size\n", "u_arrays = [np.zeros(size), np.zeros(size), np.zeros(size)]\n", "\n", "u_fields = [ps.Field.create_from_numpy_array(\"u%s\" % (name,), arr)\n", " for name, arr in zip([\"0\", \"1\", \"2\"], u_arrays)]\n", "\n", "# Nicer display for fields\n", "for i, field in enumerate(u_fields):\n", " field.latex_name = \"u^{(%d)}\" % (i,)" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "*pystencils* contains already simple rules to discretize the a diffusion term. The time discretization is done manually." ] }, { "cell_type": "code", "execution_count": 3, "metadata": {}, "outputs": [ { "data": { "text/latex": [ "$\\displaystyle \\frac{{u^{(0)}}_{(0,0)} - 2 {u^{(1)}}_{(0,0)} + {u^{(2)}}_{(0,0)}}{dt^{2}} = \\frac{- 4 {u^{(1)}}_{(0,0)} + {u^{(1)}}_{(1,0)} + {u^{(1)}}_{(0,1)} + {u^{(1)}}_{(0,-1)} + {u^{(1)}}_{(-1,0)}}{dx^{2}}$" ], "text/plain": [ "Eq((u0_C - 2*u1_C + u2_C)/dt**2, (-4*u1_C + u1_E + u1_N + u1_S + u1_W)/dx**2)" ] }, "execution_count": 3, "metadata": {}, "output_type": "execute_result" } ], "source": [ "discretize = ps.fd.Discretization2ndOrder()\n", "\n", "def central2nd_time_derivative(fields):\n", " f_next, f_current, f_last = fields\n", " return (f_next[0, 0] - 2 * f_current[0, 0] + f_last[0, 0]) / discretize.dt**2\n", "\n", "rhs = ps.fd.diffusion(u_fields[1], 1)\n", "\n", "wave_eq = sp.Eq(central2nd_time_derivative(u_fields), discretize(rhs))\n", "\n", "wave_eq = sp.simplify(wave_eq)\n", "wave_eq" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "The explicit Euler scheme is now obtained by solving above equation with respect to $u_C^{next}$." ] }, { "cell_type": "code", "execution_count": 4, "metadata": {}, "outputs": [ { "data": { "text/latex": [ "$\\displaystyle {u^{(2)}}_{(0,0)} \\leftarrow_{} \\frac{- 4 {u^{(1)}}_{(0,0)} dt^{2} + {u^{(1)}}_{(1,0)} dt^{2} + {u^{(1)}}_{(0,1)} dt^{2} + {u^{(1)}}_{(0,-1)} dt^{2} + {u^{(1)}}_{(-1,0)} dt^{2} + dx^{2} \\left(- {u^{(0)}}_{(0,0)} + 2 {u^{(1)}}_{(0,0)}\\right)}{dx^{2}}$" ], "text/plain": [ "Assignment(u2_C, (-4*u1_C*dt**2 + u1_E*dt**2 + u1_N*dt**2 + u1_S*dt**2 + u1_W*dt**2 + dx**2*(-u0_C + 2*u1_C))/dx**2)" ] }, "execution_count": 4, "metadata": {}, "output_type": "execute_result" } ], "source": [ "u_next_C = u_fields[-1][0, 0]\n", "update_rule = ps.Assignment(u_next_C, sp.solve(wave_eq, u_next_C)[0])\n", "update_rule" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "Before creating the kernel, we substitute numeric values for $dx$ and $dt$. \n", "Then a kernel is created just like in the last tutorial." ] }, { "cell_type": "code", "execution_count": 5, "metadata": {}, "outputs": [ { "data": { "text/html": [ "
void kernel (const double * RESTRICT const _data_u0, const double * RESTRICT const _data_u1, double * RESTRICT const _data_u2)\n",
       "{\n",
       "   for(ctr_0: int64 = [1: int64]; ctr_0 < [59: int64]; ctr_0 += [1: int64])\n",
       "   {\n",
       "      for(ctr_1: int64 = [1: int64]; ctr_1 < [69: int64]; ctr_1 += [1: int64])\n",
       "      {\n",
       "         _data_u2[ctr_0 * [70: int64] + ctr_1] = [0.25: float64] * _data_u1[(ctr_0 + [1: int64]) * [70: int64] + ctr_1] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [1: int64])] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [-1: int64])] + [0.25: float64] * _data_u1[(ctr_0 + [-1: int64]) * [70: int64] + ctr_1] + _data_u1[ctr_0 * [70: int64] + ctr_1] + [-1.0: float64] * _data_u0[ctr_0 * [70: int64] + ctr_1];\n",
       "      }\n",
       "   }\n",
       "}\n",
       "
\n" ], "text/latex": [ "\\begin{Verbatim}[commandchars=\\\\\\{\\}]\n", "\\PY{k+kt}{void}\\PY{+w}{ }\\PY{n+nf}{kernel}\\PY{+w}{ }\\PY{p}{(}\\PY{k}{const}\\PY{+w}{ }\\PY{k+kt}{double}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{RESTRICT}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u0}\\PY{p}{,}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{k+kt}{double}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{RESTRICT}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{,}\\PY{+w}{ }\\PY{k+kt}{double}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{RESTRICT}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u2}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\\PY{+w}{ }\\PY{k}{for}\\PY{p}{(}\\PY{n}{ctr\\PYZus{}0}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{+w}{ }\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{;}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{\\PYZlt{}}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{59}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{;}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{+}\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\n", "\\PY{+w}{ }\\PY{p}{\\PYZob{}}\n", "\\PY{+w}{ }\\PY{k}{for}\\PY{p}{(}\\PY{n}{ctr\\PYZus{}1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{+w}{ }\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{;}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{+w}{ }\\PY{o}{\\PYZlt{}}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{69}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{;}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{+w}{ }\\PY{o}{+}\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\n", "\\PY{+w}{ }\\PY{p}{\\PYZob{}}\n", "\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u2}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{p}{(}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{ctr\\PYZus{}1}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{ctr\\PYZus{}1}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{\\PYZhy{}1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{p}{(}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{\\PYZhy{}1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{\\PYZhy{}1.0}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u0}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{p}{;}\n", "\\PY{+w}{ }\\PY{p}{\\PYZcb{}}\n", "\\PY{+w}{ }\\PY{p}{\\PYZcb{}}\n", "\\PY{p}{\\PYZcb{}}\n", "\\end{Verbatim}\n" ], "text/plain": [ "void kernel (const double * RESTRICT const _data_u0, const double * RESTRICT const _data_u1, double * RESTRICT const _data_u2)\n", "{\n", " for(ctr_0: int64 = [1: int64]; ctr_0 < [59: int64]; ctr_0 += [1: int64])\n", " {\n", " for(ctr_1: int64 = [1: int64]; ctr_1 < [69: int64]; ctr_1 += [1: int64])\n", " {\n", " _data_u2[ctr_0 * [70: int64] + ctr_1] = [0.25: float64] * _data_u1[(ctr_0 + [1: int64]) * [70: int64] + ctr_1] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [1: int64])] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [-1: int64])] + [0.25: float64] * _data_u1[(ctr_0 + [-1: int64]) * [70: int64] + ctr_1] + _data_u1[ctr_0 * [70: int64] + ctr_1] + [-1.0: float64] * _data_u0[ctr_0 * [70: int64] + ctr_1];\n", " }\n", " }\n", "}" ] }, "metadata": {}, "output_type": "display_data" } ], "source": [ "update_rule = update_rule.subs({discretize.dx: 0.1, discretize.dt: 0.05})\n", "ast = ps.create_kernel(update_rule)\n", "kernel = ast.compile()\n", "\n", "ps.inspect(ast)" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "To run simulation a suitable initial condition and boundary treatment is required. We chose an initial condition which is zero at the borders of the domain. The outermost layer is not changed by the update kernel, so we have an implicit homogenous Dirichlet boundary condition." ] }, { "cell_type": "code", "execution_count": 6, "metadata": {}, "outputs": [], "source": [ "X,Y = np.meshgrid( np.linspace(0, 1, size[1]), np.linspace(0,1, size[0]))\n", "Z = np.sin(2*X*np.pi) * np.sin(2*Y*np.pi)\n", "\n", "# Initialize the previous and current values with the initial function\n", "np.copyto(u_arrays[0], Z)\n", "np.copyto(u_arrays[1], Z)\n", "# The values for the next timesteps do not matter, since they are overwritten\n", "u_arrays[2][:, :] = 0" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "One timestep now consists of applying the kernel once, then shifting the arrays." ] }, { "cell_type": "code", "execution_count": 7, "metadata": {}, "outputs": [], "source": [ "def run(timesteps=1):\n", " for t in range(timesteps):\n", " kernel(u0=u_arrays[0], u1=u_arrays[1], u2=u_arrays[2])\n", " u_arrays[0], u_arrays[1], u_arrays[2] = u_arrays[1], u_arrays[2], u_arrays[0]\n", " return u_arrays[2]" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "Lets create an animation of the solution:" ] }, { "cell_type": "code", "execution_count": 8, "metadata": {}, "outputs": [ { "data": { "text/html": [ "" ], "text/plain": [ "" ] }, "execution_count": 8, "metadata": {}, "output_type": "execute_result" } ], "source": [ "ani = psplot.surface_plot_animation(run, zlim=(-1, 1))\n", "\n", "from pystencils.jupyter import display_as_html_video\n", "display_as_html_video(ani)" ] }, { "cell_type": "code", "execution_count": 9, "metadata": {}, "outputs": [], "source": [ "assert np.isfinite(np.max(u_arrays[2]))" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "__Runing on GPU__\n", "\n", "We can also run the same kernel on the GPU, by using the ``cupy`` package." ] }, { "cell_type": "code", "execution_count": 10, "metadata": {}, "outputs": [ { "data": { "text/html": [ "
__global__ void kernel (const double * RESTRICT const _data_u0, const double * RESTRICT const _data_u1, double * RESTRICT const _data_u2)\n",
       "{\n",
       "   ctr_0: const int64 = [1: int64] + ((int64) blockIdx.y * (int64) blockDim.y + (int64) threadIdx.y);\n",
       "   ctr_1: const int64 = [1: int64] + ((int64) blockIdx.x * (int64) blockDim.x + (int64) threadIdx.x);\n",
       "   if(ctr_0 < [59: int64] && ctr_1 < [69: int64])\n",
       "   {\n",
       "      _data_u2[ctr_0 * [70: int64] + ctr_1] = [0.25: float64] * _data_u1[(ctr_0 + [1: int64]) * [70: int64] + ctr_1] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [1: int64])] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [-1: int64])] + [0.25: float64] * _data_u1[(ctr_0 + [-1: int64]) * [70: int64] + ctr_1] + _data_u1[ctr_0 * [70: int64] + ctr_1] + [-1.0: float64] * _data_u0[ctr_0 * [70: int64] + ctr_1];\n",
       "   }\n",
       "}\n",
       "
\n" ], "text/latex": [ "\\begin{Verbatim}[commandchars=\\\\\\{\\}]\n", "\\PY{n}{\\PYZus{}\\PYZus{}global\\PYZus{}\\PYZus{}}\\PY{+w}{ }\\PY{k+kt}{void}\\PY{+w}{ }\\PY{n}{kernel}\\PY{+w}{ }\\PY{p}{(}\\PY{k}{const}\\PY{+w}{ }\\PY{k+kt}{double}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{RESTRICT}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u0}\\PY{p}{,}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{k+kt}{double}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{RESTRICT}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{,}\\PY{+w}{ }\\PY{k+kt}{double}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{RESTRICT}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u2}\\PY{p}{)}\n", "\\PY{p}{\\PYZob{}}\n", "\\PY{+w}{ }\\PY{n+nl}{ctr\\PYZus{}0}\\PY{p}{:}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{int64}\\PY{+w}{ }\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{p}{(}\\PY{n}{int64}\\PY{p}{)}\\PY{+w}{ }\\PY{n}{blockIdx}\\PY{p}{.}\\PY{n}{y}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{int64}\\PY{p}{)}\\PY{+w}{ }\\PY{n}{blockDim}\\PY{p}{.}\\PY{n}{y}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{int64}\\PY{p}{)}\\PY{+w}{ }\\PY{n}{threadIdx}\\PY{p}{.}\\PY{n}{y}\\PY{p}{)}\\PY{p}{;}\n", "\\PY{+w}{ }\\PY{n+nl}{ctr\\PYZus{}1}\\PY{p}{:}\\PY{+w}{ }\\PY{k}{const}\\PY{+w}{ }\\PY{n}{int64}\\PY{+w}{ }\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{p}{(}\\PY{n}{int64}\\PY{p}{)}\\PY{+w}{ }\\PY{n}{blockIdx}\\PY{p}{.}\\PY{n}{x}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{int64}\\PY{p}{)}\\PY{+w}{ }\\PY{n}{blockDim}\\PY{p}{.}\\PY{n}{x}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{int64}\\PY{p}{)}\\PY{+w}{ }\\PY{n}{threadIdx}\\PY{p}{.}\\PY{n}{x}\\PY{p}{)}\\PY{p}{;}\n", "\\PY{+w}{ }\\PY{k}{if}\\PY{p}{(}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{\\PYZlt{}}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{59}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{\\PYZam{}}\\PY{o}{\\PYZam{}}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{+w}{ }\\PY{o}{\\PYZlt{}}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{69}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\n", "\\PY{+w}{ }\\PY{p}{\\PYZob{}}\n", "\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u2}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{=}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{p}{(}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{ctr\\PYZus{}1}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{(}\\PY{n}{ctr\\PYZus{}1}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{\\PYZhy{}1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{0.25}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{p}{(}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{\\PYZhy{}1}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{p}{)}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u1}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mf}{\\PYZhy{}1.0}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{float64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{n}{\\PYZus{}data\\PYZus{}u0}\\PY{p}{[}\\PY{n}{ctr\\PYZus{}0}\\PY{+w}{ }\\PY{o}{*}\\PY{+w}{ }\\PY{p}{[}\\PY{l+m+mi}{70}\\PY{o}{:}\\PY{+w}{ }\\PY{n}{int64}\\PY{p}{]}\\PY{+w}{ }\\PY{o}{+}\\PY{+w}{ }\\PY{n}{ctr\\PYZus{}1}\\PY{p}{]}\\PY{p}{;}\n", "\\PY{+w}{ }\\PY{p}{\\PYZcb{}}\n", "\\PY{p}{\\PYZcb{}}\n", "\\end{Verbatim}\n" ], "text/plain": [ "__global__ void kernel (const double * RESTRICT const _data_u0, const double * RESTRICT const _data_u1, double * RESTRICT const _data_u2)\n", "{\n", " ctr_0: const int64 = [1: int64] + ((int64) blockIdx.y * (int64) blockDim.y + (int64) threadIdx.y);\n", " ctr_1: const int64 = [1: int64] + ((int64) blockIdx.x * (int64) blockDim.x + (int64) threadIdx.x);\n", " if(ctr_0 < [59: int64] && ctr_1 < [69: int64])\n", " {\n", " _data_u2[ctr_0 * [70: int64] + ctr_1] = [0.25: float64] * _data_u1[(ctr_0 + [1: int64]) * [70: int64] + ctr_1] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [1: int64])] + [0.25: float64] * _data_u1[ctr_0 * [70: int64] + (ctr_1 + [-1: int64])] + [0.25: float64] * _data_u1[(ctr_0 + [-1: int64]) * [70: int64] + ctr_1] + _data_u1[ctr_0 * [70: int64] + ctr_1] + [-1.0: float64] * _data_u0[ctr_0 * [70: int64] + ctr_1];\n", " }\n", "}" ] }, "metadata": {}, "output_type": "display_data" } ], "source": [ "try:\n", " import cupy\n", "except ImportError:\n", " cupy=None\n", " print('No cupy installed')\n", "\n", "\n", "res = None\n", "if cupy:\n", " gpu_ast = ps.create_kernel(update_rule, target=ps.Target.GPU)\n", " gpu_kernel = gpu_ast.compile()\n", " res = ps.inspect(gpu_ast)\n", "res" ] }, { "cell_type": "markdown", "metadata": {}, "source": [ "The run function has to be changed now slightly, since the data has to be transfered to the GPU first, then the kernel can be executed, and in the end the data has to be transfered back" ] }, { "cell_type": "code", "execution_count": 11, "metadata": {}, "outputs": [], "source": [ "if cupy:\n", " def run_on_gpu(timesteps=1):\n", " # Transfer arrays to GPU\n", " gpuArrs = [cupy.asarray(cpu_array) for cpu_array in u_arrays]\n", "\n", " for t in range(timesteps):\n", " gpu_kernel(u0=gpuArrs[0], u1=gpuArrs[1], u2=gpuArrs[2])\n", " gpuArrs[0], gpuArrs[1], gpuArrs[2] = gpuArrs[1], gpuArrs[2], gpuArrs[0]\n", "\n", " # Transfer arrays to CPU\n", " for gpuArr, cpuArr in zip(gpuArrs, u_arrays):\n", " cpuArr[:] = gpuArr.get()\n", "assert np.isfinite(np.max(u_arrays[2])) " ] }, { "cell_type": "code", "execution_count": 12, "metadata": {}, "outputs": [ { "ename": "CompileException", "evalue": "In file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:22:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/././bits/philox_rand.h:113:9: error: no type named 'uint32_t' in namespace 'std'; did you mean '__hip_internal::uint32_t'?\n 113 | typedef std::uint32_t uint32;\n | ^~~~~~~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:22:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/././bits/philox_rand.h:114:9: error: no type named 'uint64_t' in namespace 'std'; did you mean '__hip_internal::uint64_t'?\n 114 | typedef std::uint64_t uint64;\n | ^~~~~~~~~~~~~\n | __hip_internal::uint64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1428:28: note: '__hip_internal::uint64_t' declared here\n 1428 | typedef unsigned long long uint64_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:46: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:62: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:33:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 33 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:33:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 33 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:46: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:62: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:45:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 45 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:45:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 45 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:10: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:34: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:57: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:80: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:9:10: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 9 | const int64_t ctr_1 = 1LL + ((int64_t) blockIdx.x * (int64_t) blockDim.x + (int64_t) threadIdx.x);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\nfatal error: too many errors emitted, stopping now [-ferror-limit=]\n20 errors generated when compiling for gfx1103.", "output_type": "error", "traceback": [ "\u001b[31m---------------------------------------------------------------------------\u001b[39m", "\u001b[31mNVRTCError\u001b[39m Traceback (most recent call last)", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/.venv/lib/python3.12/site-packages/cupy/cuda/compiler.py:860\u001b[39m, in \u001b[36m_NVRTCProgram.compile\u001b[39m\u001b[34m(self, options, log_stream)\u001b[39m\n\u001b[32m 859\u001b[39m nvrtc.addNameExpression(\u001b[38;5;28mself\u001b[39m.ptr, ker)\n\u001b[32m--> \u001b[39m\u001b[32m860\u001b[39m \u001b[30;43mnvrtc\u001b[39;49m\u001b[30;43m.\u001b[39;49m\u001b[30;43mcompileProgram\u001b[39;49m\u001b[30;43m(\u001b[39;49m\u001b[30;43mself\u001b[39;49m\u001b[30;43m.\u001b[39;49m\u001b[30;43mptr\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43moptions\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n\u001b[32m 861\u001b[39m mapping = \u001b[38;5;28;01mNone\u001b[39;00m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy_backends/cuda/libs/nvrtc.pyx:125\u001b[39m, in \u001b[36mcupy_backends.cuda.libs.nvrtc.compileProgram\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m--> \u001b[39m\u001b[32m125\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy_backends/cuda/libs/nvrtc.pyx:138\u001b[39m, in \u001b[36mcupy_backends.cuda.libs.nvrtc.compileProgram\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m--> \u001b[39m\u001b[32m138\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy_backends/cuda/libs/nvrtc.pyx:53\u001b[39m, in \u001b[36mcupy_backends.cuda.libs.nvrtc.check_status\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m---> \u001b[39m\u001b[32m53\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[31mNVRTCError\u001b[39m: HIPRTC_ERROR_COMPILATION (6)", "\nDuring handling of the above exception, another exception occurred:\n", "\u001b[31mCompileException\u001b[39m Traceback (most recent call last)", "\u001b[36mCell\u001b[39m\u001b[36m \u001b[39m\u001b[32mIn[12]\u001b[39m\u001b[32m, line 2\u001b[39m\n\u001b[32m 1\u001b[39m \u001b[38;5;28;01mif\u001b[39;00m cupy:\n\u001b[32m----> \u001b[39m\u001b[32m2\u001b[39m run_on_gpu(\u001b[32m400\u001b[39m)\n", "\u001b[36mCell\u001b[39m\u001b[36m \u001b[39m\u001b[32mIn[11]\u001b[39m\u001b[32m, line 7\u001b[39m, in \u001b[36mrun_on_gpu\u001b[39m\u001b[34m(timesteps)\u001b[39m\n\u001b[32m 3\u001b[39m \u001b[38;5;66;03m# Transfer arrays to GPU\u001b[39;00m\n\u001b[32m 4\u001b[39m gpuArrs = [cupy.asarray(cpu_array) \u001b[38;5;28;01mfor\u001b[39;00m cpu_array \u001b[38;5;28;01min\u001b[39;00m u_arrays]\n\u001b[32m 5\u001b[39m \n\u001b[32m 6\u001b[39m \u001b[38;5;28;01mfor\u001b[39;00m t \u001b[38;5;28;01min\u001b[39;00m range(timesteps):\n\u001b[32m----> \u001b[39m\u001b[32m7\u001b[39m gpu_kernel(u0=gpuArrs[\u001b[32m0\u001b[39m], u1=gpuArrs[\u001b[32m1\u001b[39m], u2=gpuArrs[\u001b[32m2\u001b[39m])\n\u001b[32m 8\u001b[39m gpuArrs[\u001b[32m0\u001b[39m], gpuArrs[\u001b[32m1\u001b[39m], gpuArrs[\u001b[32m2\u001b[39m] = gpuArrs[\u001b[32m1\u001b[39m], gpuArrs[\u001b[32m2\u001b[39m], gpuArrs[\u001b[32m0\u001b[39m]\n\u001b[32m 9\u001b[39m \n\u001b[32m 10\u001b[39m \u001b[38;5;66;03m# Transfer arrays to CPU\u001b[39;00m\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/pystencils/src/pystencils/jit/jit.py:29\u001b[39m, in \u001b[36mKernelWrapper.__call__\u001b[39m\u001b[34m(self, **kwargs)\u001b[39m\n\u001b[32m 27\u001b[39m \u001b[38;5;28;01mfor\u001b[39;00m name, mapper \u001b[38;5;129;01min\u001b[39;00m \u001b[38;5;28mself\u001b[39m._view_mappings.items():\n\u001b[32m 28\u001b[39m kwargs[name] = mapper.view_ndarray(kwargs[name])\n\u001b[32m---> \u001b[39m\u001b[32m29\u001b[39m \u001b[30;43mself\u001b[39;49m\u001b[30;43m.\u001b[39;49m\u001b[30;43m_invocable\u001b[39;49m\u001b[30;43m(\u001b[39;49m\u001b[30;43m*\u001b[39;49m\u001b[30;43m*\u001b[39;49m\u001b[30;43mkwargs\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/pystencils/src/pystencils/jit/gpu_cupy.py:62\u001b[39m, in \u001b[36mCupyKernelWrapper._invoke\u001b[39m\u001b[34m(self, **kwargs)\u001b[39m\n\u001b[32m 60\u001b[39m device = \u001b[38;5;28mself\u001b[39m._get_device(kernel_args)\n\u001b[32m 61\u001b[39m \u001b[38;5;28;01mwith\u001b[39;00m cp.cuda.Device(device):\n\u001b[32m---> \u001b[39m\u001b[32m62\u001b[39m \u001b[30;43mself\u001b[39;49m\u001b[30;43m.\u001b[39;49m\u001b[30;43m_raw_kernel\u001b[39;49m\u001b[30;43m(\u001b[39;49m\u001b[30;43mlaunch_grid\u001b[39;49m\u001b[30;43m.\u001b[39;49m\u001b[30;43mgrid\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mlaunch_grid\u001b[39;49m\u001b[30;43m.\u001b[39;49m\u001b[30;43mblock\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mkernel_args\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy/_core/raw.pyx:101\u001b[39m, in \u001b[36mcupy._core.raw.RawKernel.__call__\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m--> \u001b[39m\u001b[32m101\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy/_core/raw.pyx:108\u001b[39m, in \u001b[36mcupy._core.raw.RawKernel.kernel.__get__\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m--> \u001b[39m\u001b[32m108\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy/_core/raw.pyx:125\u001b[39m, in \u001b[36mcupy._core.raw.RawKernel._kernel\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m--> \u001b[39m\u001b[32m125\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy/_util.pyx:67\u001b[39m, in \u001b[36mcupy._util.memoize.decorator.ret\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m---> \u001b[39m\u001b[32m67\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy/_core/raw.pyx:549\u001b[39m, in \u001b[36mcupy._core.raw._get_raw_module\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m--> \u001b[39m\u001b[32m549\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy/_core/core.pyx:2621\u001b[39m, in \u001b[36mcupy._core.core.compile_with_cache\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m-> \u001b[39m\u001b[32m2621\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32mcupy/_core/core.pyx:2639\u001b[39m, in \u001b[36mcupy._core.core.compile_with_cache\u001b[39m\u001b[34m()\u001b[39m\n\u001b[32m-> \u001b[39m\u001b[32m2639\u001b[39m \u001b[33m'Could not get source, probably due dynamically evaluated source code.'\u001b[39m\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/.venv/lib/python3.12/site-packages/cupy/cuda/compiler.py:645\u001b[39m, in \u001b[36m_compile_module_with_cache\u001b[39m\u001b[34m(source, options, arch, extra_source, backend, enable_cooperative_groups, name_expressions, log_stream, jitify, to_ltoir)\u001b[39m\n\u001b[32m 643\u001b[39m \u001b[38;5;28;01mif\u001b[39;00m runtime.is_hip:\n\u001b[32m 644\u001b[39m backend = \u001b[33m'\u001b[39m\u001b[33mhiprtc\u001b[39m\u001b[33m'\u001b[39m \u001b[38;5;28;01mif\u001b[39;00m backend == \u001b[33m'\u001b[39m\u001b[33mnvrtc\u001b[39m\u001b[33m'\u001b[39m \u001b[38;5;28;01melse\u001b[39;00m \u001b[33m'\u001b[39m\u001b[33mhipcc\u001b[39m\u001b[33m'\u001b[39m\n\u001b[32m--> \u001b[39m\u001b[32m645\u001b[39m \u001b[38;5;28;01mreturn\u001b[39;00m \u001b[30;43m_compile_with_cache_hip\u001b[39;49m\u001b[30;43m(\u001b[39;49m\n\u001b[32m 646\u001b[39m \u001b[30;43m \u001b[39;49m\u001b[30;43msource\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43moptions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43march\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mextra_source\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mbackend\u001b[39;49m\u001b[30;43m,\u001b[39;49m\n\u001b[32m 647\u001b[39m \u001b[30;43m \u001b[39;49m\u001b[30;43mname_expressions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mlog_stream\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mcache_in_memory\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n\u001b[32m 648\u001b[39m \u001b[38;5;28;01melse\u001b[39;00m:\n\u001b[32m 649\u001b[39m \u001b[38;5;28;01mreturn\u001b[39;00m _compile_with_cache_cuda(\n\u001b[32m 650\u001b[39m source, options, arch, extra_source, backend,\n\u001b[32m 651\u001b[39m enable_cooperative_groups, name_expressions, log_stream,\n\u001b[32m 652\u001b[39m cache_in_memory, jitify, to_ltoir)\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/.venv/lib/python3.12/site-packages/cupy/cuda/compiler.py:1072\u001b[39m, in \u001b[36m_compile_with_cache_hip\u001b[39m\u001b[34m(source, options, arch, extra_source, backend, name_expressions, log_stream, cache_in_memory, use_converter)\u001b[39m\n\u001b[32m 1068\u001b[39m \u001b[38;5;28;01mpass\u001b[39;00m\n\u001b[32m 1070\u001b[39m \u001b[38;5;28;01mif\u001b[39;00m backend == \u001b[33m'\u001b[39m\u001b[33mhiprtc\u001b[39m\u001b[33m'\u001b[39m:\n\u001b[32m 1071\u001b[39m \u001b[38;5;66;03m# compile_using_nvrtc calls hiprtc for hip builds\u001b[39;00m\n\u001b[32m-> \u001b[39m\u001b[32m1072\u001b[39m binary, mapping = \u001b[30;43mcompile_using_nvrtc\u001b[39;49m\u001b[30;43m(\u001b[39;49m\n\u001b[32m 1073\u001b[39m \u001b[30;43m \u001b[39;49m\u001b[30;43msource\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43moptions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43march\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mname\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43m+\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43m'\u001b[39;49m\u001b[30;43m.cu\u001b[39;49m\u001b[30;43m'\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mname_expressions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\n\u001b[32m 1074\u001b[39m \u001b[30;43m \u001b[39;49m\u001b[30;43mlog_stream\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mcache_in_memory\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n\u001b[32m 1075\u001b[39m mod._set_mapping(mapping)\n\u001b[32m 1076\u001b[39m \u001b[38;5;28;01melse\u001b[39;00m:\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/.venv/lib/python3.12/site-packages/cupy/cuda/compiler.py:414\u001b[39m, in \u001b[36mcompile_using_nvrtc\u001b[39m\u001b[34m(source, options, arch, filename, name_expressions, log_stream, cache_in_memory, jitify, method)\u001b[39m\n\u001b[32m 411\u001b[39m \u001b[38;5;28;01mif\u001b[39;00m jitify \u001b[38;5;129;01mis\u001b[39;00m \u001b[38;5;129;01mnot\u001b[39;00m \u001b[38;5;28;01mNone\u001b[39;00m:\n\u001b[32m 412\u001b[39m _jitify_deprecation_warning(jitify)\n\u001b[32m--> \u001b[39m\u001b[32m414\u001b[39m \u001b[38;5;28;01mreturn\u001b[39;00m \u001b[30;43m_compile_using_nvrtc_no_warning\u001b[39;49m\u001b[30;43m(\u001b[39;49m\n\u001b[32m 415\u001b[39m \u001b[30;43m \u001b[39;49m\u001b[30;43msource\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43moptions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43march\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mfilename\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mname_expressions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mlog_stream\u001b[39;49m\u001b[30;43m,\u001b[39;49m\n\u001b[32m 416\u001b[39m \u001b[30;43m \u001b[39;49m\u001b[30;43mcache_in_memory\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mjitify\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mmethod\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/.venv/lib/python3.12/site-packages/cupy/cuda/compiler.py:402\u001b[39m, in \u001b[36m_compile_using_nvrtc_no_warning\u001b[39m\u001b[34m(source, options, arch, filename, name_expressions, log_stream, cache_in_memory, jitify, method)\u001b[39m\n\u001b[32m 399\u001b[39m \u001b[38;5;28;01melse\u001b[39;00m:\n\u001b[32m 400\u001b[39m cu_path = \u001b[33m'\u001b[39m\u001b[33m'\u001b[39m \u001b[38;5;28;01mif\u001b[39;00m \u001b[38;5;129;01mnot\u001b[39;00m jitify \u001b[38;5;28;01melse\u001b[39;00m filename\n\u001b[32m--> \u001b[39m\u001b[32m402\u001b[39m \u001b[38;5;28;01mreturn\u001b[39;00m \u001b[30;43m_compile\u001b[39;49m\u001b[30;43m(\u001b[39;49m\u001b[30;43msource\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43moptions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mcu_path\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mname_expressions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\n\u001b[32m 403\u001b[39m \u001b[30;43m \u001b[39;49m\u001b[30;43mlog_stream\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mjitify\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mmethod\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/.venv/lib/python3.12/site-packages/cupy/cuda/compiler.py:384\u001b[39m, in \u001b[36m_compile_using_nvrtc_no_warning.._compile\u001b[39m\u001b[34m(source, options, cu_path, name_expressions, log_stream, jitify, method)\u001b[39m\n\u001b[32m 381\u001b[39m prog = _NVRTCProgram(source, cu_path, headers, include_names,\n\u001b[32m 382\u001b[39m name_expressions=name_expressions, method=method)\n\u001b[32m 383\u001b[39m \u001b[38;5;28;01mtry\u001b[39;00m:\n\u001b[32m--> \u001b[39m\u001b[32m384\u001b[39m compiled_obj, mapping = \u001b[30;43mprog\u001b[39;49m\u001b[30;43m.\u001b[39;49m\u001b[30;43mcompile\u001b[39;49m\u001b[30;43m(\u001b[39;49m\u001b[30;43moptions\u001b[39;49m\u001b[30;43m,\u001b[39;49m\u001b[30;43m \u001b[39;49m\u001b[30;43mlog_stream\u001b[39;49m\u001b[30;43m)\u001b[39;49m\n\u001b[32m 385\u001b[39m \u001b[38;5;28;01mexcept\u001b[39;00m CompileException \u001b[38;5;28;01mas\u001b[39;00m e:\n\u001b[32m 386\u001b[39m dump = _get_bool_env_variable(\n\u001b[32m 387\u001b[39m \u001b[33m'\u001b[39m\u001b[33mCUPY_DUMP_CUDA_SOURCE_ON_ERROR\u001b[39m\u001b[33m'\u001b[39m, \u001b[38;5;28;01mFalse\u001b[39;00m)\n", "\u001b[36mFile \u001b[39m\u001b[32m/pycodegen-dev/.venv/lib/python3.12/site-packages/cupy/cuda/compiler.py:879\u001b[39m, in \u001b[36m_NVRTCProgram.compile\u001b[39m\u001b[34m(self, options, log_stream)\u001b[39m\n\u001b[32m 877\u001b[39m \u001b[38;5;28;01mexcept\u001b[39;00m nvrtc.NVRTCError:\n\u001b[32m 878\u001b[39m log = nvrtc.getProgramLog(\u001b[38;5;28mself\u001b[39m.ptr)\n\u001b[32m--> \u001b[39m\u001b[32m879\u001b[39m \u001b[38;5;28;01mraise\u001b[39;00m CompileException(log, \u001b[38;5;28mself\u001b[39m.src, \u001b[38;5;28mself\u001b[39m.name, options,\n\u001b[32m 880\u001b[39m \u001b[33m'\u001b[39m\u001b[33mnvrtc\u001b[39m\u001b[33m'\u001b[39m \u001b[38;5;28;01mif\u001b[39;00m \u001b[38;5;129;01mnot\u001b[39;00m runtime.is_hip \u001b[38;5;28;01melse\u001b[39;00m \u001b[33m'\u001b[39m\u001b[33mhiprtc\u001b[39m\u001b[33m'\u001b[39m)\n", "\u001b[31mCompileException\u001b[39m: In file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:22:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/././bits/philox_rand.h:113:9: error: no type named 'uint32_t' in namespace 'std'; did you mean '__hip_internal::uint32_t'?\n 113 | typedef std::uint32_t uint32;\n | ^~~~~~~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:22:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/././bits/philox_rand.h:114:9: error: no type named 'uint64_t' in namespace 'std'; did you mean '__hip_internal::uint64_t'?\n 114 | typedef std::uint64_t uint64;\n | ^~~~~~~~~~~~~\n | __hip_internal::uint64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1428:28: note: '__hip_internal::uint64_t' declared here\n 1428 | typedef unsigned long long uint64_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:46: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:32:62: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 32 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:33:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 33 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:33:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 33 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:46: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:44:62: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 44 | (uint32_t)ctr0, (uint32_t)ctr1, (uint32_t)ctr2, (uint32_t)ctr3, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:45:14: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 45 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\nIn file included from /tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:1:\nIn file included from /pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/generic_gpu.hpp:4:\n/pycodegen-dev/pystencils/src/pystencils/include/pystencils_runtime/./random.hpp:45:30: error: unknown type name 'uint32_t'; did you mean '__hip_internal::uint32_t'?\n 45 | (uint32_t)key0, (uint32_t)key1, //\n | ^~~~~~~~\n | __hip_internal::uint32_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1427:22: note: '__hip_internal::uint32_t' declared here\n 1427 | typedef unsigned int uint32_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:10: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:34: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:57: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:8:80: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 8 | const int64_t ctr_0 = 1LL + ((int64_t) blockIdx.y * (int64_t) blockDim.y + (int64_t) threadIdx.y);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\n/tmp/comgr-345b3a/input/tmp/tmpnz1gjn5f/c8573e542c59846fdedfebf89b822ea2cba126fc.hsaco.cu:9:10: error: unknown type name 'int64_t'; did you mean '__hip_internal::int64_t'?\n 9 | const int64_t ctr_1 = 1LL + ((int64_t) blockIdx.x * (int64_t) blockDim.x + (int64_t) threadIdx.x);\n | ^~~~~~~\n | __hip_internal::int64_t\n/tmp/comgr-345b3a/include/hiprtc_runtime.h:1432:26: note: '__hip_internal::int64_t' declared here\n 1432 | typedef signed long long int64_t;\n | ^\nfatal error: too many errors emitted, stopping now [-ferror-limit=]\n20 errors generated when compiling for gfx1103." ] } ], "source": [ "if cupy:\n", " run_on_gpu(400)" ] } ], "metadata": { "anaconda-cloud": {}, "kernelspec": { "display_name": "Python 3 (ipykernel)", "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.12.3" } }, "nbformat": 4, "nbformat_minor": 4 }