def __init__(self, nodef):
common.check_type('nodef', nodef, node.Fields)
self.directions = nodef.buffer_dict.keys()
self.gpu = gpu
self.cpu = cpu
if 'x+' in self.directions:
mf_xp = nodef.mainf_list[-1]
bf_xp = nodef.buffer_dict['x+']
mpu = getattr(self, mf_xp.device_type)
self.gf_p_h = mpu.GetFields(mf_xp, ['hy', 'hz'], (-2, 0, 0),
(-2, -1, -1))
self.sf_p_h = cpu.SetFields(bf_xp, ['hy', 'hz'], (0, 0, 0),
(0, -1, -1), True)
self.sf_p_e = mpu.SetFields(mf_xp, ['ey', 'ez'], (-1, 0, 0),
(-1, -1, -1), True)
self.gf_h = cpu.GetFields(bf_xp, ['hy', 'hz'], (1, 0, 0),
(1, -1, -1))
self.sf_e = cpu.SetFields(bf_xp, ['ey', 'ez'], (1, 0, 0),
(1, -1, -1), True)
target = 0 if rank == size - 1 else rank + 1
self.req_send_h = comm.Send_init(self.gf_h.host_array,
target,
tag=0)
self.tmp_recv_e = np.zeros(self.gf_h.host_array.shape, nodef.dtype)
self.req_recv_e = comm.Recv_init(self.tmp_recv_e, target, tag=1)
if 'x-' in self.directions:
mf_xm = nodef.mainf_list[0]
bf_xm = nodef.buffer_dict['x-']
mpu = getattr(self, mf_xm.device_type)
self.gf_m_e = mpu.GetFields(mf_xm, ['ey', 'ez'], (1, 0, 0),
(1, -1, -1))
self.sf_m_e = cpu.SetFields(bf_xm, ['ey', 'ez'], (-1, 0, 0),
(-1, -1, -1), True)
self.sf_m_h = mpu.SetFields(mf_xm, ['hy', 'hz'], (0, 0, 0),
(0, -1, -1), True)
self.gf_e = cpu.GetFields(bf_xm, ['ey', 'ez'], (0, 0, 0),
(0, -1, -1))
self.sf_h = cpu.SetFields(bf_xm, ['hy', 'hz'], (0, 0, 0),
(0, -1, -1), True)
target = size - 1 if rank == 0 else rank - 1
self.req_send_e = comm.Send_init(self.gf_e.host_array,
target,
tag=1)
self.tmp_recv_h = np.zeros(self.gf_e.host_array.shape, nodef.dtype)
self.req_recv_h = comm.Recv_init(self.tmp_recv_h, target, tag=0)
# append to the update list
self.priority_type = 'mpi'
nodef.append_instance(self)
def __init__(self, mainf_list):
"""
"""
common.check_type('mainf_list', mainf_list, \
(list, tuple), (gpu.Fields, cpu.Fields))
# local variables
cpuf = mainf_list[-1]
cpuf_dict = {cpuf.mpi_type: cpuf}
nx_list = [f.nx for f in mainf_list]
nx = sum(nx_list) - len(nx_list) + 1
ny, nz = cpuf.ny, cpuf.nz
accum_nx_list = [0] + \
list( np.add.accumulate([f.nx-1 for f in mainf_list]) )
# global variables
self.mainf_list = mainf_list
self.cpuf_dict = cpuf_dict
self.updatef_list = mainf_list[:]
self.dtype = cpuf.dtype
self.nx = nx
self.nx_list = nx_list
self.accum_nx_list = accum_nx_list
self.ns = (nx, ny, nz)
def __init__(self, direction, target_rank, ny, nz, coeff_use,
precision_float):
"""
"""
super(BufferFields, self).__init__(3,
ny,
nz,
coeff_use,
precision_float,
use_cpu_core=0)
common.check_value('direction', direction,
('x+', 'x-', 'y+', 'y-', 'z+', 'z-'))
common.check_type('target_rank', target_rank, int)
# global variables
self.direction = direction
self.target_rank = target_rank
p_or_m = direction[-1]
self.part_e_list = {'+': [''], '-': ['pre', 'post']}[p_or_m]
self.part_h_list = {'-': [''], '+': ['pre', 'post']}[p_or_m]
self.is_split_dict = { \
'+': {'e': False, 'h': True}, \
'-': {'e': True, 'h': False}}[p_or_m]
def __init__(self, gpuf, direction):
common.check_type('gpuf', gpuf, gpu.Fields)
common.check_value('direction', direction, ('+', '-', '+-'))
if '+' in direction:
self.getf_h = getf_h = gpu.GetFields(gpuf, ['hy', 'hz'],
(-1, 0, 0), (-1, -1, -1))
self.setf_e = gpu.SetFields(gpuf, ['ey', 'ez'], (-1, 0, 0),
(-1, -1, -1), True)
self.req_send_h = comm.Send_init(getf_h.host_array,
rank + 1,
tag=0)
self.tmp_recv_e = tmp_recv_e = np.zeros(getf_h.host_array.shape,
gpuf.dtype)
self.req_recv_e = comm.Recv_init(tmp_recv_e, rank + 1, tag=1)
if '-' in direction:
self.getf_e = getf_e = gpu.GetFields(gpuf, ['ey', 'ez'], (0, 0, 0),
(0, -1, -1))
self.setf_h = gpu.SetFields(gpuf, ['hy', 'hz'], (0, 0, 0),
(0, -1, -1), True)
self.req_send_e = comm.Send_init(getf_e.host_array,
rank - 1,
tag=1)
self.tmp_recv_h = tmp_recv_h = np.zeros(getf_e.host_array.shape,
gpuf.dtype)
self.req_recv_h = comm.Recv_init(tmp_recv_h, rank - 1, tag=0)
# global variables
self.direction = direction
def __init__(self, gpuf, core, direction):
common.check_type('gpuf', gpuf, Fields)
common.check_value('direction', direction, ('+', '-', '+-'))
if '+' in direction:
self.gf_h = gf_h = GetFields(gpuf, ['hy', 'hz'], (-1, 0, 0),
(-1, -1, -1))
self.sf_e = SetFields(gpuf, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1),
True)
target = 0 if rank == size - 1 else rank + 1
self.req_send_h = comm.Send_init(gf_h.host_array, target, tag=0)
self.tmp_recv_e = np.zeros(gf_h.host_array.shape, gpuf.dtype)
self.req_recv_e = comm.Recv_init(self.tmp_recv_e, target, tag=1)
if '-' in direction:
self.gf_e = gf_e = GetFields(gpuf, ['ey', 'ez'], (0, 0, 0),
(0, -1, -1))
self.sf_h = SetFields(gpuf, ['hy', 'hz'], (0, 0, 0), (0, -1, -1),
True)
target = size - 1 if rank == 0 else rank - 1
self.req_send_e = comm.Send_init(gf_e.host_array, target, tag=1)
self.tmp_recv_h = np.zeros(gf_e.host_array.shape, gpuf.dtype)
self.req_recv_h = comm.Recv_init(self.tmp_recv_h, target, tag=0)
# global variables
self.core = core
self.direction = direction
def append_buffer_fields(self, cpuf):
common.check_type('cpuf', cpuf, cpu.Fields)
common.check_value('cpuf.mpi_type', cpuf.mpi_type, \
('x+', 'x-', 'y+', 'y-', 'z+', 'z-') )
self.cpuf_dict[cpuf.mpi_type] = cpuf
self.updatef_list.append(cpuf)
def __init__(self, gpuf, direction):
common.check_type('gpuf', gpuf, gpu.Fields)
common.check_value('direction', direction, ('+', '-', '+-'))
qtask = cpu.QueueTask()
if '+' in direction:
self.cpuf_p = cpuf_p = cpu.Fields(qtask,
3,
gpuf.ny,
gpuf.nz,
gpuf.coeff_use,
gpuf.precision_float,
use_cpu_core=0)
self.gf_p_h = gpu.GetFields(gpuf, ['hy', 'hz'], (-2, 0, 0),
(-2, -1, -1))
self.sf_p_h = cpu.SetFields(cpuf_p, ['hy', 'hz'], (0, 0, 0),
(0, -1, -1), True)
self.gf_p_e = cpu.GetFields(cpuf_p, ['ey', 'ez'], (1, 0, 0),
(1, -1, -1))
self.sf_p_e = gpu.SetFields(gpuf, ['ey', 'ez'], (-1, 0, 0),
(-1, -1, -1), True)
self.gf_h = cpu.GetFields(cpuf_p, ['hy', 'hz'], (1, 0, 0),
(1, -1, -1))
self.sf_e = cpu.SetFields(cpuf_p, ['ey', 'ez'], (-1, 0, 0),
(-1, -1, -1), True)
self.tmp_recv_e = np.zeros(self.gf_h.host_array.shape, gpuf.dtype)
if '-' in direction:
self.cpuf_m = cpuf_m = cpu.Fields(qtask,
3,
gpuf.ny,
gpuf.nz,
gpuf.coeff_use,
gpuf.precision_float,
use_cpu_core=0)
self.gf_m_e = gpu.GetFields(gpuf, ['ey', 'ez'], (1, 0, 0),
(1, -1, -1))
self.sf_m_e = cpu.SetFields(cpuf_m, ['ey', 'ez'], (-1, 0, 0),
(-1, -1, -1), True)
self.gf_m_h = cpu.GetFields(cpuf_m, ['hy', 'hz'], (1, 0, 0),
(1, -1, -1))
self.sf_m_h = gpu.SetFields(gpuf, ['hy', 'hz'], (0, 0, 0),
(0, -1, -1), True)
self.gf_e = cpu.GetFields(cpuf_m, ['ey', 'ez'], (1, 0, 0),
(1, -1, -1))
self.sf_h = cpu.SetFields(cpuf_m, ['hy', 'hz'], (0, 0, 0),
(0, -1, -1), True)
self.tmp_recv_h = np.zeros(self.gf_e.host_array.shape, gpuf.dtype)
# global variables
self.direction = direction
self.qtask = qtask
def __init__(self, gpuf, core, direction, tmax):
common.check_type('gpuf', gpuf, Fields)
common.check_value('direction', direction, ('+', '-', '+-'))
if '+' in direction:
self.gf_h = gf_h = GetFields(gpuf, ['hy', 'hz'], (-1, 0, 0), (-1, -1, -1))
self.sf_e = SetFields(gpuf, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.req_send_h = comm.Send_init(gf_h.host_array, rank+1, tag=0)
self.tmp_recv_e_list = [np.zeros(gf_h.host_array.shape, gpuf.dtype) for i in range(2)]
self.req_recv_e_list = [comm.Recv_init(tmp_recv_e, rank+1, tag=1) for tmp_recv_e in self.tmp_recv_e_list]
self.switch_e = 0
if '-' in direction:
self.gf_e = gf_e = GetFields(gpuf, ['ey', 'ez'], (0, 0, 0), (0, -1, -1))
self.sf_h = SetFields(gpuf, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
self.req_send_e = comm.Send_init(gf_e.host_array, rank-1, tag=1)
self.tmp_recv_h_list = [np.zeros(gf_e.host_array.shape, gpuf.dtype) for i in range(2)]
self.req_recv_h_list = [comm.Recv_init(tmp_recv_h, rank-1, tag=0) for tmp_recv_h in self.tmp_recv_h_list]
self.switch_h = 0
# global variables
self.gpuf = gpuf
self.core = core
self.direction = direction
self.tmax = tmax
self.tstep = 1
def __init__(self, mainf_list):
"""
"""
common.check_type('mainf_list', mainf_list, \
(list, tuple), (gpu.Fields, cpu.Fields))
# local variables
cpuf = mainf_list[-1]
cpuf_dict = {cpuf.mpi_type: cpuf}
nx_list = [f.nx for f in mainf_list]
nx = sum(nx_list) - len(nx_list) + 1
ny, nz = cpuf.ny, cpuf.nz
accum_nx_list = [0] + \
list( np.add.accumulate([f.nx-1 for f in mainf_list]) )
# global variables
self.mainf_list = mainf_list
self.cpuf_dict = cpuf_dict
self.updatef_list = mainf_list[:]
self.dtype = cpuf.dtype
self.nx = nx
self.nx_list = nx_list
self.accum_nx_list = accum_nx_list
self.ns = (nx, ny, nz)
def set_ehs(self, ex, ey, ez, hx, hy, hz): # for unittest
eh_list = [ex, ey, ez, hx, hy, hz]
for eh in eh_list:
common.check_type('eh', eh, np.ndarray)
for eh, f in zip(self.ehs, eh_list):
eh[:, :, self.slice_z] = f[:]
def __init__(self, mainf_list, buffer_dict={}):
"""
"""
common.check_type('buffer_dict', buffer_dict, dict)
# local variables
nx_list = [f.nx for f in mainf_list]
nx = int(sum(nx_list) - len(nx_list) + 1)
ny, nz = [int(n) for n in mainf_list[0].ns[1:]]
accum_nx_list = np.add.accumulate([0] + [f.nx - 1 for f in mainf_list])
#accum_nx_list[-1] += 1
accum_nx_list = [int(anx) for anx in accum_nx_list]
# global variables
self.mainf_list = mainf_list
self.buffer_dict = buffer_dict
self.updatef_list = mainf_list[:] + buffer_dict.values()
self.dtype = mainf_list[0].dtype
self.nx = nx
self.nx_list = nx_list
self.accum_nx_list = accum_nx_list
self.ns = (nx, ny, nz)
# update list
self.instance_list = []
self.append_instance = lambda instance: \
common.append_instance(self.instance_list, instance)
def __init__(self, node_fields):
"""
"""
common.check_type('node_fields', node_fields, Fields)
# local variables
nodef = node_fields
mainf_list = nodef.mainf_list
# global variables
device_type_list = [f.device_type for f in nodef.mainf_list]
if 'gpu' in device_type_list:
from kemp.fdtd3d import gpu
self.gpu = gpu
if 'cpu' in device_type_list:
from kemp.fdtd3d import cpu
self.cpu = cpu
self.getf_dict = {'e': [], 'h': []}
self.setf_dict = {'e': [], 'h': []}
self.getf_block_dict = {'e': [], 'h': []}
self.setf_block_dict = {'e': [], 'h': []}
for f0, f1 in zip(mainf_list[:-1], mainf_list[1:]):
for eh in ['e', 'h']:
self.put_getf_setf_list(eh, f0, f1)
# append to the update list
self.priority_type = 'exchange'
nodef.append_instance(self)
def append_buffer_fields(self, cpuf):
common.check_type('cpuf', cpuf, cpu.Fields)
common.check_value('cpuf.mpi_type', cpuf.mpi_type, \
('x+', 'x-', 'y+', 'y-', 'z+', 'z-') )
self.cpuf_dict[cpuf.mpi_type] = cpuf
self.updatef_list.append(cpuf)
def __init__(self, fields):
common.check_type('fields', fields, Fields)
# global variables
self.mainf = fields
if fields.ce_on:
cex, cey, cez = fields.ces
self.ces = cex[:, :-1, :-1], cey[:-1, :, :-1], cez[:-1, :-1, :]
else:
self.ces = fields.ces
if fields.ch_on:
chx, chy, chz = fields.chs
self.chs = chx[:, 1:, 1:], cey[1:, :, 1:], cez[1:, 1:, :]
else:
self.chs = fields.chs
if fields.rd_on:
erdx, erdy, erdz = fields.erds
hrdx, hrdy, hrdz = fields.hrds
sln = slice(None, None)
nax = np.newaxis
self.erds = erdx[sln, nax, nax], erdy[sln, nax], erdz[sln]
self.hrds = hrdx[sln, nax, nax], hrdy[sln, nax], hrdz[sln]
else:
self.erds = fields.erds
self.hrds = fields.hrds
# append to the update list
self.priority_type = 'core'
fields.append_instance(self)
def __init__(self, nodef, direction):
common.check_type('nodef', nodef, node.Fields)
common.check_value('direction', direction, ('+', '-', '+-'))
self.gpu = gpu
self.cpu = cpu
if '+' in direction:
mf_p = nodef.mainf_list[-1]
mpu = getattr(self, mf_p.device_type)
self.getf_h = mpu.GetFields(mf_p, ['hy', 'hz'], (-1, 0, 0), (-1, -1, -1))
self.setf_e = mpu.SetFields(mf_p, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.tmp_recv_e = np.zeros(self.getf_h.host_array.shape, nodef.dtype)
if '-' in direction:
mf_m = nodef.mainf_list[0]
mpu = getattr(self, mf_m.device_type)
self.getf_e = mpu.GetFields(mf_m, ['ey', 'ez'], (0, 0, 0), (0, -1, -1))
self.setf_h = mpu.SetFields(mf_m, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
self.tmp_recv_h = np.zeros(self.getf_e.host_array.shape, nodef.dtype)
# global variables
self.direction = direction
self.target_p = 0 if rank == size-1 else rank+1
self.target_m = size-1 if rank == 0 else rank-1
# append to the update list
self.priority_type = 'mpi'
nodef.append_instance(self)
def __init__(self, gpuf, core, direction):
common.check_type('gpuf', gpuf, Fields)
common.check_value('direction', direction, ('+', '-', '+-'))
if '+' in direction:
self.gf_h = gf_h = GetFields(gpuf, ['hy', 'hz'], (-1, 0, 0), (-1, -1, -1))
self.sf_e = SetFields(gpuf, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
target = 0 if rank == size-1 else rank+1
self.req_send_h = comm.Send_init(gf_h.host_array, target, tag=0)
self.tmp_recv_e = np.zeros(gf_h.host_array.shape, gpuf.dtype)
self.req_recv_e = comm.Recv_init(self.tmp_recv_e, target, tag=1)
if '-' in direction:
self.gf_e = gf_e = GetFields(gpuf, ['ey', 'ez'], (0, 0, 0), (0, -1, -1))
self.sf_h = SetFields(gpuf, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
target = size-1 if rank == 0 else rank-1
self.req_send_e = comm.Send_init(gf_e.host_array, target, tag=1)
self.tmp_recv_h = np.zeros(gf_e.host_array.shape, gpuf.dtype)
self.req_recv_h = comm.Recv_init(self.tmp_recv_h, target, tag=0)
# global variables
self.core = core
self.direction = direction
def set_ehs(self, ex, ey, ez, hx, hy, hz): # for unittest
eh_list = [ex, ey, ez, hx, hy, hz]
for eh in eh_list:
common.check_type('eh', eh, np.ndarray)
for eh, f in zip(self.ehs, eh_list):
eh[:,:,self.slice_z] = f[:]
def __init__(self, node_fields, axes):
"""
"""
common.check_type('node_fields', node_fields, Fields)
common.check_type('axes', axes, str)
assert len(set(axes).intersection(
set('xyz'))) > 0, 'axes option is wrong: %s is given' % repr(axes)
# local variables
nodef = node_fields
mainf_list = nodef.mainf_list
buffer_dict = nodef.buffer_dict
self.cpu = cpu
if 'gpu' in [f.device_type for f in nodef.updatef_list]:
from kemp.fdtd3d import gpu
self.gpu = gpu
# create pbc instances
f0 = mainf_list[-1]
f1 = mainf_list[0]
if f0 is f1:
getattr(self, f0.device_type).Pbc(f0, axes)
else:
if 'x' in axes:
nx, ny, nz = f0.ns
self.getf_e = getattr(self, f1.device_type).GetFields(
f1, ['ey', 'ez'], (0, 0, 0), (0, ny - 1, nz - 1))
self.setf_e = getattr(self, f0.device_type).SetFields(
f0, ['ey', 'ez'], (nx - 1, 0, 0), (nx - 1, ny - 1, nz - 1),
True)
self.getf_h = getattr(self, f0.device_type).GetFields(
f0, ['hy', 'hz'], (nx - 1, 0, 0), (nx - 1, ny - 1, nz - 1))
self.setf_h = getattr(self, f1.device_type).SetFields(
f1, ['hy', 'hz'], (0, 0, 0), (0, ny - 1, nz - 1), True)
# append to the update list
self.priority_type = 'pbc'
nodef.append_instance(self)
axs = axes.strip('x')
if len(set(axs).intersection(set('yz'))) > 0:
for f in mainf_list:
getattr(self, f.device_type).Pbc(f, axs)
# for buffer fields
for direction, buffer in buffer_dict.items():
replaced_axes = map(
lambda ax: common_buffer.axes_dict[direction[0]][ax],
list(axes))
axs = ''.join(replaced_axes).strip('x')
if axs != '':
cpu.Pbc(buffer, axs)
def __init__(self, fields, axes):
"""
"""
common.check_type('fields', fields, Fields)
common.check_type('axes', axes, str)
assert len(set(axes).intersection(
set('xyz'))) > 0, 'axes option is wrong: %s is given' % repr(axes)
# local variables
nx, ny, nz = fields.ns
dtype_str_list = fields.dtype_str_list
# program
macros = ['NMAX', 'IDX0', 'IDX1', 'DTYPE', 'PRAGMA_fp64']
program_dict = {}
gs_dict = {}
for axis in list(axes):
program_dict[axis] = {}
for e_or_h in ['e', 'h']:
pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['get']
pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['get']
nmaxi_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(
pt0, pt1)
idx0_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str,
zid_str)
pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['set']
pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['set']
nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(
pt0, pt1)
idx1_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str,
zid_str)
values = [nmax_str, idx0_str, idx1_str] + dtype_str_list
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'copy_self.cl').read(), \
macros, values)
program = cl.Program(fields.context, ksrc).build()
program_dict[axis][e_or_h] = program
nmax = int(nmax_str)
remainder = nmax % fields.ls
gs_dict[
axis] = nmax if remainder == 0 else nmax - remainder + fields.ls
# global variables
self.mainf = fields
self.axes = axes
self.program_dict = program_dict
self.gs_dict = gs_dict
# append to the update list
self.priority_type = 'pbc'
self.mainf.append_instance(self)
def __init__(self,
fields,
directions,
npml=10,
sigma_max=4,
kappa_max=1,
alpha_max=0,
m_sigma=3,
m_alpha=1):
common.check_type('fields', fields, Fields)
common.check_type('directions', directions, (list, tuple), str)
assert len(directions) == 3
for axis in directions:
assert axis in ['+', '-', '+-', '']
# local variables
dt = fields.dt
nx, ny, nz = fields.ns
dtype = fields.dtype
# allocations
psi_xs = [np.zeros((2 * npml + 2, ny, nz), dtype) for i in range(4)]
psi_ys = [np.zeros((nx, 2 * npml + 2, nz), dtype) for i in range(4)]
psi_zs = [np.zeros((nx, 2 * npml + 2, nz), dtype) for i in range(4)]
i_e = np.arange(0.5, npml)
i_h = np.arange(1, npml + 1)
sigma_e = sigma_max # * (i_e / npml) ** m_sigma
sigma_h = sigma_max # * (i_h / npml) ** m_sigma
print 'sigma_e', sigma_e
print 'sigma_h', sigma_h
kappa_e = 1 + (kappa_max - 1) * (i_e / npml)**m_sigma
kappa_h = 1 + (kappa_max - 1) * (i_h / npml)**m_sigma
alpha_e = alpha_max * ((npml - i_e) / npml)**m_alpha
alpha_h = alpha_max * ((npml - i_h) / npml)**m_alpha
com_e = (kappa_e * alpha_e + sigma_e) * dt + 2 * kappa_e
com_h = (kappa_h * alpha_h + sigma_h) * dt + 2 * kappa_h
pca_e = 4 * kappa_e / com_e - 1
pca_h = 4 * kappa_h / com_h - 1
pcb_e = (alpha_e * dt - 2 + 4 * kappa_e) / com_e - 1
pcb_h = (alpha_h * dt - 2 + 4 * kappa_h) / com_h - 1
pcc_e = (alpha_e * dt + 2) / com_e - 1
pcc_h = (alpha_h * dt + 2) / com_h - 1
# global variables
self.mainf = fields
self.directions = directions
self.npml = npml
self.psi_xs = psi_xs
self.psi_ys = psi_ys
self.psi_zs = psi_zs
self.pcs_e = [pca_e, pcb_e, pcc_e]
self.pcs_h = [pca_h, pcb_h, pcc_h]
# append to the update list
self.priority_type = 'pml'
fields.append_instance(self)
def set_eh_bufs(self, ex, ey, ez, hx, hy, hz): # for unittest
eh_list = [ex, ey, ez, hx, hy, hz]
for eh in eh_list:
common.check_type('eh', eh, np.ndarray)
pad_arr = np.zeros(self.ns_pad, dtype=self.dtype)
for f_buf, f in zip(self.eh_bufs, eh_list):
f_pitch = np.append(f, pad_arr, 2).copy('C')
cuda.memcpy_htod(f_buf, f_pitch)
def __init__(self,
fields,
directions,
npml=10,
sigma_max=4,
kappa_max=1,
alpha_max=0,
m_sigma=3,
m_alpha=1):
common.check_type('fields', fields, Fields)
common.check_type('directions', directions, (list, tuple), str)
assert len(directions) == 3
for axis in directions:
assert axis in ['+', '-', '+-', '']
# local variables
dt = fields.dt
nx, ny, nz = fields.ns
dtype = fields.dtype
# allocations
psi_xs = [np.zeros((2 * npml + 2, ny, nz), dtype) for i in range(4)]
psi_ys = [np.zeros((nx, 2 * npml + 2, nz), dtype) for i in range(4)]
psi_zs = [np.zeros((nx, 2 * npml + 2, nz), dtype) for i in range(4)]
i_half = np.arange(0.5, npml)
i_one = np.arange(1, npml + 1)
sigma_half = sigma_max * (i_half / npml)**m_sigma
sigma_one = sigma_max * (i_one / npml)**m_sigma
kappa_half = 1 + (kappa_max - 1) * (i_half / npml)**m_sigma
kappa_one = 1 + (kappa_max - 1) * (i_one / npml)**m_sigma
alpha_half = alpha_max * ((npml - i_half) / npml)**m_alpha
alpha_one = alpha_max * ((npml - i_one) / npml)**m_alpha
g_half = (kappa_half * alpha_half + sigma_half) * dt + 2 * kappa_half
g_one = (kappa_one * alpha_one + sigma_one) * dt + 2 * kappa_one
pca_half = 4 * kappa_half / g_half - 1
pca_one = 4 * kappa_one / g_one - 1
pcb_half = (alpha_half * dt - 2 + 4 * kappa_half) / g_half - 1
pcb_one = (alpha_one * dt - 2 + 4 * kappa_one) / g_one - 1
pcc_half = (alpha_half * dt + 2) / g_half - 1
pcc_one = (alpha_one * dt + 2) / g_one - 1
# global variables
self.mainf = fields
self.directions = directions
self.npml = npml
self.psi_xs = psi_xs
self.psi_ys = psi_ys
self.psi_zs = psi_zs
self.pcs_half = [pca_half, pcb_half, pcc_half]
self.pcs_one = [pca_one, pcb_one, pcc_one]
# append to the update list
self.priority_type = 'pml'
fields.append_instance(self)
def set_eh_bufs(self, ex, ey, ez, hx, hy, hz): # for unittest
eh_list = [ex, ey, ez, hx, hy, hz]
for eh in eh_list:
common.check_type('eh', eh, np.ndarray)
pad_arr = np.zeros(self.ns_pad, dtype=self.dtype)
for f_buf, f in zip(self.eh_bufs, eh_list):
f_pitch = np.append(f, pad_arr, 2).copy('C')
cl.enqueue_copy(self.queue, f_buf, f_pitch)
def __init__(self, node_fields, axis):
"""
"""
common.check_type('node_fields', node_fields, NodeFields)
common.check_value('axis', axis, ['x', 'y', 'z'])
# local variables
self.gpu = gpu
self.cpu = cpu
mainf_list = node_fields.mainf_list
cpuf_dict = node_fields.cpuf_dict
axis_id = {'x': 0, 'y': 1, 'z': 2}[axis]
set_cpuf = set(cpuf_dict.keys())
for ax in ['x', 'y', 'z']:
if not set_cpuf.isdisjoint([ax + '+', ax + '-']):
raise ValueError, 'There are %s-axis buffer instances. The pbc_internal operation along %s-axis is not allowed.' % (
ax, ax)
# create pbc instances
f0 = mainf_list[0]
f1 = mainf_list[-1]
if axis == 'x':
if f0 is not f1:
setf_e = cpu.SetFields(f1, ['ey', 'ez'], \
(f1.nx-1, 0, 0), (f1.nx-1, f1.ny-2, f1.nz-2), True)
getf_e = gpu.GetFields(f0, ['ey', 'ez'], \
(0, 0, 0), (0, f0.ny-2, f0.nz-2) )
setf_h = gpu.SetFields(f0, ['hy', 'hz'], \
(0, 1, 1), (0, f0.ny-1, f0.nz-1), True )
getf_h = cpu.GetFields(f1, ['hy', 'hz'], \
(f1.nx-1, 1, 1), (f1.nx-1, f1.ny-1, f1.nz-1) )
else:
getattr(self, f0.device_type).Pbc(f0, axis)
elif axis in ['y', 'z']:
for f in mainf_list:
getattr(self, f.device_type).Pbc(f, axis)
# global variables and functions
if axis == 'x' and f0 is not f1:
self.setf_e = setf_e
self.getf_e = getf_e
self.setf_h = setf_h
self.getf_h = getf_h
self.update_e = self.update_e_actual
self.update_h = self.update_h_actual
else:
self.update_e = lambda: None
self.update_h = lambda: None
def __init__(self, buffer_fields, max_tstep):
common.check_type('buffer_fields', buffer_fields, BufferFields)
# local variables
mainf = buffer_fields
nx, ny, nz = mainf.ns
dtype = mainf.dtype
direction = mainf.direction
target_rank = mainf.target_rank
# create instances (getf, setf and mpi requests)
if '+' in direction: # split h
getf = GetFields(mainf, ['hy', 'hz'], (1, 0, 0),
(1, ny - 1, nz - 1))
setf = SetFields(mainf, ['ey', 'ez'], (2, 0, 0),
(2, ny - 1, nz - 1), True)
if rank < target_rank:
tag_send, tag_recv = 0, 1
elif rank > target_rank: # pbc
tag_send, tag_recv = 2, 3
elif '-' in direction: # split e
getf = GetFields(mainf, ['ey', 'ez'], (1, 0, 0),
(1, ny - 1, nz - 1))
setf = SetFields(mainf, ['hy', 'hz'], (0, 0, 0),
(0, ny - 1, nz - 1), True)
if rank > target_rank:
tag_send, tag_recv = 1, 0
elif rank < target_rank: # pbc
tag_send, tag_recv = 3, 2
req_send = comm.Send_init(getf.host_array, target_rank, tag=tag_send)
tmp_recv_list = [
np.zeros(getf.host_array.shape, dtype) for i in range(2)
]
req_recv_list = [
comm.Recv_init(tmp_recv, target_rank, tag=tag_recv)
for tmp_recv in tmp_recv_list
]
# global variables and functions
self.getf = getf
self.setf = setf
self.req_send = req_send
self.req_recv_list = req_recv_list
self.tmp_recv_list = tmp_recv_list
self.switch = 0
self.max_tstep = max_tstep
self.tstep = 1
# append to the update list
self.priority_type = 'mpi'
mainf.append_instance(self)
def __init__(self, node_fields, axes):
"""
"""
common.check_type('node_fields', node_fields, Fields)
common.check_type('axes', axes, str)
assert len( set(axes).intersection(set('xyz')) ) > 0, 'axes option is wrong: %s is given' % repr(axes)
# local variables
nodef = node_fields
mainf_list = nodef.mainf_list
buffer_dict = nodef.buffer_dict
device_type_list = [f.device_type for f in nodef.updatef_list]
if 'gpu' in device_type_list:
from kemp.fdtd3d import gpu
self.gpu = gpu
if 'cpu' in device_type_list:
from kemp.fdtd3d import cpu
self.cpu = cpu
# create Pbc instances
if len(mainf_list) == 1:
f0 = mainf_list[0]
getattr(self, f0.device_type).Pbc(f0, axes)
elif len(mainf_list) > 1:
f0 = mainf_list[0]
f1 = mainf_list[-1]
if 'x' in axes:
nx, ny, nz = f1.ns
self.getf_e = getattr(self, f0.device_type).GetFields(f0, ['ey', 'ez'], (0, 0, 0), (0, -1, -1) )
self.setf_e = getattr(self, f1.device_type).SetFields(f1, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.getf_h = getattr(self, f1.device_type).GetFields(f1, ['hy', 'hz'], (-1, 0, 0), (-1, -1, -1) )
self.setf_h = getattr(self, f0.device_type).SetFields(f0, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True )
self.endf_is_cpu = True if f1.device_type == 'cpu' else False
# append to the update list
self.priority_type = 'pbc'
nodef.append_instance(self)
axs = axes.strip('x')
if len( set(axs).intersection(set('yz')) ) > 0:
for f in mainf_list:
getattr(self, f.device_type).Pbc(f, axs)
# for buffer fields
for direction, buf in buffer_dict.items():
axs = axes.strip( direction[0] )
if axs != '':
cpu.Pbc(buf, axs)
def set_chs(self, chx, chy, chz):
for ch in [chx, chy, chz]:
common.check_type('ch', ch, np.ndarray)
if self.ch_on:
self.chx[:] = chx[:]
self.chy[:] = chy[:]
self.chz[:] = chz[:]
else:
raise AttributeError("The Fields instance has no ch arrays.")
def set_ces(self, cex, cey, cez):
for ce in [cex, cey, cez]:
common.check_type('ce', ce, np.ndarray)
if self.ce_on:
self.cex[:] = cex[:]
self.cey[:] = cey[:]
self.cez[:] = cez[:]
else:
raise AttributeError("The Fields instance has no ce arrays.")
def __init__(self, node_fields):
"""
"""
common.check_type('node_fields', node_fields, NodeFields)
# local variables
nodef = node_fields
mainf_list = nodef.mainf_list
cpuf_dict = nodef.cpuf_dict
self.setf_dict = {'e': [], 'h': []}
self.getf_dict = {'e': [], 'h': []}
self.setf_block_dict = {'e': [], 'h': []}
self.getf_block_dict = {'e': [], 'h': []}
anx_list = nodef.accum_nx_list
# main and x+
for f0, f1 in zip(mainf_list[:-1], mainf_list[1:]):
self.append_setf_getf_pair('x', f0, f1)
# x-
if cpuf_dict.has_key('x-'):
cf = cpuf_dict['x-']
gf = mainf_list[0]
self.append_setf_getf_pair('x', cf, gf, True)
# y+
if cpuf_dict.has_key('y+'):
cf = cpuf_dict['y+']
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair('y', gf, cf, True, [anx0, anx1 - 1])
# y-
if cpuf_dict.has_key('y-'):
cf = cpuf_dict['y-']
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair('y', cf, gf, True, [anx0, anx1 - 1])
# z+
if cpuf_dict.has_key('z+'):
cf = cpuf_dict['z+']
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair('z', gf, cf, True, [anx0, anx1 - 1])
# z-
if cpuf_dict.has_key('z-'):
cf = cpuf_dict['z-']
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair('z', cf, gf, True, [anx0, anx1 - 1])
# global variables
self.mainf_list = mainf_list
self.cpuf_dict = cpuf_dict
def __init__(self, node_fields, axis):
"""
"""
common.check_type('node_fields', node_fields, NodeFields)
common.check_value('axis', axis, ['x', 'y', 'z'])
# local variables
self.gpu = gpu
self.cpu = cpu
mainf_list = node_fields.mainf_list
cpuf_dict = node_fields.cpuf_dict
axis_id = {'x':0, 'y':1, 'z':2}[axis]
set_cpuf = set( cpuf_dict.keys() )
for ax in ['x', 'y', 'z']:
if not set_cpuf.isdisjoint( [ax+'+', ax+'-'] ):
raise ValueError, 'There are %s-axis buffer instances. The pbc_internal operation along %s-axis is not allowed.' % (ax, ax)
# create pbc instances
f0 = mainf_list[0]
f1 = mainf_list[-1]
if axis == 'x':
if f0 is not f1:
setf_e = cpu.SetFields(f1, ['ey', 'ez'], \
(f1.nx-1, 0, 0), (f1.nx-1, f1.ny-2, f1.nz-2), True)
getf_e = gpu.GetFields(f0, ['ey', 'ez'], \
(0, 0, 0), (0, f0.ny-2, f0.nz-2) )
setf_h = gpu.SetFields(f0, ['hy', 'hz'], \
(0, 1, 1), (0, f0.ny-1, f0.nz-1), True )
getf_h = cpu.GetFields(f1, ['hy', 'hz'], \
(f1.nx-1, 1, 1), (f1.nx-1, f1.ny-1, f1.nz-1) )
else:
getattr(self, f0.device_type).Pbc(f0, axis)
elif axis in ['y', 'z']:
for f in mainf_list:
getattr(self, f.device_type).Pbc(f, axis)
# global variables and functions
if axis == 'x' and f0 is not f1:
self.setf_e = setf_e
self.getf_e = getf_e
self.setf_h = setf_h
self.getf_h = getf_h
self.update_e = self.update_e_actual
self.update_h = self.update_h_actual
else:
self.update_e = lambda : None
self.update_h = lambda : None
def __init__(self, node_fields):
"""
"""
common.check_type("node_fields", node_fields, NodeFields)
# local variables
nodef = node_fields
mainf_list = nodef.mainf_list
cpuf_dict = nodef.cpuf_dict
self.setf_dict = {"e": [], "h": []}
self.getf_dict = {"e": [], "h": []}
self.setf_block_dict = {"e": [], "h": []}
self.getf_block_dict = {"e": [], "h": []}
anx_list = nodef.accum_nx_list
# main and x+
for f0, f1 in zip(mainf_list[:-1], mainf_list[1:]):
self.append_setf_getf_pair("x", f0, f1)
# x-
if cpuf_dict.has_key("x-"):
cf = cpuf_dict["x-"]
gf = mainf_list[0]
self.append_setf_getf_pair("x", cf, gf, True)
# y+
if cpuf_dict.has_key("y+"):
cf = cpuf_dict["y+"]
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair("y", gf, cf, True, [anx0, anx1 - 1])
# y-
if cpuf_dict.has_key("y-"):
cf = cpuf_dict["y-"]
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair("y", cf, gf, True, [anx0, anx1 - 1])
# z+
if cpuf_dict.has_key("z+"):
cf = cpuf_dict["z+"]
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair("z", gf, cf, True, [anx0, anx1 - 1])
# z-
if cpuf_dict.has_key("z-"):
cf = cpuf_dict["z-"]
for gf, anx0, anx1 in zip(mainf_list, anx_list[:-1], anx_list[1:]):
self.append_setf_getf_pair("z", cf, gf, True, [anx0, anx1 - 1])
# global variables
self.mainf_list = mainf_list
self.cpuf_dict = cpuf_dict
def __init__(self, fields, axis):
"""
"""
common.check_type('fields', fields, Fields)
common.check_value('axis', axis, ['x', 'y', 'z'])
# local variables
nx, ny, nz = fields.ns
dtype_str_list = fields.dtype_str_list
# program
replace = lambda lst, idx, val: lst[:idx] + [val] + lst[idx+1:]
base0 = {'e': [0, 0, 0], 'h': [1, 1, 1]}
base1 = {'e': [nx-2, ny-2, nz-2], 'h': [nx-1, ny-1, nz-1]}
axis_id = {'x':0, 'y':1, 'z':2}[axis]
nn = fields.ns[axis_id]
value_dict = {'e': [], 'h': []}
for eh in ['e', 'h']:
for idx in {'e': [0, nn-1], 'h':[nn-1, 0]}[eh]:
pt0 = replace(base0[eh], axis_id, idx)
pt1 = replace(base1[eh], axis_id, idx)
nmax, xid, yid, zid = \
common_gpu.macro_replace_list(pt0, pt1)
value_dict[eh].append( \
'%s*ny*nz + %s*nz + %s' % (xid, yid, zid) )
macros = ['NMAX', 'IDX0', 'IDX1', 'DTYPE', 'PRAGMA_fp64']
values_e = [nmax] + value_dict['e'] + dtype_str_list
values_h = [nmax] + value_dict['h'] + dtype_str_list
ksrc_e = common.replace_template_code( \
open(common_gpu.src_path + 'copy_self.cl').read(), \
macros, values_e)
ksrc_h = common.replace_template_code( \
open(common_gpu.src_path + 'copy_self.cl').read(), \
macros, values_h)
program_e = cl.Program(fields.context, ksrc_e).build()
program_h = cl.Program(fields.context, ksrc_h).build()
# global variables
self.mainf = fields
self.program_e = program_e
self.program_h = program_h
self.strfs_e = {'x':['ey','ez'], 'y':['ex','ez'], 'z':['ex','ey']}[axis]
self.strfs_h = {'x':['hy','hz'], 'y':['hx','hz'], 'z':['hx','hy']}[axis]
# append to the update list
self.priority_type = 'pbc'
self.mainf.append_instance(self)
def __init__(self, node_fields, target_rank, direction):
common.check_type('node_fields', node_fields, Fields)
common.check_type('target_rank', target_rank, int)
# local variables
nodef = node_fields
dtype = nodef.dtype
assert rank != target_rank, 'The target_rank %d is same as the my_rank %d.' % (
target_rank, rank)
# create instances (getf, setf and mpi requests)
if '+' in direction:
mainf = nodef.mainf_list[-1]
nx, ny, nz = mainf.ns
getf = GetFields(mainf, ['hy', 'hz'], (nx - 1, 0, 0),
(nx - 1, ny - 1, nz - 1))
setf = SetFields(mainf, ['ey', 'ez'], (nx - 1, 0, 0),
(nx - 1, ny - 1, nz - 1), True)
if rank < target_rank:
tag_send, tag_recv = 0, 1
elif rank > target_rank:
tag_send, tag_recv = 2, 3
elif '-' in direction:
mainf = nodef.mainf_list[0]
nx, ny, nz = mainf.ns
getf = GetFields(mainf, ['ey', 'ez'], (0, 0, 0),
(0, ny - 1, nz - 1))
setf = SetFields(mainf, ['hy', 'hz'], (0, 0, 0),
(0, ny - 1, nz - 1), True)
if rank > target_rank:
tag_send, tag_recv = 1, 0
elif rank < target_rank: # pbc
tag_send, tag_recv = 3, 2
# global variables
self.target_rank = target_rank
self.getf = getf
self.setf = setf
self.tag_send = tag_send
self.tag_recv = tag_recv
self.tmp_recv = np.zeros(getf.host_array.shape, dtype)
# global functions
self.update_e = self.recv if '+' in direction else self.send
self.update_h = self.send if '+' in direction else self.recv
# append instance
self.priority_type = 'mpi'
nodef.append_instance(self)
def __init__(self, fields, axes):
"""
"""
common.check_type('fields', fields, Fields)
common.check_type('axes', axes, str)
assert len( set(axes).intersection(set('xyz')) ) > 0, 'axes option is wrong: %s is given' % repr(axes)
# local variables
nx, ny, nz = fields.ns
dtype_str_list = fields.dtype_str_list
# program
macros = ['NMAX', 'IDX0', 'IDX1', 'DTYPE', 'PRAGMA_fp64']
program_dict = {}
gs_dict = {}
for axis in list(axes):
program_dict[axis] = {}
for e_or_h in ['e', 'h']:
pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['get']
pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['get']
nmaxi_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1)
idx0_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str, zid_str)
pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['set']
pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['set']
nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1)
idx1_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str, zid_str)
values = [nmax_str, idx0_str, idx1_str] + dtype_str_list
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'copy_self.cl').read(), \
macros, values)
program = cl.Program(fields.context, ksrc).build()
program_dict[axis][e_or_h] = program
nmax = int(nmax_str)
remainder = nmax % fields.ls
gs_dict[axis] = nmax if remainder == 0 else nmax - remainder + fields.ls
# global variables
self.mainf = fields
self.axes = axes
self.program_dict = program_dict
self.gs_dict = gs_dict
# append to the update list
self.priority_type = 'pbc'
self.mainf.append_instance(self)
def __init__(self, buffer_fields, target_rank):
common.check_type('buffer_fields', buffer_fields, BufferFields)
common.check_type('target_rank', target_rank, int)
# local variables
mainf = buffer_fields
nx, ny, nz = mainf.ns
dtype = mainf.dtype
direction = mainf.direction
assert rank != target_rank, 'The target_rank %d is same as the my_rank %d.' % (
target_rank, rank)
# create instances (getf, setf and mpi requests)
if '+' in direction: # split h
getf = GetFields(mainf, ['hy', 'hz'], (1, 0, 0),
(1, ny - 1, nz - 1))
setf = SetFields(mainf, ['ey', 'ez'], (2, 0, 0),
(2, ny - 1, nz - 1), True)
if rank < target_rank:
tag_send, tag_recv = 0, 1
elif rank > target_rank: # pbc
tag_send, tag_recv = 2, 3
elif '-' in direction: # split e
getf = GetFields(mainf, ['ey', 'ez'], (1, 0, 0),
(1, ny - 1, nz - 1))
setf = SetFields(mainf, ['hy', 'hz'], (0, 0, 0),
(0, ny - 1, nz - 1), True)
if rank > target_rank:
tag_send, tag_recv = 1, 0
elif rank < target_rank: # pbc
tag_send, tag_recv = 3, 2
# global variables and functions
self.mainf = mainf
self.target_rank = target_rank
self.getf = getf
self.setf = setf
self.tag_send = tag_send
self.tag_recv = tag_recv
self.tmp_recv = np.zeros(getf.host_array.shape, dtype)
# global functions
self.update_e = self.recv if '+' in direction else self.send
self.update_h = self.send if '+' in direction else self.recv
# append to the update list
self.priority_type = 'mpi'
mainf.append_instance(self)
def __init__(self, fields):
"""
"""
common.check_type('fields', fields, Fields)
# local variables
dtype = fields.dtype
nx, ny, nz = ns = fields.ns
ce_on = fields.ce_on
ch_on = fields.ch_on
rd_on = fields.rd_on
# program
cf_values = {True: ['', 'OOOOOOOOO', ', &Cx, &Cy, &Cz', 'cx[idx]', 'cy[idx]', 'cz[idx]'], \
False: ['// ', 'OOOOOO', '', '0.5', '0.5', '0.5']}
ce_macros = [
'COEFF_E ', 'PARSE_ARGS_CE', 'PYARRAYOBJECT_CE', 'CEX', 'CEY',
'CEZ'
]
ce_values = cf_values[ce_on]
ch_macros = [
'COEFF_H ', 'PARSE_ARGS_CH', 'PYARRAYOBJECT_CH', 'CHX', 'CHY',
'CHZ'
]
ch_values = cf_values[ch_on]
macros = fields.dtype_omp_macros + ce_macros + ch_macros
values = fields.dtype_omp_values + ce_values + ch_values
ksrc = common.replace_template_code( \
open(common_cpu.src_path + 'core.c').read(), macros, values)
program = common_cpu.build_clib(ksrc, 'core')
# arguments
e_args = fields.ehs
h_args = fields.ehs
if ce_on:
e_args += fields.ces
if ch_on:
h_args += fields.chs
# global variables
self.mainf = fields
self.program = program
self.e_args = e_args
self.h_args = h_args
# append to the update list
self.priority_type = 'core'
fields.append_instance(self)
def __init__(self, fields, axes):
common.check_type('fields', fields, Fields)
common.check_type('axes', axes, str)
assert len(set(axes).intersection(set('xyz'))) > 0
# global variables
self.mainf = fields
self.axes = axes
# append to the update list
self.priority_type = 'pbc'
fields.append_instance(self)
def __init__(self, node_fields, axes):
"""
"""
common.check_type('node_fields', node_fields, Fields)
common.check_type('axes', axes, str)
assert len( set(axes).intersection(set('xyz')) ) > 0, 'axes option is wrong: %s is given' % repr(axes)
# local variables
nodef = node_fields
mainf_list = nodef.mainf_list
buffer_dict = nodef.buffer_dict
self.cpu = cpu
if 'gpu' in [f.device_type for f in nodef.updatef_list]:
from kemp.fdtd3d import gpu
self.gpu = gpu
# create pbc instances
f0 = mainf_list[-1]
f1 = mainf_list[0]
if f0 is f1:
getattr(self, f0.device_type).Pbc(f0, axes)
else:
if 'x' in axes:
nx, ny, nz = f0.ns
self.getf_e = getattr(self, f1.device_type).GetFields(f1, ['ey', 'ez'], (0, 0, 0), (0, ny-1, nz-1) )
self.setf_e = getattr(self, f0.device_type).SetFields(f0, ['ey', 'ez'], (nx-1, 0, 0), (nx-1, ny-1, nz-1), True)
self.getf_h = getattr(self, f0.device_type).GetFields(f0, ['hy', 'hz'], (nx-1, 0, 0), (nx-1, ny-1, nz-1) )
self.setf_h = getattr(self, f1.device_type).SetFields(f1, ['hy', 'hz'], (0, 0, 0), (0, ny-1, nz-1), True )
# append to the update list
self.priority_type = 'pbc'
nodef.append_instance(self)
axs = axes.strip('x')
if len( set(axs).intersection(set('yz')) ) > 0:
for f in mainf_list:
getattr(self, f.device_type).Pbc(f, axs)
# for buffer fields
for direction , buffer in buffer_dict.items():
replaced_axes = map(lambda ax: common_buffer.axes_dict[direction[0]][ax], list(axes))
axs = ''.join(replaced_axes).strip('x')
if axs != '':
cpu.Pbc(buffer, axs)
def __init__(self, fields, directions, npml=50, sigma_max=0.5, kappa_max=1, alpha_max=0, m_sigma=3, m_alpha=1):
common.check_type('fields', fields, Fields)
common.check_type('directions', directions, (list, tuple), str)
assert len(directions) == 3
for axis in directions:
assert axis in ['+', '-', '+-', '']
# local variables
dt = fields.dt
nx, ny, nz = fields.ns
dtype = fields.dtype
# allocations
psi_xs = [np.zeros((2*npml + 2, ny, nz), dtype) for i in range(4)]
psi_ys = [np.zeros((nx, 2*npml + 2, nz), dtype) for i in range(4)]
psi_zs = [np.zeros((nx, 2*npml + 2, nz), dtype) for i in range(4)]
i_e = np.arange(0.5, npml)
i_h = np.arange(1, npml+1)
sigma_e = sigma_max# * (i_e / npml) ** m_sigma
sigma_h = sigma_max# * (i_h / npml) ** m_sigma
print 'sigma_e', sigma_e
print 'sigma_h', sigma_h
kappa_e = 1 + (kappa_max - 1) * (i_e / npml) ** m_sigma
kappa_h = 1 + (kappa_max - 1) * (i_h / npml) ** m_sigma
alpha_e = alpha_max * ((npml - i_e) / npml) ** m_alpha
alpha_h = alpha_max * ((npml - i_h) / npml) ** m_alpha
com_e = (kappa_e * alpha_e + sigma_e) * dt + 2 * kappa_e
com_h = (kappa_h * alpha_h + sigma_h) * dt + 2 * kappa_h
pca_e = 4 * kappa_e / com_e - 1
pca_h = 4 * kappa_h / com_h - 1
pcb_e = (alpha_e * dt - 2 + 4 * kappa_e) / com_e - 1
pcb_h = (alpha_h * dt - 2 + 4 * kappa_h) / com_h - 1
pcc_e = (alpha_e * dt + 2) / com_e - 1
pcc_h = (alpha_h * dt + 2) / com_h - 1
# global variables
self.mainf = fields
self.directions = directions
self.npml = npml
self.psi_xs = psi_xs
self.psi_ys = psi_ys
self.psi_zs = psi_zs
self.pcs_e = [pca_e, pcb_e, pcc_e]
self.pcs_h = [pca_h, pcb_h, pcc_h]
# append to the update list
self.priority_type = 'pml'
fields.append_instance(self)
def macro_replace_list(pt0, pt1):
"""
Return the replace string list correspond to macro
This is used to generate the opencl kernel from the template.
"""
common.check_type('pt0', pt0, (list, tuple), int)
common.check_type('pt1', pt1, (list, tuple), int)
x0, y0, z0 = pt0
x1, y1, z1 = pt1
snx = abs(x1 - x0) + 1
sny = abs(y1 - y0) + 1
snz = abs(z1 - z0) + 1
nmax = snx * sny * snz
xid, yid, zid = x0, y0, z0
if x0 == x1 and y0 == y1 and z0 == z1:
pass
elif x0 != x1 and y0 == y1 and z0 == z1:
xid = '(gid + %d)' % x0
elif x0 == x1 and y0 != y1 and z0 == z1:
yid = '(gid + %d)' % y0
elif x0 == x1 and y0 == y1 and z0 != z1:
zid = '(gid + %d)' % z0
elif x0 != x1 and y0 != y1 and z0 == z1:
xid = '(gid/%d + %d)' % (sny, x0)
yid = '(gid%%%d + %d)' % (sny, y0)
elif x0 == x1 and y0 != y1 and z0 != z1:
yid = '(gid/%d + %d)' % (snz, y0)
zid = '(gid%%%d + %d)' % (snz, z0)
elif x0 != x1 and y0 == y1 and z0 != z1:
xid = '(gid/%d + %d)' % (snz, x0)
zid = '(gid%%%d + %d)' % (snz, z0)
elif x0 != x1 and y0 != y1 and z0 != z1:
xid = '(gid/%d + %d)' % (sny*snz, x0)
yid = '((gid/%d)%%%d + %d)' % (snz, sny, y0)
zid = '(gid%%%d + %d)' % (snz, z0)
return [str(nmax), str(xid), str(yid), str(zid)]
def __init__(self, fields):
common.check_type('fields', fields, Fields)
# global variables
self.mainf = fields
pad = fields.pad
slice_z0 = slice(1, None) if pad == 0 else slice(1, -pad)
slice_z1 = slice(None, -pad - 1)
self.slice_z_list = [fields.slice_z, slice_z0, slice_z1]
# append to the update list
self.priority_type = 'core'
fields.append_instance(self)
def __init__(self, fields):
common.check_type('fields', fields, Fields)
# global variables
self.mainf = fields
pad = fields.pad
slice_z0 = slice(1, None) if pad == 0 else slice(1, -pad)
slice_z1 = slice(None, -pad-1)
self.slice_z_list = [fields.slice_z, slice_z0, slice_z1]
# append to the update list
self.priority_type = 'core'
fields.append_instance(self)
def __init__(self, fields, str_f, pt0, pt1):
"""
"""
common.check_type("fields", fields, Fields)
common.check_type("str_f", str_f, (str, list, tuple), str)
common.check_type("pt0", pt0, (list, tuple), int)
common.check_type("pt1", pt1, (list, tuple), int)
# local variables
str_fs = common.convert_to_tuple(str_f)
for strf in str_fs:
strf_list = ["ex", "ey", "ez", "hx", "hy", "hz"]
common.check_value("str_f", strf, strf_list)
for axis, n, p0, p1 in zip(["x", "y", "z"], fields.ns, pt0, pt1):
common.check_value("pt0 %s" % axis, p0, range(n))
common.check_value("pt1 %s" % axis, p1, range(n))
# allocation
shape = common.shape_two_points(pt0, pt1, len(str_fs))
host_array = np.zeros(shape, dtype=fields.dtype)
split_host_array = np.split(host_array, len(str_fs))
split_host_array_dict = dict(zip(str_fs, split_host_array))
# global variables
self.mainf = fields
self.str_fs = str_fs
self.slice_xyz = common.slice_index_two_points(pt0, pt1)
self.host_array = host_array
self.split_host_array_dict = split_host_array_dict
def __init__(self, fields, str_f, pt0, pt1):
"""
"""
common.check_type('fields', fields, (Fields, BufferFields))
common.check_type('str_f', str_f, (str, list, tuple), str)
common.check_type('pt0', pt0, (list, tuple), int)
common.check_type('pt1', pt1, (list, tuple), int)
# local variables
str_fs = common.convert_to_tuple(str_f)
for strf in str_fs:
strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz']
common.check_value('str_f', strf, strf_list)
for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1):
common.check_value('pt0 %s' % axis, p0, range(n))
common.check_value('pt1 %s' % axis, p1, range(n))
# allocation
shape = common.shape_two_points(pt0, pt1, len(str_fs))
host_array = np.zeros(shape, dtype=fields.dtype)
split_host_array = np.split(host_array, len(str_fs))
split_host_array_dict = dict( zip(str_fs, split_host_array) )
# global variables
self.mainf = fields
self.str_fs = str_fs
self.slice_xyz = common.slices_two_points(pt0, pt1)
self.host_array = host_array
self.split_host_array_dict = split_host_array_dict
def __init__(self, node_fields):
common.check_type('node_fields', node_fields, Fields)
# create Core instances
f_list = node_fields.updatef_list
self.cpu = cpu
if 'gpu' in [f.device_type for f in f_list]:
from kemp.fdtd3d import gpu
self.gpu = gpu
for f in f_list:
getattr(self, f.device_type).Core(f)
def __init__(self, node_fields, target_rank, direction):
common.check_type('node_fields', node_fields, Fields)
common.check_type('target_rank', target_rank, int)
# local variables
nodef = node_fields
dtype = nodef.dtype
assert rank != target_rank, 'The target_rank %d is same as the my_rank %d.' % (target_rank, rank)
# create instances (getf, setf and mpi requests)
if '+' in direction:
mainf = nodef.mainf_list[-1]
nx, ny, nz = mainf.ns
getf = GetFields(mainf, ['hy', 'hz'], (nx-1, 0, 0), (nx-1, ny-1, nz-1))
setf = SetFields(mainf, ['ey', 'ez'], (nx-1, 0, 0), (nx-1, ny-1, nz-1), True)
if rank < target_rank:
tag_send, tag_recv = 0, 1
elif rank > target_rank:
tag_send, tag_recv = 2, 3
elif '-' in direction:
mainf = nodef.mainf_list[0]
nx, ny, nz = mainf.ns
getf = GetFields(mainf, ['ey', 'ez'], (0, 0, 0), (0, ny-1, nz-1))
setf = SetFields(mainf, ['hy', 'hz'], (0, 0, 0), (0, ny-1, nz-1), True)
if rank > target_rank:
tag_send, tag_recv = 1, 0
elif rank < target_rank: # pbc
tag_send, tag_recv = 3, 2
# global variables
self.target_rank = target_rank
self.getf = getf
self.setf = setf
self.tag_send = tag_send
self.tag_recv = tag_recv
self.tmp_recv = np.zeros(getf.host_array.shape, dtype)
# global functions
self.update_e = self.recv if '+' in direction else self.send
self.update_h = self.send if '+' in direction else self.recv
# append instance
self.priority_type = 'mpi'
nodef.append_instance(self)
def macro_replace_list(pt0, pt1):
"""
Return the replace string list correspond to macro
This is used to generate the opencl kernel from the template.
"""
common.check_type('pt0', pt0, (list, tuple), int)
common.check_type('pt1', pt1, (list, tuple), int)
x0, y0, z0 = pt0
x1, y1, z1 = pt1
snx = abs(x1 - x0) + 1
sny = abs(y1 - y0) + 1
snz = abs(z1 - z0) + 1
nmax = snx * sny * snz
xid, yid, zid = x0, y0, z0
if x0 == x1 and y0 == y1 and z0 == z1:
pass
elif x0 != x1 and y0 == y1 and z0 == z1:
xid = '(gid + %d)' % x0
elif x0 == x1 and y0 != y1 and z0 == z1:
yid = '(gid + %d)' % y0
elif x0 == x1 and y0 == y1 and z0 != z1:
zid = '(gid + %d)' % z0
elif x0 != x1 and y0 != y1 and z0 == z1:
xid = '(gid/%d + %d)' % (sny, x0)
yid = '(gid%%%d + %d)' % (sny, y0)
elif x0 == x1 and y0 != y1 and z0 != z1:
yid = '(gid/%d + %d)' % (snz, y0)
zid = '(gid%%%d + %d)' % (snz, z0)
elif x0 != x1 and y0 == y1 and z0 != z1:
xid = '(gid/%d + %d)' % (snz, x0)
zid = '(gid%%%d + %d)' % (snz, z0)
elif x0 != x1 and y0 != y1 and z0 != z1:
xid = '(gid/%d + %d)' % (sny * snz, x0)
yid = '((gid/%d)%%%d + %d)' % (snz, sny, y0)
zid = '(gid%%%d + %d)' % (snz, z0)
return [str(nmax), str(xid), str(yid), str(zid)]
def __init__(self, fields, str_f, pt0, pt1):
"""
"""
common.check_type('fields', fields, (Fields, BufferFields))
common.check_type('str_f', str_f, (str, list, tuple), str)
common.check_type('pt0', pt0, (list, tuple), int)
common.check_type('pt1', pt1, (list, tuple), int)
# local variables
str_fs = common.convert_to_tuple(str_f)
for strf in str_fs:
strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz']
common.check_value('str_f', strf, strf_list)
for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1):
common.check_value('pt0 %s' % axis, p0, range(n))
common.check_value('pt1 %s' % axis, p1, range(n))
# allocation
shape = common.shape_two_points(pt0, pt1, len(str_fs))
host_array = np.zeros(shape, dtype=fields.dtype)
split_host_array = np.split(host_array, len(str_fs))
split_host_array_dict = dict(zip(str_fs, split_host_array))
# global variables
self.mainf = fields
self.str_fs = str_fs
self.slice_xyz = common.slices_two_points(pt0, pt1)
self.host_array = host_array
self.split_host_array_dict = split_host_array_dict
def __init__(self, buffer_fields, target_rank, max_tstep):
common.check_type('buffer_fields', buffer_fields, BufferFields)
common.check_type('target_rank', target_rank, int)
# local variables
mainf = buffer_fields
nx, ny, nz = mainf.ns
dtype = mainf.dtype
direction = mainf.direction
assert rank != target_rank, 'The target_rank %d is same as the my_rank %d.' % (target_rank, rank)
# create instances (getf, setf and mpi requests)
if '+' in direction: # split h
getf = GetFields(mainf, ['hy', 'hz'], (1, 0, 0), (1, ny-1, nz-1))
setf = SetFields(mainf, ['ey', 'ez'], (2, 0, 0), (2, ny-1, nz-1), True)
if rank < target_rank:
tag_send, tag_recv = 0, 1
elif rank > target_rank: # pbc
tag_send, tag_recv = 2, 3
elif '-' in direction: # split e
getf = GetFields(mainf, ['ey', 'ez'], (1, 0, 0), (1, ny-1, nz-1))
setf = SetFields(mainf, ['hy', 'hz'], (0, 0, 0), (0, ny-1, nz-1), True)
if rank > target_rank:
tag_send, tag_recv = 1, 0
elif rank < target_rank: # pbc
tag_send, tag_recv = 3, 2
req_send = comm.Send_init(getf.host_array, target_rank, tag=tag_send)
tmp_recv = np.zeros(getf.host_array.shape, dtype)
req_recv = comm.Recv_init(tmp_recv, target_rank, tag=tag_recv)
# global variables and functions
self.mainf = mainf
self.getf = getf
self.setf = setf
self.req_send = req_send
self.tmp_recv = tmp_recv
self.req_recv = req_recv
self.max_tstep = max_tstep
self.tstep = 1
# append to the update list
self.priority_type = 'mpi'
mainf.append_instance(self)
def __init__(self, fields, axes):
"""
"""
common.check_type('fields', fields, (Fields, BufferFields))
common.check_type('axes', axes, str)
assert len( set(axes).intersection(set('xyz')) ) > 0, 'axes option is wrong: %s is given' % repr(axes)
# global variables
self.mainf = fields
self.axes = axes
# append to the update list
self.priority_type = 'pbc'
self.mainf.append_instance(self)
def __init__(self, node_fields):
common.check_type('node_fields', node_fields, NodeFields)
# local variables
nodef = node_fields
updatef_list = nodef.updatef_list
# create Core instances
self.gpu, self.cpu = gpu, cpu
core_list = []
for f in updatef_list:
core_list.append( getattr(self, f.device_type).Core(f) )
# global variables
self.core_list = core_list
def __init__(self, buffer_fields, max_tstep):
common.check_type('buffer_fields', buffer_fields, BufferFields)
# local variables
mainf = buffer_fields
nx, ny, nz = mainf.ns
dtype = mainf.dtype
direction = mainf.direction
target_rank = mainf.target_rank
# create instances (getf, setf and mpi requests)
if '+' in direction: # split h
getf = GetFields(mainf, ['hy', 'hz'], (1, 0, 0), (1, ny-1, nz-1))
setf = SetFields(mainf, ['ey', 'ez'], (2, 0, 0), (2, ny-1, nz-1), True)
if rank < target_rank:
tag_send, tag_recv = 0, 1
elif rank > target_rank: # pbc
tag_send, tag_recv = 2, 3
elif '-' in direction: # split e
getf = GetFields(mainf, ['ey', 'ez'], (1, 0, 0), (1, ny-1, nz-1))
setf = SetFields(mainf, ['hy', 'hz'], (0, 0, 0), (0, ny-1, nz-1), True)
if rank > target_rank:
tag_send, tag_recv = 1, 0
elif rank < target_rank: # pbc
tag_send, tag_recv = 3, 2
req_send = comm.Send_init(getf.host_array, target_rank, tag=tag_send)
tmp_recv_list = [np.zeros(getf.host_array.shape, dtype) for i in range(2)]
req_recv_list = [comm.Recv_init(tmp_recv, target_rank, tag=tag_recv) for tmp_recv in tmp_recv_list]
# global variables and functions
self.getf = getf
self.setf = setf
self.req_send = req_send
self.req_recv_list = req_recv_list
self.tmp_recv_list = tmp_recv_list
self.switch = 0
self.max_tstep = max_tstep
self.tstep = 1
# append to the update list
self.priority_type = 'mpi'
mainf.append_instance(self)
def __init__(self, nodef):
common.check_type('nodef', nodef, node.Fields)
self.directions = nodef.buffer_dict.keys()
self.gpu = gpu
self.cpu = cpu
if 'x+' in self.directions:
mf_xp = nodef.mainf_list[-1]
bf_xp = nodef.buffer_dict['x+']
mpu = getattr(self, mf_xp.device_type)
self.gf_p_h = mpu.GetFields(mf_xp, ['hy', 'hz'], (-2, 0, 0), (-2, -1, -1))
self.sf_p_h = cpu.SetFields(bf_xp, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
self.sf_p_e = mpu.SetFields(mf_xp, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.gf_h = cpu.GetFields(bf_xp, ['hy', 'hz'], (1, 0, 0), (1, -1, -1))
self.sf_e = cpu.SetFields(bf_xp, ['ey', 'ez'], (1, 0, 0), (1, -1, -1), True)
target = 0 if rank == size-1 else rank+1
self.req_send_h = comm.Send_init(self.gf_h.host_array, target, tag=0)
self.tmp_recv_e = np.zeros(self.gf_h.host_array.shape, nodef.dtype)
self.req_recv_e = comm.Recv_init(self.tmp_recv_e, target, tag=1)
if 'x-' in self.directions:
mf_xm = nodef.mainf_list[0]
bf_xm = nodef.buffer_dict['x-']
mpu = getattr(self, mf_xm.device_type)
self.gf_m_e = mpu.GetFields(mf_xm, ['ey', 'ez'], (1, 0, 0), (1, -1, -1))
self.sf_m_e = cpu.SetFields(bf_xm, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.sf_m_h = mpu.SetFields(mf_xm, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
self.gf_e = cpu.GetFields(bf_xm, ['ey', 'ez'], (0, 0, 0), (0, -1, -1))
self.sf_h = cpu.SetFields(bf_xm, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
target = size-1 if rank == 0 else rank-1
self.req_send_e = comm.Send_init(self.gf_e.host_array, target, tag=1)
self.tmp_recv_h = np.zeros(self.gf_e.host_array.shape, nodef.dtype)
self.req_recv_h = comm.Recv_init(self.tmp_recv_h, target, tag=0)
# append to the update list
self.priority_type = 'mpi'
nodef.append_instance(self)
def __init__(self, fields, str_f, pt0, pt1):
"""
"""
common.check_type('fields', fields, Fields)
common.check_type('str_f', str_f, (str, list, tuple), str)
common.check_type('pt0', pt0, (list, tuple), (int, float))
common.check_type('pt1', pt1, (list, tuple), (int, float))
pt0 = list( common.convert_indices(fields.ns, pt0) )
pt1 = list( common.convert_indices(fields.ns, pt1) )
# local variables
str_fs = common.convert_to_tuple(str_f)
dtype_str_list = fields.dtype_str_list
for strf in str_fs:
strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz']
common.check_value('str_f', strf, strf_list)
for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1):
common.check_value('pt0 %s' % axis, p0, range(n))
common.check_value('pt1 %s' % axis, p1, range(n))
# program
macros = ['NMAX', 'XID', 'YID', 'ZID', \
'ARGS', \
'TARGET', 'SOURCE', 'OVERWRITE', \
'DTYPE']
values = common_gpu.macro_replace_list(pt0, pt1) + \
['DTYPE *source', \
'target[sub_idx]', 'source[idx]', '='] + \
dtype_str_list
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'copy.cu').read(), macros, values)
program = SourceModule(ksrc)
kernel_copy = program.get_function('copy')
# allocation
source_bufs = [fields.get_buf(str_f) for str_f in str_fs]
shape = common.shape_two_points(pt0, pt1, len(str_fs))
host_array = np.zeros(shape, fields.dtype)
split_host_array = np.split(host_array, len(str_fs))
split_host_array_dict = dict( zip(str_fs, split_host_array) )
target_buf = cuda.to_device(host_array)
# global variables
self.mainf = fields
self.kernel_copy = kernel_copy
self.source_bufs = source_bufs
self.target_buf = target_buf
self.host_array = host_array
self.split_host_array_dict = split_host_array_dict
def __init__(self, fields, target_rank, tmax):
common.check_type('fields', fields, Fields)
# local variables
nx, ny, nz = fields.ns
dtype = fields.dtype
mpi_type = fields.mpi_type
common.check_value('mpi_type', mpi_type, \
['x+', 'x-', 'y+', 'y-', 'z+', 'z-'])
# create instances (getf, setf and mpi requests)
if '+' in mpi_type: # split h
getf = GetFields(fields, ['hy', 'hz'], \
(1, 1, 1), (1, ny-1, nz-1))
setf = SetFields(fields, ['ey', 'ez'], \
(nx-1, 0, 0), (nx-1, ny-2, nz-2), True)
req_send = comm.Send_init(getf.host_array, target_rank, tag=1)
tmp_recv = np.zeros(getf.host_array.shape, dtype)
req_recv = comm.Recv_init(tmp_recv, target_rank, tag=2)
elif '-' in mpi_type: # split e
getf = GetFields(fields, ['ey', 'ez'], \
(nx-2, 0, 0), (nx-2, ny-2, nz-2))
setf = SetFields(fields, ['hy', 'hz'], \
(0, 1, 1), (0, ny-1, nz-1), True)
req_send = comm.Send_init(getf.host_array, target_rank, tag=2)
tmp_recv = np.zeros(getf.host_array.shape, dtype)
req_recv = comm.Recv_init(tmp_recv, target_rank, tag=1)
# global variables and functions
self.mainf = fields
self.getf = getf
self.setf = setf
self.tmp_recv = tmp_recv
self.req_send = req_send
self.req_recv = req_recv
self.tmax = tmax
self.tstep = 1
# append to the update list
self.priority_type = 'mpi'
self.mainf.append_instance(self)
def __init__(self, fields, str_f, pt0, pt1, source_is_array=False):
"""
"""
common.check_type('fields', fields, Fields)
common.check_type('str_f', str_f, (str, list, tuple), str)
common.check_type('pt0', pt0, (list, tuple), int)
common.check_type('pt1', pt1, (list, tuple), int)
common.check_type('source_is_array', source_is_array, bool)
self.mainf = mainf = fields
str_fs = common.convert_to_tuple(str_f)
# program
macros = ['NMAX', 'XID', 'YID', 'ZID', \
'ARGS', \
'TARGET', 'SOURCE', \
'DTYPE', 'PRAGMA_fp64']
if source_is_array:
values = macro_replace_list(pt0, pt1) + \
['__global DTYPE *source', \
'target[idx]', 'source[sub_idx]'] + mainf.dtype_str_list[:2]
else:
values = macro_replace_list(pt0, pt1) + \
['DTYPE source', \
'target[idx]', 'source'] + mainf.dtype_str_list[:2]
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'subdomain.cl').read(), macros, values)
self.program = cl.Program(mainf.context, ksrc).build()
# allocation
self.target_bufs = [mainf.get_buf(str_f) for str_f in str_fs]
shape = list( common.shape_two_points(pt0, pt1) )
shape[0] *= len(str_fs)
if source_is_array:
tmp_array = np.zeros(shape, dtype=mainf.dtype)
self.source_buf = cl.Buffer( \
mainf.context, \
cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, \
hostbuf=tmp_array)
self.set_fields = self.set_fields_spatial_value
else:
self.set_fields = self.set_fields_single_value
def __init__(self, fields, axis):
"""
"""
common.check_type('fields', fields, Fields)
common.check_value('axis', axis, ['x', 'y', 'z'])
mtype = fields.mpi_type
if axis == 'x' and mtype in ['x+', 'x-', 'y+', 'y-', 'z+', 'z-']:
raise ValueError, 'The fields.mpi_type is \'%s\'. The buffer instance is only permit the pbc operation along y and z axes' % mtype
# local variables
nx, ny, nz = fields.ns
axis_id = {'x':0, 'y':1, 'z':2}[axis]
# slice indices
replace = lambda lst, idx, val: lst[:idx] + [val] + lst[idx+1:]
slices_e = [slice(None, -1), slice(None, -1), slice(None, nz-1)]
slices_h = [slice(1, None), slice(1, None), slice(1, nz)]
slices_e_src = replace(slices_e, axis_id, slice(None, 1))
slices_h_dest = replace(slices_h, axis_id, slice(None, 1))
if axis == 'z':
slices_e_dest = replace(slices_e, axis_id, slice(nz-1, nz))
slices_h_src = replace(slices_h, axis_id, slice(nz-1, nz))
else:
slices_e_dest = replace(slices_e, axis_id, slice(-1, None))
slices_h_src = replace(slices_h, axis_id, slice(-1, None))
# global variables
self.mainf = fields
self.slices_dict = { \
'e_src': fields.split_slices_dict('e', slices_e_src), \
'e_dest': fields.split_slices_dict('e', slices_e_dest), \
'h_src': fields.split_slices_dict('h', slices_h_src), \
'h_dest': fields.split_slices_dict('h', slices_h_dest) }
self.strfs = {\
'x': {'e': ['ey','ez'], 'h': ['hy','hz']}, \
'y': {'e': ['ex','ez'], 'h': ['hx','hz']}, \
'z': {'e': ['ex','ey'], 'h': ['hx','hy']} }[axis]
# append to the update list
self.priority_type = 'pbc'
self.mainf.append_instance(self)
def __init__(self, gpuf, direction, tmax):
common.check_type('gpuf', gpuf, gpu.Fields)
common.check_value('direction', direction, ('+', '-', '+-'))
qtask = cpu.QueueTask()
if '+' in direction:
self.cpuf_p = cpuf_p = cpu.Fields(qtask, 3, gpuf.ny, gpuf.nz, gpuf.coeff_use, gpuf.precision_float, use_cpu_core=1)
self.gf_p_h = gpu.GetFields(gpuf, ['hy', 'hz'], (-2, 0, 0), (-2, -1, -1))
self.sf_p_h = cpu.SetFields(cpuf_p, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
self.gf_p_e = cpu.GetFields(cpuf_p, ['ey', 'ez'], (1, 0, 0), (1, -1, -1))
self.sf_p_e = gpu.SetFields(gpuf, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.gf_h = gf_h = cpu.GetFields(cpuf_p, ['hy', 'hz'], (1, 0, 0), (1, -1, -1))
self.sf_e = cpu.SetFields(cpuf_p, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.req_send_h = comm.Send_init(gf_h.host_array, rank+1, tag=0)
self.tmp_recv_e_list = [np.zeros(gf_h.host_array.shape, gpuf.dtype) for i in range(2)]
self.req_recv_e_list = [comm.Recv_init(tmp_recv_e, rank+1, tag=1) for tmp_recv_e in self.tmp_recv_e_list]
self.switch_e = 0
if '-' in direction:
self.cpuf_m = cpuf_m = cpu.Fields(qtask, 3, gpuf.ny, gpuf.nz, gpuf.coeff_use, gpuf.precision_float, use_cpu_core=1)
self.gf_m_e = gpu.GetFields(gpuf, ['ey', 'ez'], (1, 0, 0), (1, -1, -1))
self.sf_m_e = cpu.SetFields(cpuf_m, ['ey', 'ez'], (-1, 0, 0), (-1, -1, -1), True)
self.gf_m_h = cpu.GetFields(cpuf_m, ['hy', 'hz'], (1, 0, 0), (1, -1, -1))
self.sf_m_h = gpu.SetFields(gpuf, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
self.gf_e = gf_e = cpu.GetFields(cpuf_m, ['ey', 'ez'], (1, 0, 0), (1, -1, -1))
self.sf_h = cpu.SetFields(cpuf_m, ['hy', 'hz'], (0, 0, 0), (0, -1, -1), True)
self.req_send_e = comm.Send_init(gf_e.host_array, rank-1, tag=1)
self.tmp_recv_h_list = [np.zeros(gf_e.host_array.shape, gpuf.dtype) for i in range(2)]
self.req_recv_h_list = [comm.Recv_init(tmp_recv_h, rank-1, tag=0) for tmp_recv_h in self.tmp_recv_h_list]
self.switch_h = 0
# global variables
self.direction = direction
self.qtask = qtask
self.tmax = tmax
self.tstep = 1
def __init__(self, fields, axes):
"""
"""
common.check_type("fields", fields, Fields)
common.check_type("axes", axes, str)
assert len(set(axes).intersection(set("xyz"))) > 0, "axes option is wrong: %s is given" % repr(axes)
# local variables
nx, ny, nz = fields.ns
dtype_str_list = fields.dtype_str_list
# program
macros = ["NMAX", "IDX0", "IDX1", "DTYPE", "PRAGMA_fp64"]
program_dict = {}
for axis in list(axes):
program_dict[axis] = {}
for e_or_h in ["e", "h"]:
pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]["get"]
pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]["get"]
nmax, xid, yid, zid = common_gpu.macro_replace_list(pt0, pt1)
idx0_str = "%s*ny*nz + %s*nz + %s" % (xid, yid, zid)
pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]["set"]
pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]["set"]
nmax, xid, yid, zid = common_gpu.macro_replace_list(pt0, pt1)
idx1_str = "%s*ny*nz + %s*nz + %s" % (xid, yid, zid)
values = [nmax, idx0_str, idx1_str] + dtype_str_list
ksrc = common.replace_template_code(open(common_gpu.src_path + "copy_self.cl").read(), macros, values)
program = cl.Program(fields.context, ksrc).build()
program_dict[axis][e_or_h] = program
# global variables
self.mainf = fields
self.axes = axes
self.program_dict = program_dict
# append to the update list
self.priority_type = "pbc"
self.mainf.append_instance(self)
