2016-03-30 15:00:00 -07:00
|
|
|
"""
|
|
|
|
Class for constructing and holding the basic FDTD operations and fields
|
|
|
|
"""
|
|
|
|
|
|
|
|
from typing import List, Dict, Callable
|
2016-09-01 14:39:44 -07:00
|
|
|
from collections import OrderedDict
|
2016-03-30 15:00:00 -07:00
|
|
|
import numpy
|
2016-09-01 14:39:44 -07:00
|
|
|
import jinja2
|
2016-03-30 15:00:00 -07:00
|
|
|
import warnings
|
|
|
|
|
|
|
|
import pyopencl
|
|
|
|
import pyopencl.array
|
|
|
|
from pyopencl.elementwise import ElementwiseKernel
|
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
from fdfd_tools import vec
|
|
|
|
|
|
|
|
|
|
|
|
__author__ = 'Jan Petykiewicz'
|
|
|
|
|
|
|
|
|
|
|
|
# Create jinja2 env on module load
|
|
|
|
jinja_env = jinja2.Environment(loader=jinja2.PackageLoader(__name__, 'kernels'))
|
2016-03-30 15:00:00 -07:00
|
|
|
|
|
|
|
|
|
|
|
class Simulation(object):
|
|
|
|
"""
|
|
|
|
Constructs and holds the basic FDTD operations and related fields
|
2017-08-24 11:28:03 -07:00
|
|
|
|
2019-07-15 00:11:09 -07:00
|
|
|
After constructing this object, call the (update_E, update_H) members
|
2017-08-24 11:28:03 -07:00
|
|
|
to perform FDTD updates on the stored (E, H, S) fields:
|
|
|
|
|
2017-10-01 12:34:11 -07:00
|
|
|
pmls = [{'axis': a, 'polarity': p} for a in 'xyz' for p in 'np']
|
|
|
|
sim = Simulation(grid.grids, do_poynting=True, pmls=pmls)
|
2017-08-24 11:28:03 -07:00
|
|
|
with open('sources.c', 'w') as f:
|
|
|
|
f.write('{}'.format(sim.sources))
|
|
|
|
|
|
|
|
for t in range(max_t):
|
|
|
|
sim.update_E([]).wait()
|
|
|
|
|
|
|
|
# Find the linear index for the center point, for Ey
|
|
|
|
ind = numpy.ravel_multi_index(tuple(grid.shape//2), dims=grid.shape, order='C') + \
|
|
|
|
numpy.prod(grid.shape) * 1
|
|
|
|
# Perturb the field (i.e., add a soft current source)
|
|
|
|
sim.E[ind] += numpy.sin(omega * t * sim.dt)
|
|
|
|
event = sim.update_H([])
|
|
|
|
event.wait()
|
|
|
|
|
|
|
|
with lzma.open('saved_simulation', 'wb') as f:
|
|
|
|
dill.dump(fdfd_tools.unvec(sim.E.get(), grid.shape), f)
|
|
|
|
|
|
|
|
Code in the form
|
|
|
|
event2 = sim.update_H([event0, event1])
|
|
|
|
indicates that the update_H operation should be prepared immediately, but wait for
|
|
|
|
event0 and event1 to occur (i.e. previous operations to finish) before starting execution.
|
|
|
|
event2 can then be used to prepare further operations to be run after update_H.
|
2016-03-30 15:00:00 -07:00
|
|
|
"""
|
2018-11-30 00:13:27 -08:00
|
|
|
E = None # type: pyopencl.array.Array
|
|
|
|
H = None # type: pyopencl.array.Array
|
|
|
|
S = None # type: pyopencl.array.Array
|
|
|
|
eps = None # type: pyopencl.array.Array
|
|
|
|
dt = None # type: float
|
2018-11-30 00:36:11 -08:00
|
|
|
inv_dxes = None # type: List[pyopencl.array.Array]
|
2016-03-30 15:00:00 -07:00
|
|
|
|
|
|
|
arg_type = None # type: numpy.float32 or numpy.float64
|
|
|
|
|
|
|
|
context = None # type: pyopencl.Context
|
|
|
|
queue = None # type: pyopencl.CommandQueue
|
|
|
|
|
2017-08-24 11:28:03 -07:00
|
|
|
update_E = None # type: Callable[[List[pyopencl.Event]], pyopencl.Event]
|
|
|
|
update_H = None # type: Callable[[List[pyopencl.Event]], pyopencl.Event]
|
|
|
|
update_S = None # type: Callable[[List[pyopencl.Event]], pyopencl.Event]
|
2018-11-29 02:01:28 -08:00
|
|
|
update_J = None # type: Callable[[List[pyopencl.Event]], pyopencl.Event]
|
2016-09-01 14:39:44 -07:00
|
|
|
sources = None # type: Dict[str, str]
|
2016-03-30 15:00:00 -07:00
|
|
|
|
|
|
|
def __init__(self,
|
|
|
|
epsilon: List[numpy.ndarray],
|
2017-10-01 12:34:11 -07:00
|
|
|
pmls: List[Dict[str, int or float]],
|
2018-11-29 02:01:48 -08:00
|
|
|
bloch_boundaries: List[Dict[str, int or float]] = (),
|
2018-11-30 00:36:11 -08:00
|
|
|
dxes: List[List[numpy.ndarray]] or float = None,
|
|
|
|
dt: float = None,
|
2018-11-29 01:59:05 -08:00
|
|
|
initial_fields: Dict[str, List[numpy.ndarray]] = None,
|
2016-09-01 14:39:44 -07:00
|
|
|
context: pyopencl.Context = None,
|
|
|
|
queue: pyopencl.CommandQueue = None,
|
|
|
|
float_type: numpy.float32 or numpy.float64 = numpy.float32,
|
2018-11-29 02:01:28 -08:00
|
|
|
do_poynting: bool = True,
|
|
|
|
do_fieldsrc: bool = False):
|
2016-03-30 15:00:00 -07:00
|
|
|
"""
|
|
|
|
Initialize the simulation.
|
|
|
|
|
|
|
|
:param epsilon: List containing [eps_r,xx, eps_r,yy, eps_r,zz], where each element is a Yee-shifted ndarray
|
|
|
|
spanning the simulation domain. Relative epsilon is used.
|
2017-10-01 12:34:11 -07:00
|
|
|
:param pmls: List of dicts with keys:
|
|
|
|
'axis': One of 'x', 'y', 'z'.
|
|
|
|
'direction': One of 'n', 'p'.
|
|
|
|
'thickness': Number of layers, default 8.
|
|
|
|
'epsilon_eff': Effective epsilon to match to. Default 1.0.
|
|
|
|
'mu_eff': Effective mu to match to. Default 1.0.
|
|
|
|
'ln_R_per_layer': Desired (ln(R) / thickness) value. Default -1.6.
|
|
|
|
'm': Polynomial grading exponent. Default 3.5.
|
|
|
|
'ma': Exponent for alpha. Default 1.
|
2018-11-30 01:03:20 -08:00
|
|
|
:param bloch_boundaries: List of dicts with keys:
|
|
|
|
'axis': One of 'x', 'y', 'z'.
|
|
|
|
'real': Real part of bloch phase factor (i.e. real(exp(i * phase)))
|
|
|
|
'imag': Imaginary part of bloch phase factor (i.e. imag(exp(i * phase)))
|
2018-11-30 00:36:11 -08:00
|
|
|
:param dt: Time step. Default is min(dxes) * .99/sqrt(3).
|
2018-11-30 01:03:20 -08:00
|
|
|
:param initial_fields: Dict with optional keys ('E', 'H', 'F', 'G') containing initial values for the
|
|
|
|
specified fields (default is 0 everywhere). Fields have same format as epsilon.
|
2016-03-30 15:00:00 -07:00
|
|
|
:param context: pyOpenCL context. If not given, pyopencl.create_some_context(False) is called.
|
|
|
|
:param queue: pyOpenCL command queue. If not given, pyopencl.CommandQueue(context) is called.
|
|
|
|
:param float_type: numpy.float32 or numpy.float64. Default numpy.float32.
|
2017-08-24 11:28:03 -07:00
|
|
|
:param do_poynting: If true, enables calculation of the poynting vector, S.
|
|
|
|
Poynting vector calculation adds the following computational burdens:
|
2019-07-15 00:11:09 -07:00
|
|
|
****INACCURATE, TODO FIXME*****
|
2017-08-24 11:28:03 -07:00
|
|
|
* During update_H, ~6 extra additions/cell are performed in order to spatially
|
|
|
|
average E and temporally average H. These quantities are multiplied
|
|
|
|
(6 multiplications/cell) and then stored (6 writes/cell, cache-friendly).
|
|
|
|
* update_S performs a discrete cross product using the precalculated products
|
|
|
|
from update_H. This is not nice to the cache and similar to e.g. update_E
|
|
|
|
in complexity.
|
|
|
|
* GPU memory requirements are approximately doubled, since S and the intermediate
|
|
|
|
products must be stored.
|
2016-03-30 15:00:00 -07:00
|
|
|
"""
|
2018-11-29 01:59:05 -08:00
|
|
|
if initial_fields is None:
|
|
|
|
initial_fields = {}
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2018-11-29 01:59:05 -08:00
|
|
|
self.shape = epsilon[0].shape
|
|
|
|
self.arg_type = float_type
|
|
|
|
self.sources = {}
|
|
|
|
self._create_context(context, queue)
|
|
|
|
self._create_eps(epsilon)
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2018-11-30 00:36:11 -08:00
|
|
|
if dxes is None:
|
|
|
|
dxes = 1.0
|
|
|
|
|
|
|
|
if isinstance(dxes, (float, int)):
|
|
|
|
uniform_dx = dxes
|
|
|
|
min_dx = dxes
|
|
|
|
else:
|
|
|
|
uniform_dx = False
|
|
|
|
self.inv_dxes = [self._create_field(1 / dxn) for dxn in dxes[0] + dxes[1]]
|
|
|
|
min_dx = min(min(dxn) for dxn in dxes[0] + dxes[1])
|
|
|
|
|
|
|
|
max_dt = min_dx * .99 / numpy.sqrt(3)
|
|
|
|
|
|
|
|
if dt is None:
|
|
|
|
self.dt = max_dt
|
|
|
|
elif dt > max_dt:
|
2016-03-30 15:00:00 -07:00
|
|
|
warnings.warn('Warning: unstable dt: {}'.format(dt))
|
|
|
|
elif dt <= 0:
|
|
|
|
raise Exception('Invalid dt: {}'.format(dt))
|
|
|
|
else:
|
|
|
|
self.dt = dt
|
|
|
|
|
2018-11-29 01:59:05 -08:00
|
|
|
self.E = self._create_field(initial_fields.get('E', None))
|
|
|
|
self.H = self._create_field(initial_fields.get('H', None))
|
2018-11-29 02:01:48 -08:00
|
|
|
if bloch_boundaries:
|
|
|
|
self.F = self._create_field(initial_fields.get('F', None))
|
|
|
|
self.G = self._create_field(initial_fields.get('G', None))
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2017-10-01 12:34:11 -07:00
|
|
|
for pml in pmls:
|
|
|
|
pml.setdefault('thickness', 8)
|
|
|
|
pml.setdefault('epsilon_eff', 1.0)
|
|
|
|
pml.setdefault('mu_eff', 1.0)
|
|
|
|
pml.setdefault('ln_R_per_layer', -1.6)
|
|
|
|
pml.setdefault('m', 3.5)
|
2018-11-30 00:36:32 -08:00
|
|
|
pml.setdefault('cfs_alpha', 0)
|
2017-10-01 12:34:11 -07:00
|
|
|
pml.setdefault('ma', 1)
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
ctype = type_to_C(self.arg_type)
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
def ptr(arg: str) -> str:
|
2018-11-30 00:37:07 -08:00
|
|
|
return ctype + ' * restrict ' + arg
|
2016-09-01 14:39:44 -07:00
|
|
|
|
|
|
|
base_fields = OrderedDict()
|
|
|
|
base_fields[ptr('E')] = self.E
|
|
|
|
base_fields[ptr('H')] = self.H
|
|
|
|
base_fields[ctype + ' dt'] = self.dt
|
2018-11-30 00:36:11 -08:00
|
|
|
if uniform_dx == False:
|
|
|
|
inv_dx_names = ['inv_d' + eh + r for eh in 'eh' for r in 'xyz']
|
|
|
|
for name, field in zip(inv_dx_names, self.inv_dxes):
|
|
|
|
base_fields[ptr(name)] = field
|
2016-09-01 14:39:44 -07:00
|
|
|
|
|
|
|
eps_field = OrderedDict()
|
|
|
|
eps_field[ptr('eps')] = self.eps
|
|
|
|
|
2018-11-29 02:01:48 -08:00
|
|
|
if bloch_boundaries:
|
|
|
|
base_fields[ptr('F')] = self.F
|
|
|
|
base_fields[ptr('G')] = self.G
|
|
|
|
|
|
|
|
bloch_fields = OrderedDict()
|
|
|
|
bloch_fields[ptr('E')] = self.F
|
|
|
|
bloch_fields[ptr('H')] = self.G
|
|
|
|
bloch_fields[ctype + ' dt'] = self.dt
|
|
|
|
bloch_fields[ptr('F')] = self.E
|
|
|
|
bloch_fields[ptr('G')] = self.H
|
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
common_source = jinja_env.get_template('common.cl').render(
|
|
|
|
ftype=ctype,
|
2018-11-29 01:59:05 -08:00
|
|
|
shape=self.shape,
|
2016-09-01 14:39:44 -07:00
|
|
|
)
|
|
|
|
jinja_args = {
|
|
|
|
'common_header': common_source,
|
|
|
|
'pmls': pmls,
|
|
|
|
'do_poynting': do_poynting,
|
2018-11-29 02:01:48 -08:00
|
|
|
'bloch': bloch_boundaries,
|
2018-11-30 00:36:11 -08:00
|
|
|
'uniform_dx': uniform_dx,
|
2016-09-01 14:39:44 -07:00
|
|
|
}
|
|
|
|
E_source = jinja_env.get_template('update_e.cl').render(**jinja_args)
|
|
|
|
H_source = jinja_env.get_template('update_h.cl').render(**jinja_args)
|
|
|
|
self.sources['E'] = E_source
|
|
|
|
self.sources['H'] = H_source
|
|
|
|
|
2018-11-29 02:01:48 -08:00
|
|
|
if bloch_boundaries:
|
|
|
|
bloch_args = jinja_args.copy()
|
|
|
|
bloch_args['do_poynting'] = False
|
|
|
|
bloch_args['bloch'] = [{'axis': b['axis'],
|
|
|
|
'real': b['imag'],
|
|
|
|
'imag': b['real']}
|
|
|
|
for b in bloch_boundaries]
|
|
|
|
F_source = jinja_env.get_template('update_e.cl').render(**bloch_args)
|
|
|
|
G_source = jinja_env.get_template('update_h.cl').render(**bloch_args)
|
|
|
|
self.sources['F'] = F_source
|
|
|
|
self.sources['G'] = G_source
|
2018-11-29 01:59:05 -08:00
|
|
|
|
|
|
|
|
|
|
|
S_fields = OrderedDict()
|
2016-09-01 14:39:44 -07:00
|
|
|
if do_poynting:
|
|
|
|
self.S = pyopencl.array.zeros_like(self.E)
|
|
|
|
S_fields[ptr('S')] = self.S
|
|
|
|
|
2018-11-29 02:01:28 -08:00
|
|
|
J_fields = OrderedDict()
|
|
|
|
if do_fieldsrc:
|
|
|
|
J_source = jinja_env.get_template('update_j.cl').render(**jinja_args)
|
|
|
|
self.sources['J'] = J_source
|
|
|
|
|
|
|
|
self.Ji = pyopencl.array.zeros_like(self.E)
|
|
|
|
self.Jr = pyopencl.array.zeros_like(self.E)
|
|
|
|
J_fields[ptr('Jr')] = self.Jr
|
|
|
|
J_fields[ptr('Ji')] = self.Ji
|
|
|
|
|
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
'''
|
|
|
|
PML
|
|
|
|
'''
|
2018-11-29 01:59:05 -08:00
|
|
|
pml_e_fields, pml_h_fields = self._create_pmls(pmls)
|
2018-11-29 02:01:48 -08:00
|
|
|
if bloch_boundaries:
|
|
|
|
pml_f_fields, pml_g_fields = self._create_pmls(pmls)
|
2018-11-29 01:59:05 -08:00
|
|
|
|
|
|
|
'''
|
|
|
|
Create operations
|
|
|
|
'''
|
|
|
|
self.update_E = self._create_operation(E_source, (base_fields, eps_field, pml_e_fields))
|
|
|
|
self.update_H = self._create_operation(H_source, (base_fields, pml_h_fields, S_fields))
|
2018-11-29 02:01:48 -08:00
|
|
|
if bloch_boundaries:
|
|
|
|
self.update_F = self._create_operation(F_source, (bloch_fields, eps_field, pml_f_fields))
|
|
|
|
self.update_G = self._create_operation(G_source, (bloch_fields, pml_g_fields))
|
2018-11-29 02:01:28 -08:00
|
|
|
if do_fieldsrc:
|
|
|
|
args = OrderedDict()
|
|
|
|
[args.update(d) for d in (base_fields, J_fields)]
|
|
|
|
var_args = [ctype + ' ' + v for v in 'cs'] + ['uint ' + r + m for r in 'xyz' for m in ('min', 'max')]
|
|
|
|
update = ElementwiseKernel(self.context, operation=J_source,
|
|
|
|
arguments=', '.join(list(args.keys()) + var_args))
|
|
|
|
self.update_J = lambda e, *a: update(*args.values(), *a, wait_for=e)
|
2018-11-29 01:59:05 -08:00
|
|
|
|
|
|
|
|
|
|
|
def _create_pmls(self, pmls):
|
|
|
|
ctype = type_to_C(self.arg_type)
|
|
|
|
def ptr(arg: str) -> str:
|
|
|
|
return ctype + ' *' + arg
|
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
pml_e_fields = OrderedDict()
|
|
|
|
pml_h_fields = OrderedDict()
|
|
|
|
for pml in pmls:
|
2017-10-01 12:34:11 -07:00
|
|
|
a = 'xyz'.find(pml['axis'])
|
|
|
|
|
2018-11-29 02:00:30 -08:00
|
|
|
sigma_max = -pml['ln_R_per_layer'] / 2 * (pml['m'] + 1)
|
|
|
|
kappa_max = numpy.sqrt(pml['mu_eff'] * pml['epsilon_eff'])
|
2018-11-30 00:36:32 -08:00
|
|
|
alpha_max = pml['cfs_alpha']
|
2017-10-01 12:34:11 -07:00
|
|
|
|
|
|
|
def par(x):
|
2019-07-14 23:51:57 -07:00
|
|
|
scaling = (x / pml['thickness']) ** pml['m']
|
2018-11-29 02:00:30 -08:00
|
|
|
sigma = scaling * sigma_max
|
|
|
|
kappa = 1 + scaling * (kappa_max - 1)
|
2017-10-01 12:34:11 -07:00
|
|
|
alpha = ((1 - x / pml['thickness']) ** pml['ma']) * alpha_max
|
2018-11-29 02:00:30 -08:00
|
|
|
p0 = numpy.exp(-(sigma / kappa + alpha) * self.dt)
|
|
|
|
p1 = sigma / (sigma + kappa * alpha) * (p0 - 1)
|
2019-07-14 23:51:57 -07:00
|
|
|
p2 = 1 / kappa
|
2018-11-29 02:00:30 -08:00
|
|
|
return p0, p1, p2
|
2017-10-01 12:34:11 -07:00
|
|
|
|
2018-11-29 01:59:05 -08:00
|
|
|
xe, xh = (numpy.arange(1, pml['thickness'] + 1, dtype=self.arg_type)[::-1] for _ in range(2))
|
2017-10-01 12:34:11 -07:00
|
|
|
if pml['polarity'] == 'p':
|
|
|
|
xe -= 0.5
|
|
|
|
elif pml['polarity'] == 'n':
|
|
|
|
xh -= 0.5
|
|
|
|
|
2018-11-29 02:00:30 -08:00
|
|
|
pml_p_names = [['p' + pml['axis'] + i + eh + pml['polarity'] for i in '012'] for eh in 'eh']
|
2017-10-01 12:34:11 -07:00
|
|
|
for name_e, name_h, pe, ph in zip(pml_p_names[0], pml_p_names[1], par(xe), par(xh)):
|
|
|
|
pml_e_fields[ptr(name_e)] = pyopencl.array.to_device(self.queue, pe)
|
|
|
|
pml_h_fields[ptr(name_h)] = pyopencl.array.to_device(self.queue, ph)
|
|
|
|
|
|
|
|
uv = 'xyz'.replace(pml['axis'], '')
|
|
|
|
psi_base = 'Psi_' + pml['axis'] + pml['polarity'] + '_'
|
2016-09-01 14:39:44 -07:00
|
|
|
psi_names = [[psi_base + eh + c for c in uv] for eh in 'EH']
|
|
|
|
|
2018-11-29 01:59:05 -08:00
|
|
|
psi_shape = list(self.shape)
|
2017-10-01 12:34:11 -07:00
|
|
|
psi_shape[a] = pml['thickness']
|
2016-09-01 14:39:44 -07:00
|
|
|
|
|
|
|
for ne, nh in zip(*psi_names):
|
|
|
|
pml_e_fields[ptr(ne)] = pyopencl.array.zeros(self.queue, tuple(psi_shape), dtype=self.arg_type)
|
|
|
|
pml_h_fields[ptr(nh)] = pyopencl.array.zeros(self.queue, tuple(psi_shape), dtype=self.arg_type)
|
2018-11-29 01:59:05 -08:00
|
|
|
return pml_e_fields, pml_h_fields
|
2017-08-14 15:41:20 -07:00
|
|
|
|
2018-11-29 01:59:05 -08:00
|
|
|
def _create_operation(self, source, args_fields):
|
|
|
|
args = OrderedDict()
|
|
|
|
[args.update(d) for d in args_fields]
|
|
|
|
update = ElementwiseKernel(self.context, operation=source,
|
|
|
|
arguments=', '.join(args.keys()))
|
|
|
|
return lambda e: update(*args.values(), wait_for=e)
|
2016-09-01 14:39:44 -07:00
|
|
|
|
2018-11-29 01:59:05 -08:00
|
|
|
def _create_context(self, context: pyopencl.Context = None,
|
|
|
|
queue: pyopencl.CommandQueue = None):
|
|
|
|
if context is None:
|
|
|
|
self.context = pyopencl.create_some_context()
|
|
|
|
else:
|
|
|
|
self.context = context
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2018-11-29 01:59:05 -08:00
|
|
|
if queue is None:
|
|
|
|
self.queue = pyopencl.CommandQueue(self.context)
|
|
|
|
else:
|
|
|
|
self.queue = queue
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2018-11-29 01:38:28 -08:00
|
|
|
def _create_eps(self, epsilon: List[numpy.ndarray]):
|
|
|
|
if len(epsilon) != 3:
|
|
|
|
raise Exception('Epsilon must be a list with length of 3')
|
|
|
|
if not all((e.shape == epsilon[0].shape for e in epsilon[1:])):
|
|
|
|
raise Exception('All epsilon grids must have the same shape. Shapes are {}', [e.shape for e in epsilon])
|
|
|
|
if not epsilon[0].shape == self.shape:
|
|
|
|
raise Exception('Epsilon shape mismatch. Expected {}, got {}'.format(self.shape, epsilon[0].shape))
|
|
|
|
self.eps = pyopencl.array.to_device(self.queue, vec(epsilon).astype(self.arg_type))
|
|
|
|
|
|
|
|
def _create_field(self, initial_value: List[numpy.ndarray] = None):
|
|
|
|
if initial_value is None:
|
|
|
|
return pyopencl.array.zeros_like(self.eps)
|
|
|
|
else:
|
|
|
|
if len(initial_value) != 3:
|
|
|
|
Exception('Initial field value must be a list of length 3')
|
|
|
|
if not all((f.shape == self.shape for f in initial_value)):
|
|
|
|
Exception('Initial field list elements must have same shape as epsilon elements')
|
|
|
|
return pyopencl.array.to_device(self.queue, vec(initial_value).astype(self.arg_type))
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2018-11-30 00:37:16 -08:00
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
def type_to_C(float_type) -> str:
|
|
|
|
"""
|
|
|
|
Returns a string corresponding to the C equivalent of a numpy type.
|
|
|
|
Only works for float16, float32, float64.
|
2016-03-30 15:00:00 -07:00
|
|
|
|
2016-09-01 14:39:44 -07:00
|
|
|
:param float_type: e.g. numpy.float32
|
|
|
|
:return: string containing the corresponding C type (eg. 'double')
|
|
|
|
"""
|
|
|
|
if float_type == numpy.float16:
|
|
|
|
arg_type = 'half'
|
|
|
|
elif float_type == numpy.float32:
|
|
|
|
arg_type = 'float'
|
|
|
|
elif float_type == numpy.float64:
|
|
|
|
arg_type = 'double'
|
|
|
|
else:
|
|
|
|
raise Exception('Unsupported type')
|
|
|
|
return arg_type
|