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, 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, 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)
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, 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, 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)
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', 'PRAGMA_fp64']
values = common_gpu.macro_replace_list(pt0, pt1) + \
['__global DTYPE *source', \
'target[sub_idx]', 'source[idx]', '='] + \
dtype_str_list
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'copy.cl').read(), macros, values)
program = cl.Program(fields.context, ksrc).build()
# 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, dtype=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 = cl.Buffer( \
fields.context, \
cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, \
hostbuf=host_array)
# global variables
self.mainf = fields
self.program = program
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, 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, 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)
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", "PRAGMA_fp64"]
values = (
common_gpu.macro_replace_list(pt0, pt1)
+ ["__global DTYPE *source", "target[sub_idx]", "source[idx]", "="]
+ dtype_str_list
)
ksrc = common.replace_template_code(open(common_gpu.src_path + "copy.cl").read(), macros, values)
program = cl.Program(fields.context, ksrc).build()
# 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, dtype=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 = cl.Buffer(fields.context, cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, hostbuf=host_array)
# global variables
self.mainf = fields
self.program = program
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,
str_f,
pt0,
pt1,
is_array=False,
is_overwrite=True):
"""
"""
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))
common.check_type('is_array', is_array, bool)
common.check_type('is_overwrite', is_overwrite, bool)
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
overwrite_str = {True: '=', False: '+='}[is_overwrite]
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']
if is_array:
values = common_gpu.macro_replace_list(pt0, pt1) + \
['DTYPE *source', \
'target[idx]', 'source[sub_idx]', overwrite_str] + \
dtype_str_list
else:
values = common_gpu.macro_replace_list(pt0, pt1) + \
['DTYPE source', \
'target[idx]', 'source', overwrite_str] + \
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
target_bufs = [fields.get_buf(str_f) for str_f in str_fs]
shape = common.shape_two_points(pt0, pt1, len(str_fs))
if is_array:
tmp_array = np.zeros(shape, fields.dtype)
source_buf = cuda.to_device(tmp_array)
# global variabels and functions
self.mainf = fields
self.kernel_copy = kernel_copy
self.target_bufs = target_bufs
self.shape = shape
if is_array:
self.source_buf = source_buf
self.set_fields = self.set_fields_spatial_value
else:
self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True):
"""
"""
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))
common.check_type('is_array', is_array, bool)
common.check_type('is_overwrite', is_overwrite, bool)
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
overwrite_str = {True: '=', False: '+='}[is_overwrite]
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']
if is_array:
values = common_gpu.macro_replace_list(pt0, pt1) + \
['DTYPE *source', \
'target[idx]', 'source[sub_idx]', overwrite_str] + \
dtype_str_list
else:
values = common_gpu.macro_replace_list(pt0, pt1) + \
['DTYPE source', \
'target[idx]', 'source', overwrite_str] + \
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
target_bufs = [fields.get_buf(str_f) for str_f in str_fs]
shape = common.shape_two_points(pt0, pt1, len(str_fs))
if is_array:
tmp_array = np.zeros(shape, fields.dtype)
source_buf = cuda.to_device(tmp_array)
# global variabels and functions
self.mainf = fields
self.kernel_copy = kernel_copy
self.target_bufs = target_bufs
self.shape = shape
if is_array:
self.source_buf = source_buf
self.set_fields = self.set_fields_spatial_value
else:
self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True):
"""
"""
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))
common.check_type('is_array', is_array, bool)
common.check_type('is_overwrite', is_overwrite, bool)
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
overwrite_str = {True: '=', False: '+='}[is_overwrite]
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', 'PRAGMA_fp64']
nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1)
if is_array:
values = [nmax_str, xid_str, yid_str, zid_str, \
'__global DTYPE *source', \
'target[idx]', 'source[sub_idx]', overwrite_str] + \
dtype_str_list
else:
values = [nmax_str, xid_str, yid_str, zid_str, \
'DTYPE source', \
'target[idx]', 'source', overwrite_str] + \
dtype_str_list
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'copy.cl').read(), macros, values)
program = cl.Program(fields.context, ksrc).build()
# allocation
target_bufs = [fields.get_buf(str_f) for str_f in str_fs]
shape = common.shape_two_points(pt0, pt1, len(str_fs))
if is_array:
tmp_array = np.zeros(shape, dtype=fields.dtype)
source_buf = cl.Buffer( \
fields.context, \
cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, \
hostbuf=tmp_array)
# global variabels and functions
self.mainf = fields
self.program = program
self.target_bufs = target_bufs
self.shape = shape
nmax = int(nmax_str)
remainder = nmax % fields.ls
self.gs = nmax if remainder == 0 else nmax - remainder + fields.ls
if is_array:
self.source_buf = source_buf
self.set_fields = self.set_fields_spatial_value
else:
self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True):
"""
"""
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))
common.check_type('is_array', is_array, bool)
common.check_type('is_overwrite', is_overwrite, bool)
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
overwrite_str = {True: '=', False: '+='}[is_overwrite]
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', 'PRAGMA_fp64']
nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1)
if is_array:
values = [nmax_str, xid_str, yid_str, zid_str, \
'__global DTYPE *source', \
'target[idx]', 'source[sub_idx]', overwrite_str] + \
dtype_str_list
else:
values = [nmax_str, xid_str, yid_str, zid_str, \
'DTYPE source', \
'target[idx]', 'source', overwrite_str] + \
dtype_str_list
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'copy.cl').read(), macros, values)
program = cl.Program(fields.context, ksrc).build()
# allocation
target_bufs = [fields.get_buf(str_f) for str_f in str_fs]
shape = common.shape_two_points(pt0, pt1, len(str_fs))
if is_array:
tmp_array = np.zeros(shape, dtype=fields.dtype)
source_buf = cl.Buffer( \
fields.context, \
cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, \
hostbuf=tmp_array)
# global variabels and functions
self.mainf = fields
self.program = program
self.target_bufs = target_bufs
self.shape = shape
nmax = int(nmax_str)
remainder = nmax % fields.ls
self.gs = nmax if remainder == 0 else nmax - remainder + fields.ls
if is_array:
self.source_buf = source_buf
self.set_fields = self.set_fields_spatial_value
else:
self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True):
"""
"""
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("is_array", is_array, bool)
common.check_type("is_overwrite", is_overwrite, bool)
# local variables
str_fs = common.convert_to_tuple(str_f)
dtype_str_list = fields.dtype_str_list
overwrite_str = {True: "=", False: "+="}[is_overwrite]
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", "PRAGMA_fp64"]
if is_array:
values = (
common_gpu.macro_replace_list(pt0, pt1)
+ ["__global DTYPE *source", "target[idx]", "source[sub_idx]", overwrite_str]
+ dtype_str_list
)
else:
values = (
common_gpu.macro_replace_list(pt0, pt1)
+ ["DTYPE source", "target[idx]", "source", overwrite_str]
+ dtype_str_list
)
ksrc = common.replace_template_code(open(common_gpu.src_path + "copy.cl").read(), macros, values)
program = cl.Program(fields.context, ksrc).build()
# allocation
target_bufs = [fields.get_buf(str_f) for str_f in str_fs]
shape = common.shape_two_points(pt0, pt1, len(str_fs))
if is_array:
tmp_array = np.zeros(shape, dtype=fields.dtype)
source_buf = cl.Buffer(
fields.context, cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, hostbuf=tmp_array
)
# global variabels and functions
self.mainf = fields
self.program = program
self.target_bufs = target_bufs
self.shape = shape
if is_array:
self.source_buf = source_buf
self.set_fields = self.set_fields_spatial_value
else:
self.set_fields = self.set_fields_single_value
def __init__(self,
fields,
pt0,
pt1,
ep_inf,
drude_freq,
gamma,
mask_arrays=(1, 1, 1)):
common.check_type('fields', fields, Fields)
common.check_type('pt0', pt0, (list, tuple), (int, float))
common.check_type('pt1', pt1, (list, tuple), (int, float))
common.check_type('ep_inf', ep_inf, (int, float))
common.check_type('drude_freq', drude_freq, (int, float))
common.check_type('gamma', gamma, (int, float))
common.check_type('mask_arrays', mask_arrays, (list, tuple),
(np.ndarray, int))
# local variables
pt0 = common.convert_indices(fields.ns, pt0)
pt1 = common.convert_indices(fields.ns, pt1)
context = fields.context
queue = fields.queue
dtype = fields.dtype
shape = common.shape_two_points(pt0, pt1, is_dummy=True)
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))
for mask_array in mask_arrays:
if isinstance(mask_array, np.ndarray):
assert common.shape_two_points(pt0, pt1) == mask_array.shape, \
'shape mismatch : %s, %s' % (shape, mask_array.shape)
# allocations
psis = [np.zeros(shape, dtype) for i in range(3)]
psi_bufs = [
cl.Buffer(context, cl.mem_flags.READ_WRITE, psi.nbytes)
for psi in psis
]
for psi_buf, psi in zip(psi_bufs, psis):
cl.enqueue_copy(queue, psi_buf, psi)
dt = fields.dt
aa = (2 - gamma * dt) / (2 + gamma * dt)
bb = drude_freq**2 * dt / (2 + gamma * dt)
comm = 2 * ep_inf + bb * dt
ca = 2 * dt / comm
cb = -(aa + 3) * bb * dt / comm
cc = -(aa + 1) * dt / comm
cas = [ca * mask for mask in mask_arrays]
shape = common.shape_two_points(pt0, pt1, is_dummy=True)
f = np.zeros(shape, dtype)
psi_bufs = [
cl.Buffer(context, cl.mem_flags.READ_WRITE, f.nbytes)
for i in range(3)
]
for psi_buf in psi_bufs:
cl.enqueue_copy(queue, psi_buf, f)
cf = np.ones(shape, dtype)
mask_bufs = [
cl.Buffer(context, cl.mem_flags.READ_ONLY, cf.nbytes)
for i in range(3)
]
for mask_buf, mask in zip(mask_bufs, mask_arrays):
cl.enqueue_copy(queue, mask_buf, cf * mask)
# modify ce arrays
slices = common.slices_two_points(pt0, pt1)
for ce, ca in zip(fields.get_ces(), cas):
ce[slices] = ca * mask + ce[slices] * mask.__invert__()
# program
nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(
pt0, pt1)
macros = ['NMAX', 'XID', 'YID', 'ZID', 'DX', 'DTYPE', 'PRAGMA_fp64']
values = [nmax_str, xid_str, yid_str, zid_str,
str(fields.ls)] + fields.dtype_str_list
ksrc = common.replace_template_code( \
open(common_gpu.src_path + 'drude.cl').read(), macros, values)
program = cl.Program(fields.context, ksrc).build()
# arguments
pca = aa
pcb = (aa + 1) * bb
args = fields.ns + [dtype(cb), dtype(cc), dtype(pca), dtype(pcb)] \
+ fields.eh_bufs[:3] + psi_bufs + mask_bufs
# global variables
self.mainf = fields
self.program = program
self.args = args
nx, ny, nz = fields.ns
nmax = int(nmax_str)
remainder = nmax % fields.ls
self.gs = nmax if remainder == 0 else nmax - remainder + fields.ls
# append to the update list
self.priority_type = 'material'
fields.append_instance(self)
