[pyfr] 19/88: Simplify kernel generation.
Ghislain Vaillant
ghisvail-guest at moszumanska.debian.org
Wed Nov 16 12:05:26 UTC 2016
This is an automated email from the git hooks/post-receive script.
ghisvail-guest pushed a commit to branch master
in repository pyfr.
commit 989131e8ff7fcf6b6d1f2155a0a54126f78fd859
Author: Freddie Witherden <freddie at witherden.org>
Date: Sat Apr 23 15:49:21 2016 -0700
Simplify kernel generation.
---
pyfr/backends/base/generator.py | 79 +++++++++++++++++++++++++++++++++-----
pyfr/backends/cuda/generator.py | 71 ++--------------------------------
pyfr/backends/mic/generator.py | 80 ++-------------------------------------
pyfr/backends/opencl/generator.py | 71 ++--------------------------------
pyfr/backends/openmp/generator.py | 79 ++------------------------------------
5 files changed, 85 insertions(+), 295 deletions(-)
diff --git a/pyfr/backends/base/generator.py b/pyfr/backends/base/generator.py
index 1dd8ef3..9a7e7ca 100644
--- a/pyfr/backends/base/generator.py
+++ b/pyfr/backends/base/generator.py
@@ -6,15 +6,6 @@ import re
import numpy as np
-def procbody(body, fpdtype):
- # At single precision suffix all floating point constants by 'f'
- if fpdtype == np.float32:
- body = re.sub(r'(?=\d*[.eE])(?=\.?\d)\d*\.?\d*(?:[eE][+-]?\d+)?',
- r'\g<0>f', body)
-
- return body
-
-
class Arg(object):
def __init__(self, name, spec, body):
self.name = name
@@ -63,7 +54,6 @@ class BaseKernelGenerator(object, metaclass=ABCMeta):
def __init__(self, name, ndim, args, body, fpdtype):
self.name = name
self.ndim = ndim
- self.body = procbody(body, fpdtype)
self.fpdtype = fpdtype
# Parse and sort our argument list
@@ -90,6 +80,12 @@ class BaseKernelGenerator(object, metaclass=ABCMeta):
if ndim == 2 and any(v.ismpi for v in self.vectargs):
raise ValueError('MPI matrices are not supported for 2D kernels')
+ # Render the main body of our kernel
+ self.body = self._render_body(body)
+
+ # Determine the dimensions to be iterated over
+ self._dims = ['_nx'] if ndim == 1 else ['_ny', '_nx']
+
def argspec(self):
# Argument names and types
argn, argt = [], []
@@ -129,3 +125,66 @@ class BaseKernelGenerator(object, metaclass=ABCMeta):
@abstractmethod
def render(self):
pass
+
+ def _deref_arg_view(self, arg):
+ ptns = ['{0}_v[{0}_vix[_x]]',
+ r'{0}_v[{0}_vix[_x] + {0}_vcstri[_x]*\1]',
+ r'{0}_v[{0}_vix[_x] + {0}_vrstri[_x]*\1 + {0}_vcstri[_x]*\2]']
+
+ return ptns[arg.ncdim].format(arg.name)
+
+ def _deref_arg_array_1d(self, arg):
+ # Leading (sub) dimension
+ lsdim = 'lsd' + arg.name if not arg.ismpi else '_nx'
+
+ # Vector: name_v[_x]
+ if arg.ncdim == 0:
+ ix = '_x'
+ # Stacked vector: name_v[lsdim*\1 + _x]
+ elif arg.ncdim == 1:
+ ix = r'{0}*\1 + _x'.format(lsdim)
+ # Doubly stacked vector: name_v[(nv*\1 + \2)*lsdim + _x]
+ else:
+ ix = r'({0}*\1 + \2)*{1} + _x'.format(arg.cdims[1], lsdim)
+
+ return '{0}_v[{1}]'.format(arg.name, ix)
+
+ def _deref_arg_array_2d(self, arg):
+ # Broadcast vector: name_v[_x]
+ if arg.isbroadcast:
+ ix = '_x'
+ # Matrix: name_v[lsdim*_y + _x]
+ elif arg.ncdim == 0:
+ ix = 'lsd{}*_y + _x'.format(arg.name)
+ # Stacked matrix: name_v[(_y*nv + \1)*lsdim + _x]
+ elif arg.ncdim == 1:
+ ix = r'(_y*{0} + \1)*lsd{1} + _x'.format(arg.cdims[0], arg.name)
+ # Doubly stacked matrix: name_v[((\1*_ny + _y)*nv + \2)*lsdim + _x]
+ else:
+ ix = (r'((\1*_ny + _y)*{0} + \2)*lsd{1} + _x'
+ .format(arg.cdims[1], arg.name))
+
+ return '{0}_v[{1}]'.format(arg.name, ix)
+
+ def _render_body(self, body):
+ ptns = [r'\b{0}\b', r'\b{0}\[(\d+)\]', r'\b{0}\[(\d+)\]\[(\d+)\]']
+
+ # At single precision suffix all floating point constants by 'f'
+ if self.fpdtype == np.float32:
+ body = re.sub(r'(?=\d*[.eE])(?=\.?\d)\d*\.?\d*(?:[eE][+-]?\d+)?',
+ r'\g<0>f', body)
+
+ # Dereference vector arguments
+ for va in self.vectargs:
+ if va.isview:
+ darg = self._deref_arg_view(va)
+ else:
+ if self.ndim == 1:
+ darg = self._deref_arg_array_1d(va)
+ else:
+ darg = self._deref_arg_array_2d(va)
+
+ # Substitute
+ body = re.sub(ptns[va.ncdim].format(va.name), darg, body)
+
+ return body
diff --git a/pyfr/backends/cuda/generator.py b/pyfr/backends/cuda/generator.py
index 83c229d..46ae875 100644
--- a/pyfr/backends/cuda/generator.py
+++ b/pyfr/backends/cuda/generator.py
@@ -1,7 +1,5 @@
# -*- coding: utf-8 -*-
-import re
-
from pyfr.backends.base.generator import BaseKernelGenerator
@@ -11,18 +9,13 @@ class CUDAKernelGenerator(BaseKernelGenerator):
# Specialise
if self.ndim == 1:
- self._dims = ['_nx']
self._limits = 'if (_x < _nx)'
- self._deref_arg_array = self._deref_arg_array_1d
else:
- self._dims = ['_ny', '_nx']
self._limits = 'for (int _y = 0; _y < _ny && _x < _nx; ++_y)'
- self._deref_arg_array = self._deref_arg_array_2d
def render(self):
- # Get the kernel specification and main body
- spec = self._emit_spec()
- body = self._emit_body()
+ # Kernel spec
+ spec = self._render_spec()
# Iteration limits (if statement/for loop)
limits = self._limits
@@ -35,9 +28,9 @@ class CUDAKernelGenerator(BaseKernelGenerator):
{{
{body}
}}
- }}'''.format(spec=spec, limits=limits, body=body)
+ }}'''.format(spec=spec, limits=limits, body=self.body)
- def _emit_spec(self):
+ def _render_spec(self):
# We first need the argument list; starting with the dimensions
kargs = ['int ' + d for d in self._dims]
@@ -70,59 +63,3 @@ class CUDAKernelGenerator(BaseKernelGenerator):
kargs.append('int lsd{0.name}'.format(va))
return '__global__ void {0}({1})'.format(self.name, ', '.join(kargs))
-
- def _deref_arg_view(self, arg):
- ptns = ['{0}_v[{0}_vix[_x]]',
- r'{0}_v[{0}_vix[_x] + {0}_vcstri[_x]*\1]',
- r'{0}_v[{0}_vix[_x] + {0}_vrstri[_x]*\1 + {0}_vcstri[_x]*\2]']
-
- return ptns[arg.ncdim].format(arg.name)
-
- def _deref_arg_array_1d(self, arg):
- # Leading (sub) dimension
- lsdim = 'lsd' + arg.name if not arg.ismpi else '_nx'
-
- # Vector name_v[_x]
- if arg.ncdim == 0:
- ix = '_x'
- # Stacked vector; name_v[lsdim*\1 + _x]
- elif arg.ncdim == 1:
- ix = r'{0}*\1 + _x'.format(lsdim)
- # Doubly stacked vector; name_v[(nv*\1 + \2)*lsdim + _x]
- else:
- ix = r'({0}*\1 + \2)*{1} + _x'.format(arg.cdims[1], lsdim)
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _deref_arg_array_2d(self, arg):
- # Broadcast vector: name_v[_x]
- if arg.isbroadcast:
- ix = '_x'
- # Matrix: name_v[lsdim*_y + _x]
- elif arg.ncdim == 0:
- ix = 'lsd{0}*_y + _x'.format(arg.name)
- # Stacked matrix: name_v[(_y*nv + \1)*lsdim + _x]
- elif arg.ncdim == 1:
- ix = r'(_y*{0} + \1)*lsd{1} + _x'.format(arg.cdims[0], arg.name)
- # Doubly stacked matrix: name_v[((\1*_ny + _y)*nv + \2)*lsdim + _x]
- else:
- ix = (r'((\1*_ny + _y)*{0} + \2)*lsd{1} + _x'
- .format(arg.cdims[1], arg.name))
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _emit_body(self):
- body = self.body
- ptns = [r'\b{0}\b', r'\b{0}\[(\d+)\]', r'\b{0}\[(\d+)\]\[(\d+)\]']
-
- for va in self.vectargs:
- # Dereference the argument
- if va.isview:
- darg = self._deref_arg_view(va)
- else:
- darg = self._deref_arg_array(va)
-
- # Substitute
- body = re.sub(ptns[va.ncdim].format(va.name), darg, body)
-
- return body
diff --git a/pyfr/backends/mic/generator.py b/pyfr/backends/mic/generator.py
index f551ba4..ade6573 100644
--- a/pyfr/backends/mic/generator.py
+++ b/pyfr/backends/mic/generator.py
@@ -1,27 +1,12 @@
# -*- coding: utf-8 -*-
-import re
-
from pyfr.backends.base.generator import BaseKernelGenerator
-from pyfr.util import ndrange
class MICKernelGenerator(BaseKernelGenerator):
- def __init__(self, *args, **kwargs):
- super().__init__(*args, **kwargs)
-
- # Specialise
- if self.ndim == 1:
- self._dims = ['_nx']
- self._deref_arg_array = self._deref_arg_array_1d
- else:
- self._dims = ['_ny', '_nx']
- self._deref_arg_array = self._deref_arg_array_2d
-
def render(self):
- # Kernel spec, unpacking code, and body
- spec, unpack = self._emit_spec_unpack()
- body = self._emit_body()
+ # Kernel spec and unpacking code
+ spec, unpack = self._render_spec_unpack()
if self.ndim == 1:
tpl = '''{spec}
@@ -58,9 +43,9 @@ class MICKernelGenerator(BaseKernelGenerator):
}}
}}'''
- return tpl.format(spec=spec, unpack=unpack, body=body)
+ return tpl.format(spec=spec, unpack=unpack, body=self.body)
- def _emit_spec_unpack(self):
+ def _render_spec_unpack(self):
# Start by unpacking the dimensions
kspec = ['long *arg{0}' for d in self._dims]
kpack = ['int {0} = *arg{{0}};'.format(d) for d in self._dims]
@@ -108,60 +93,3 @@ class MICKernelGenerator(BaseKernelGenerator):
unpack = '\n'.join(a.format(i) for i, a in enumerate(kpack))
return 'void {0}({1})'.format(self.name, params), unpack
-
-
- def _deref_arg_view(self, arg):
- ptns = ['{0}_v[{0}_vix[_x]]',
- r'{0}_v[{0}_vix[_x] + {0}_vcstri[_x]*\1]',
- r'{0}_v[{0}_vix[_x] + {0}_vrstri[_x]*\1 + {0}_vcstri[_x]*\2]']
-
- return ptns[arg.ncdim].format(arg.name)
-
- def _deref_arg_array_1d(self, arg):
- # Leading (sub) dimension
- lsdim = 'lsd' + arg.name if not arg.ismpi else '_nx'
-
- # Vector: name_v[_x]
- if arg.ncdim == 0:
- ix = '_x'
- # Stacked vector: name_v[lsdim*\1 + _x]
- elif arg.ncdim == 1:
- ix = r'{0}*\1 + _x'.format(lsdim)
- # Doubly stacked vector: name_v[(nv*\1 + \2)*lsdim + _x]
- else:
- ix = r'({0}*\1 + \2)*{1} + _x'.format(arg.cdims[1], lsdim)
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _deref_arg_array_2d(self, arg):
- # Broadcast vector: name_v[_x]
- if arg.isbroadcast:
- ix = '_x'
- # Matrix: name_v[lsdim*_y + _x]
- elif arg.ncdim == 0:
- ix = 'lsd{}*_y + _x'.format(arg.name)
- # Stacked matrix: name_v[(_y*nv + \1)*lsdim + _x]
- elif arg.ncdim == 1:
- ix = r'(_y*{0} + \1)*lsd{1} + _x'.format(arg.cdims[0], arg.name)
- # Doubly stacked matrix: name_v[((\1*_ny + _y)*nv + \2)*lsdim + _x]
- else:
- ix = (r'((\1*_ny + _y)*{0} + \2)*lsd{1} + _x'
- .format(arg.cdims[1], arg.name))
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _emit_body(self):
- body = self.body
- ptns = [r'\b{0}\b', r'\b{0}\[(\d+)\]', r'\b{0}\[(\d+)\]\[(\d+)\]']
-
- for va in self.vectargs:
- # Dereference the argument
- if va.isview:
- darg = self._deref_arg_view(va)
- else:
- darg = self._deref_arg_array(va)
-
- # Substitute
- body = re.sub(ptns[va.ncdim].format(va.name), darg, body)
-
- return body
diff --git a/pyfr/backends/opencl/generator.py b/pyfr/backends/opencl/generator.py
index 398db7b..94320ba 100644
--- a/pyfr/backends/opencl/generator.py
+++ b/pyfr/backends/opencl/generator.py
@@ -1,7 +1,5 @@
# -*- coding: utf-8 -*-
-import re
-
from pyfr.backends.base.generator import BaseKernelGenerator
@@ -11,18 +9,13 @@ class OpenCLKernelGenerator(BaseKernelGenerator):
# Specialise
if self.ndim == 1:
- self._dims = ['_nx']
self._limits = 'if (_x < _nx)'
- self._deref_arg_array = self._deref_arg_array_1d
else:
- self._dims = ['_ny', '_nx']
self._limits = 'for (int _y = 0; _y < _ny && _x < _nx; ++_y)'
- self._deref_arg_array = self._deref_arg_array_2d
def render(self):
- # Get the kernel specification and main body
- spec = self._emit_spec()
- body = self._emit_body()
+ # Kernel spec
+ spec = self._render_spec()
# Iteration limits (if statement/for loop)
limits = self._limits
@@ -35,9 +28,9 @@ class OpenCLKernelGenerator(BaseKernelGenerator):
{{
{body}
}}
- }}'''.format(spec=spec, limits=limits, body=body)
+ }}'''.format(spec=spec, limits=limits, body=self.body)
- def _emit_spec(self):
+ def _render_spec(self):
# We first need the argument list; starting with the dimensions
kargs = ['int ' + d for d in self._dims]
@@ -71,59 +64,3 @@ class OpenCLKernelGenerator(BaseKernelGenerator):
kargs.extend(k.format(va) for k in ka)
return '__kernel void {0}({1})'.format(self.name, ', '.join(kargs))
-
- def _deref_arg_view(self, arg):
- ptns = ['{0}_v[{0}_vix[_x]]',
- r'{0}_v[{0}_vix[_x] + {0}_vcstri[_x]*\1]',
- r'{0}_v[{0}_vix[_x] + {0}_vrstri[_x]*\1 + {0}_vcstri[_x]*\2]']
-
- return ptns[arg.ncdim].format(arg.name)
-
- def _deref_arg_array_1d(self, arg):
- # Leading (sub) dimension
- lsdim = 'lsd' + arg.name if not arg.ismpi else '_nx'
-
- # Vector name_v[_x]
- if arg.ncdim == 0:
- ix = '_x'
- # Stacked vector; name_v[lsdim*\1 + _x]
- elif arg.ncdim == 1:
- ix = r'{0}*\1 + _x'.format(lsdim)
- # Doubly stacked vector; name_v[(nv*\1 + \2)*lsdim + _x]
- else:
- ix = r'({0}*\1 + \2)*{1} + _x'.format(arg.cdims[1], lsdim)
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _deref_arg_array_2d(self, arg):
- # Broadcast vector: name_v[_x]
- if arg.isbroadcast:
- ix = '_x'
- # Matrix: name_v[lsdim*_y + _x]
- elif arg.ncdim == 0:
- ix = 'lsd{0}*_y + _x'.format(arg.name)
- # Stacked matrix: name_v[(_y*nv + \1)*lsdim + _x]
- elif arg.ncdim == 1:
- ix = r'(_y*{0} + \1)*lsd{1} + _x'.format(arg.cdims[0], arg.name)
- # Doubly stacked matrix: name_v[((\1*_ny + _y)*nv + \2)*lsdim + _x]
- else:
- ix = (r'((\1*_ny + _y)*{0} + \2)*lsd{1} + _x'
- .format(arg.cdims[1], arg.name))
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _emit_body(self):
- body = self.body
- ptns = [r'\b{0}\b', r'\b{0}\[(\d+)\]', r'\b{0}\[(\d+)\]\[(\d+)\]']
-
- for va in self.vectargs:
- # Dereference the argument
- if va.isview:
- darg = self._deref_arg_view(va)
- else:
- darg = self._deref_arg_array(va)
-
- # Substitute
- body = re.sub(ptns[va.ncdim].format(va.name), darg, body)
-
- return body
diff --git a/pyfr/backends/openmp/generator.py b/pyfr/backends/openmp/generator.py
index aeecdc9..4ed7b05 100644
--- a/pyfr/backends/openmp/generator.py
+++ b/pyfr/backends/openmp/generator.py
@@ -1,26 +1,12 @@
# -*- coding: utf-8 -*-
-import re
-
from pyfr.backends.base.generator import BaseKernelGenerator
class OpenMPKernelGenerator(BaseKernelGenerator):
- def __init__(self, *args, **kwargs):
- super().__init__(*args, **kwargs)
-
- # Specialise
- if self.ndim == 1:
- self._dims = ['_nx']
- self._deref_arg_array = self._deref_arg_array_1d
- else:
- self._dims = ['_ny', '_nx']
- self._deref_arg_array = self._deref_arg_array_2d
-
def render(self):
- # Kernel spec and body
- spec = self._emit_spec()
- body = self._emit_body()
+ # Kernel spec
+ spec = self._render_spec()
if self.ndim == 1:
tpl = '''
@@ -56,9 +42,9 @@ class OpenMPKernelGenerator(BaseKernelGenerator):
}}
}}'''
- return tpl.format(spec=spec, body=body)
+ return tpl.format(spec=spec, body=self.body)
- def _emit_spec(self):
+ def _render_spec(self):
# We first need the argument list; starting with the dimensions
kargs = ['int ' + d for d in self._dims]
@@ -91,60 +77,3 @@ class OpenMPKernelGenerator(BaseKernelGenerator):
kargs.append('int lsd{0.name}'.format(va))
return 'void {0}({1})'.format(self.name, ', '.join(kargs))
-
-
- def _deref_arg_view(self, arg):
- ptns = ['{0}_v[{0}_vix[_x]]',
- r'{0}_v[{0}_vix[_x] + {0}_vcstri[_x]*\1]',
- r'{0}_v[{0}_vix[_x] + {0}_vrstri[_x]*\1 + {0}_vcstri[_x]*\2]']
-
- return ptns[arg.ncdim].format(arg.name)
-
- def _deref_arg_array_1d(self, arg):
- # Leading (sub) dimension
- lsdim = 'lsd' + arg.name if not arg.ismpi else '_nx'
-
- # Vector: name_v[_x]
- if arg.ncdim == 0:
- ix = '_x'
- # Stacked vector: name_v[lsdim*\1 + _x]
- elif arg.ncdim == 1:
- ix = r'{0}*\1 + _x'.format(lsdim)
- # Doubly stacked vector: name_v[(nv*\1 + \2)*lsdim + _x]
- else:
- ix = r'({0}*\1 + \2)*{1} + _x'.format(arg.cdims[1], lsdim)
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _deref_arg_array_2d(self, arg):
- # Broadcast vector: name_v[_x]
- if arg.isbroadcast:
- ix = '_x'
- # Matrix: name_v[lsdim*_y + _x]
- elif arg.ncdim == 0:
- ix = 'lsd{}*_y + _x'.format(arg.name)
- # Stacked matrix: name_v[(_y*nv + \1)*lsdim + _x]
- elif arg.ncdim == 1:
- ix = r'(_y*{0} + \1)*lsd{1} + _x'.format(arg.cdims[0], arg.name)
- # Doubly stacked matrix: name_v[((\1*_ny + _y)*nv + \2)*lsdim + _x]
- else:
- ix = (r'((\1*_ny + _y)*{0} + \2)*lsd{1} + _x'
- .format(arg.cdims[1], arg.name))
-
- return '{0}_v[{1}]'.format(arg.name, ix)
-
- def _emit_body(self):
- body = self.body
- ptns = [r'\b{0}\b', r'\b{0}\[(\d+)\]', r'\b{0}\[(\d+)\]\[(\d+)\]']
-
- for va in self.vectargs:
- # Dereference the argument
- if va.isview:
- darg = self._deref_arg_view(va)
- else:
- darg = self._deref_arg_array(va)
-
- # Substitute
- body = re.sub(ptns[va.ncdim].format(va.name), darg, body)
-
- return body
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/pyfr.git
More information about the debian-science-commits
mailing list