def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True):
"""
"""
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)
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)
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))
# global variables and functions
self.mainf = fields
self.str_fs = str_fs
self.slice_xyz = common.slices_two_points(pt0, pt1)
self.shape = common.shape_two_points(pt0, pt1, len(str_fs))
self.is_overwrite = is_overwrite
if is_array:
self.func = self.set_fields_spatial_value
else:
self.func = self.set_fields_single_value
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 runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slidx = common.slice_index_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
device = gpu_devices[0]
fields = Fields(context, device, nx, ny, nz, '', 'single')
getf = GetFields(fields, str_f, pt0, pt1)
# host allocations
ehs = common_update.generate_random_ehs(nx, ny, nz, fields.dtype)
eh_dict = dict( zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], ehs) )
fields.set_eh_bufs(*ehs)
# verify
getf.get_event().wait()
for str_f in str_fs:
original = eh_dict[str_f][slidx]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
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 runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slidx = common.slice_index_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
device = gpu_devices[0]
fields = Fields(context, device, nx, ny, nz, '')
getf = GetFields(fields, str_f, pt0, pt1)
# host allocations
eh_dict = {}
for sf in str_fs:
eh_dict[sf] = np.random.rand(*fields.ns).astype(fields.dtype)
cl.enqueue_copy(fields.queue, fields.get_buf(sf), eh_dict[sf])
# verify
getf.get_event().wait()
for str_f in str_fs:
original = eh_dict[str_f][slidx]
copy = getf.get_fields(str_f)
self.assertEqual(np.abs(eh_dict[str_f][slidx] - getf.get_fields(str_f)).max(), 0, self.args)
def runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slidx = common.slices_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
fields = Fields(0, nx, ny, nz, '', 'single')
getf = GetFields(fields, str_f, pt0, pt1)
# host allocations
ehs = common_random.generate_ehs(nx, ny, nz, fields.dtype)
eh_dict = dict( zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], ehs) )
fields.set_eh_bufs(*ehs)
# verify
getf.wait()
for str_f in str_fs:
original = eh_dict[str_f][slidx]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
fields.context_pop()
def runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slidx = common.slice_index_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
device = gpu_devices[0]
fields = Fields(context, device, nx, ny, nz, '')
getf = GetFields(fields, str_f, pt0, pt1)
# host allocations
eh_dict = {}
for sf in str_fs:
eh_dict[sf] = np.random.rand(*fields.ns).astype(fields.dtype)
cl.enqueue_copy(fields.queue, fields.get_buf(sf), eh_dict[sf])
# verify
getf.get_event().wait()
for str_f in str_fs:
original = eh_dict[str_f][slidx]
copy = getf.get_fields(str_f)
self.assertEqual(
np.abs(eh_dict[str_f][slidx] - getf.get_fields(str_f)).max(),
0, self.args)
def runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slidx = common.slices_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
device = gpu_devices[0]
fields = Fields(context, device, nx, ny, nz, '', 'single')
getf = GetFields(fields, str_f, pt0, pt1)
# host allocations
ehs = common_random.generate_ehs(nx, ny, nz, fields.dtype)
eh_dict = dict( zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], ehs) )
fields.set_eh_bufs(*ehs)
# verify
getf.get_event().wait()
for str_f in str_fs:
original = eh_dict[str_f][slidx]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
def __init__(self, node_fields, str_f, pt0, pt1, dir_path, tag, max_tstep, is_mpi=False, rank=0):
"""
"""
common.check_type('node_fields', node_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('dir_path', dir_path, str)
common.check_type('tag', tag, int)
common.check_type('max_tstep', max_tstep, int)
# local variables
nodef = node_fields
str_fs = common.convert_to_tuple(str_f)
# setup
getf = GetFields(nodef, str_fs, pt0, pt1)
ndigit = int( ('%e' % max_tstep).split('+')[1] ) + 1
fpath_form = dir_path + '%%.%dd_tag%d.h5' % (ndigit, tag)
if is_mpi:
fpath_form = fpath_form.rstrip('.h5') + '_mpi%d.h5' % rank
# global variables
self.str_fs = str_fs
self.pt0 = pt0
self.pt0 = pt1
self.getf = getf
self.fpath_form = fpath_form
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)
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))
# global variables and functions
self.mainf = fields
self.str_fs = str_fs
self.slice_xyz = common.slice_index_two_points(pt0, pt1)
self.shape = common.shape_two_points(pt0, pt1, len(str_fs))
self.is_overwrite = is_overwrite
if is_array:
self.func = self.set_fields_spatial_value
else:
self.func = self.set_fields_single_value
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 runTest(self):
nx, ny, nz, str_f, pt0, pt1, is_array = self.args
slices = common.slices_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
mainf_list = [gpu.Fields(context, device, nx, ny, nz) \
for device in gpu_devices]
mainf_list.append(cpu.Fields(nx, ny, nz))
nodef = Fields(mainf_list)
dtype = nodef.dtype
anx = nodef.accum_nx_list
getf = GetFields(nodef, str_f, (0, 0, 0),
(nodef.nx - 1, ny - 1, nz - 1))
setf = SetFields(nodef, str_f, pt0, pt1, is_array)
# generate random source
if is_array:
shape = common.shape_two_points(pt0, pt1, len(str_fs))
value = np.random.rand(*shape).astype(nodef.dtype)
split_value = np.split(value, len(str_fs))
split_value_dict = dict(zip(str_fs, split_value))
else:
value = np.random.ranf()
# host allocations
global_ehs = [np.zeros(nodef.ns, dtype) for i in range(6)]
eh_dict = dict(zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], global_ehs))
# verify
for str_f in str_fs:
if is_array:
eh_dict[str_f][slices] = split_value_dict[str_f]
else:
eh_dict[str_f][slices] = value
setf.set_fields(value)
gpu_getf = gpu.GetFields(mainf_list[0], str_fs, (0, 0, 0),
(nx - 1, ny - 1, nz - 1))
gpu_getf.get_event().wait()
getf.wait()
for str_f in str_fs:
original = eh_dict[str_f]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
#if norm != 0:
#print '\ngpu getf\n', gpu_getf.get_fields(str_f)
#print original[slices]
#print copy[slices]
self.assertEqual(norm, 0, '%s, %g, %s' % (self.args, norm, str_f))
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 set_savefields(self, str_f, pt0, pt1, tstep_range, dir_path, tag=0):
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('tstep_range', tstep_range, (list, tuple), int)
common.check_type('dir_path', dir_path, str)
common.check_type('tag', tag, int)
pt0 = common.convert_indices(self.ns, pt0)
pt1 = common.convert_indices(self.ns, pt1)
t0 = common.convert_index(self.max_tstep, tstep_range[0])
t1 = common.convert_index(self.max_tstep, tstep_range[1]) + 1
tgap = tstep_range[2]
tmax = self.max_tstep
tag = tag if tag not in self.savef_tag_list else max(self.savef_tag_list)+1
self.savef_tag_list.append(tag)
if is_mpi:
overlap = common.overlap_two_regions(self.node_pt0, self.node_pt1, pt0, pt1)
if overlap != None:
sx0, sy0, sz0 = self.node_pt0
ox0, oy0, oz0 = overlap[0]
ox1, oy1, oz1 = overlap[1]
local_pt0 = (ox0-sx0, oy0-sy0, oz0-sz0)
local_pt1 = (ox1-sx0, oy1-sy0, oz1-sz0)
savef = node.SaveFields(self.nodef, str_f, local_pt0, local_pt1, dir_path, tag, tmax, True, rank)
self.savef_list.append(savef)
else:
savef = node.SaveFields(self.nodef, str_f, pt0, pt1, dir_path, tag, tmax)
self.savef_list.append(savef)
# save the subdomain informations as pickle
if self.is_master:
if is_mpi:
divide_info_dict = common_mpi.divide_info_dict(size, self.mpi_shape, pt0, pt1, self.asn_dict)
else:
divide_info_dict = { \
'shape': common.shape_two_points(pt0, pt1), \
'pt0': pt0, \
'pt1': pt1, \
'anx_list': self.nodef.accum_nx_list }
divide_info_dict['str_fs'] = common.convert_to_tuple(str_f)
divide_info_dict['tmax'] = tmax
divide_info_dict['tgap'] = tgap
pkl_file = open(dir_path + 'divide_info.pkl', 'wb')
pickle.dump(divide_info_dict, pkl_file)
pkl_file.close()
self.t0 = t0
self.t1 = t1
self.tgap = tgap
def runTest(self):
nx, ny, nz, str_f, pt0, pt1, is_array = self.args
slices = common.slices_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
mainf_list = [gpu.Fields(context, device, nx, ny, nz) \
for device in gpu_devices]
mainf_list.append( cpu.Fields(nx, ny, nz) )
nodef = Fields(mainf_list)
dtype = nodef.dtype
anx = nodef.accum_nx_list
getf = GetFields(nodef, str_f, (0, 0, 0), (nodef.nx-1, ny-1, nz-1))
setf = SetFields(nodef, str_f, pt0, pt1, is_array)
# generate random source
if is_array:
shape = common.shape_two_points(pt0, pt1, len(str_fs))
value = np.random.rand(*shape).astype(nodef.dtype)
split_value = np.split(value, len(str_fs))
split_value_dict = dict( zip(str_fs, split_value) )
else:
value = np.random.ranf()
# host allocations
global_ehs = [np.zeros(nodef.ns, dtype) for i in range(6)]
eh_dict = dict( zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], global_ehs) )
# verify
for str_f in str_fs:
if is_array:
eh_dict[str_f][slices] = split_value_dict[str_f]
else:
eh_dict[str_f][slices] = value
setf.set_fields(value)
gpu_getf = gpu.GetFields(mainf_list[0], str_fs, (0, 0, 0), (nx-1, ny-1, nz-1))
gpu_getf.get_event().wait()
getf.wait()
for str_f in str_fs:
original = eh_dict[str_f]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
#if norm != 0:
#print '\ngpu getf\n', gpu_getf.get_fields(str_f)
#print original[slices]
#print copy[slices]
self.assertEqual(norm, 0, '%s, %g, %s' % (self.args, norm, str_f))
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 runTest(self):
if len(self.args) == 6:
nx, ny, nz, str_f, pt0, pt1 = self.args
src_is_array = False
elif len(self.args) == 7:
nx, ny, nz, str_f, pt0, pt1, src_is_array = self.args
slidx = common.slice_index_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
device = gpu_devices[0]
fields = Fields(context, device, nx, ny, nz, '')
setf = SetFields(fields, str_f, pt0, pt1, src_is_array)
# generate random source
if src_is_array:
shape = list(common.shape_two_points(pt0, pt1))
shape[0] *= len(str_fs)
value = np.random.rand(*shape).astype(fields.dtype)
split_value = np.split(value, len(str_fs))
split_value_dict = dict(zip(str_fs, split_value))
else:
value = np.random.ranf()
# host allocations
eh_dict = {}
for sf in str_fs:
eh_dict[sf] = np.zeros(fields.ns, dtype=fields.dtype)
gpu_eh = np.zeros(fields.ns, dtype=fields.dtype)
# verify
for str_f in str_fs:
if src_is_array:
eh_dict[str_f][slidx] = split_value_dict[str_f]
else:
eh_dict[str_f][slidx] = value
setf.set_fields(value)
for str_f in str_fs:
cl.enqueue_copy(fields.queue, gpu_eh, fields.get_buf(str_f))
self.assertEqual(
np.abs(eh_dict[str_f] - gpu_eh).max(), 0, self.args)
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, 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 runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slices = common.slice_index_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
mainf_list = [gpu.Fields(context, device, nx, ny, nz) \
for device in gpu_devices]
mainf_list.append(cpu.Fields(nx, ny, nz))
nodef = NodeFields(mainf_list)
dtype = nodef.dtype
anx = nodef.accum_nx_list
getf = NodeGetFields(nodef, str_f, pt0, pt1)
# generate random source
global_ehs = [np.zeros(nodef.ns, dtype) for i in range(6)]
eh_dict = dict(zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], global_ehs))
for i, f in enumerate(mainf_list[:-1]):
nx, ny, nz = f.ns
ehs = common_update.generate_random_ehs(nx, ny, nz, dtype)
f.set_eh_bufs(*ehs)
for eh, geh in zip(ehs, global_ehs):
geh[anx[i]:anx[i + 1], :, :] = eh[:-1, :, :]
f = mainf_list[-1]
nx, ny, nz = f.ns
ehs = common_update.generate_random_ehs(nx, ny, nz, dtype)
f.set_ehs(*ehs)
for eh, geh in zip(ehs, global_ehs):
geh[anx[-2]:anx[-1] + 1, :, :] = eh[:]
# verify
getf.wait()
for str_f in str_fs:
original = eh_dict[str_f][slices]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
def runTest(self):
if len(self.args) == 6:
nx, ny, nz, str_f, pt0, pt1 = self.args
src_is_array = False
elif len(self.args) == 7:
nx, ny, nz, str_f, pt0, pt1, src_is_array = self.args
slidx = common.slice_index_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
device = gpu_devices[0]
fields = Fields(context, device, nx, ny, nz, '')
setf = SetFields(fields, str_f, pt0, pt1, src_is_array)
# generate random source
if src_is_array:
shape = list( common.shape_two_points(pt0, pt1) )
shape[0] *= len(str_fs)
value = np.random.rand(*shape).astype(fields.dtype)
split_value = np.split(value, len(str_fs))
split_value_dict = dict( zip(str_fs, split_value) )
else:
value = np.random.ranf()
# host allocations
eh_dict = {}
for sf in str_fs:
eh_dict[sf] = np.zeros(fields.ns, dtype=fields.dtype)
gpu_eh = np.zeros(fields.ns, dtype=fields.dtype)
# verify
for str_f in str_fs:
if src_is_array:
eh_dict[str_f][slidx] = split_value_dict[str_f]
else:
eh_dict[str_f][slidx] = value
setf.set_fields(value)
for str_f in str_fs:
cl.enqueue_copy(fields.queue, gpu_eh, fields.get_buf(str_f))
self.assertEqual(np.abs(eh_dict[str_f] - gpu_eh).max(), 0, self.args)
def runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slices = common.slices_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
mainf_list = [gpu.Fields(context, device, nx, ny, nz) \
for device in gpu_devices]
mainf_list.append( cpu.Fields(nx, ny, nz) )
nodef = Fields(mainf_list)
dtype = nodef.dtype
anx = nodef.accum_nx_list
getf = GetFields(nodef, str_f, pt0, pt1)
# generate random source
global_ehs = [np.zeros(nodef.ns, dtype) for i in range(6)]
eh_dict = dict( zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], global_ehs) )
for i, f in enumerate(mainf_list[:-1]):
nx, ny, nz = f.ns
ehs = common_random.generate_ehs(nx, ny, nz, dtype)
f.set_eh_bufs(*ehs)
for eh, geh in zip(ehs, global_ehs):
geh[anx[i]:anx[i+1],:,:] = eh[:-1,:,:]
f = mainf_list[-1]
nx, ny, nz = f.ns
ehs = common_random.generate_ehs(nx, ny, nz, dtype)
f.set_ehs(*ehs)
for eh, geh in zip(ehs, global_ehs):
geh[anx[-2]:anx[-1]+1,:,:] = eh[:]
# verify
getf.wait()
for str_f in str_fs:
original = eh_dict[str_f][slices]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
def runTest(self):
nx, ny, nz, str_f, pt0, pt1, is_array = self.args
slidx = common.slices_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
gpu_devices = common_gpu.gpu_device_list(print_info=False)
context = cl.Context(gpu_devices)
device = gpu_devices[0]
qtask = QueueTask()
fields = Fields(context, device, qtask, nx, ny, nz, '', 'single')
setf = SetFields(fields, str_f, pt0, pt1, is_array)
# generate random source
if is_array:
shape = common.shape_two_points(pt0, pt1, len(str_fs))
value = np.random.rand(*shape).astype(fields.dtype)
split_value = np.split(value, len(str_fs))
split_value_dict = dict( zip(str_fs, split_value) )
else:
value = np.random.ranf()
# host allocations
ehs = [np.zeros(fields.ns, dtype=fields.dtype) for i in range(6)]
eh_dict = dict( zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], ehs) )
gpu_eh = np.zeros(fields.ns_pitch, dtype=fields.dtype)
# verify
for str_f in str_fs:
if is_array:
eh_dict[str_f][slidx] = split_value_dict[str_f]
else:
eh_dict[str_f][slidx] = value
setf.set_fields(value)
setf.mainf.enqueue_barrier()
for str_f in str_fs:
cl.enqueue_copy(fields.queue, gpu_eh, fields.get_buf(str_f))
original = eh_dict[str_f]
copy = gpu_eh[:,:,fields.slice_z]
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
def runTest(self):
nx, ny, nz, str_f, pt0, pt1, is_array = self.args
slidx = common.slices_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
fields = Fields(0, nx, ny, nz, '', 'single')
setf = SetFields(fields, str_f, pt0, pt1, is_array)
# generate random source
if is_array:
shape = common.shape_two_points(pt0, pt1, len(str_fs))
value = np.random.rand(*shape).astype(fields.dtype)
split_value = np.split(value, len(str_fs))
split_value_dict = dict( zip(str_fs, split_value) )
else:
value = np.random.ranf()
# host allocations
ehs = [np.zeros(fields.ns, dtype=fields.dtype) for i in range(6)]
eh_dict = dict( zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], ehs) )
gpu_eh = np.zeros(fields.ns_pitch, dtype=fields.dtype)
# verify
for str_f in str_fs:
if is_array:
eh_dict[str_f][slidx] = split_value_dict[str_f]
else:
eh_dict[str_f][slidx] = value
setf.set_fields(value)
for str_f in str_fs:
cuda.memcpy_dtoh(gpu_eh, fields.get_buf(str_f))
original = eh_dict[str_f]
copy = gpu_eh[:,:,fields.slice_z]
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
fields.context_pop()
def __init__(self,
node_fields,
str_f,
pt0,
pt1,
dir_path,
tag,
max_tstep,
is_mpi=False,
rank=0):
"""
"""
common.check_type('node_fields', node_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('dir_path', dir_path, str)
common.check_type('tag', tag, int)
common.check_type('max_tstep', max_tstep, int)
# local variables
nodef = node_fields
str_fs = common.convert_to_tuple(str_f)
# setup
getf = GetFields(nodef, str_fs, pt0, pt1)
ndigit = int(('%e' % max_tstep).split('+')[1]) + 1
fpath_form = dir_path + '%%.%dd_tag%d.h5' % (ndigit, tag)
if is_mpi:
fpath_form = fpath_form.rstrip('.h5') + '_mpi%d.h5' % rank
# global variables
self.str_fs = str_fs
self.pt0 = pt0
self.pt0 = pt1
self.getf = getf
self.fpath_form = fpath_form
def runTest(self):
nx, ny, nz, str_f, pt0, pt1 = self.args
slice_xyz = common.slice_index_two_points(pt0, pt1)
str_fs = common.convert_to_tuple(str_f)
# instance
fields = Fields(nx, ny, nz, '', 'single')
getf = GetFields(fields, str_f, pt0, pt1)
# host allocations
ehs = common_update.generate_random_ehs(nx, ny, nz, fields.dtype)
eh_dict = dict(zip(['ex', 'ey', 'ez', 'hx', 'hy', 'hz'], ehs))
fields.set_ehs(*ehs)
# verify
getf.get_event().wait()
for str_f in str_fs:
original = eh_dict[str_f][slice_xyz]
copy = getf.get_fields(str_f)
norm = np.linalg.norm(original - copy)
self.assertEqual(norm, 0, '%s, %g' % (self.args, norm))
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):
"""
"""
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
nodef = fields
str_fs = common.convert_to_tuple(str_f)
mainf_list = nodef.mainf_list
anx = nodef.accum_nx_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'], nodef.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))
dummied_shape = common.shape_two_points(pt0, pt1, is_dummy=True)
host_array = np.zeros(shape, dtype=nodef.dtype)
split_host_array = np.split(host_array, len(str_fs))
split_host_array_dict = dict(zip(str_fs, split_host_array))
getf_list = []
slices_list = []
self.gpu, self.cpu = gpu, cpu
for i, mainf in enumerate(mainf_list):
nx0 = anx[i]
nx1 = anx[i + 1] - 1 if i < len(mainf_list) - 1 else anx[i + 1]
overlap = common.intersection_two_lines((nx0, nx1),
(pt0[0], pt1[0]))
if overlap != None:
x0, y0, z0 = pt0
x1, y1, z1 = pt1
slice_pt0 = (overlap[0] - x0, 0, 0)
slice_pt1 = (overlap[1] - x0, y1 - y0, z1 - z0)
slices = []
for j, p0, p1 in zip([0, 1, 2], slice_pt0, slice_pt1):
if dummied_shape[j] != 1:
slices.append(slice(p0, p1 + 1))
slices_list.append(slices if slices != [] else [slice(0, 1)])
local_pt0 = (overlap[0] - nx0, y0, z0)
local_pt1 = (overlap[1] - nx0, y1, z1)
getf_list.append( getattr(self, mainf.device_type). \
GetFields(mainf, str_fs, local_pt0, local_pt1) )
# global variables
self.str_fs = str_fs
self.host_array = host_array
self.split_host_array_dict = split_host_array_dict
self.getf_list = getf_list
self.slices_list = slices_list
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
nodef = fields
str_fs = common.convert_to_tuple(str_f)
mainf_list = nodef.mainf_list
anx = nodef.accum_nx_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'], nodef.ns, pt0, pt1):
common.check_value('pt0 %s' % axis, p0, range(n))
common.check_value('pt1 %s' % axis, p1, range(n))
# allocation
dummied_shape = common.shape_two_points(pt0, pt1, is_dummy=True)
setf_list = []
slices_list = []
self.gpu, self.cpu = gpu, cpu
for i, mainf in enumerate(mainf_list):
nx0 = anx[i]
nx1 = anx[i + 1]
overlap = common.intersection_two_lines((nx0, nx1),
(pt0[0], pt1[0]))
if overlap != None:
x0, y0, z0 = pt0
x1, y1, z1 = pt1
slice_pt0 = (overlap[0] - x0, 0, 0)
slice_pt1 = (overlap[1] - x0, y1 - y0, z1 - z0)
slices = []
for j, p0, p1 in zip([0, 1, 2], slice_pt0, slice_pt1):
if dummied_shape[j] != 1:
slices.append(slice(p0, p1 + 1))
slices_list.append(slices if slices != [] else [slice(0, 1)])
local_pt0 = (overlap[0] - nx0, y0, z0)
local_pt1 = (overlap[1] - nx0, y1, z1)
setf_list.append( getattr(self, mainf.device_type). \
SetFields(mainf, str_fs, local_pt0, local_pt1, \
is_array, is_overwrite) )
# global variables and functions
self.str_fs = str_fs
self.dtype = nodef.dtype
self.shape = common.shape_two_points(pt0, pt1, len(str_fs))
self.setf_list = setf_list
self.slices_list = slices_list
if is_array:
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)
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
nodef = fields
str_fs = common.convert_to_tuple(str_f)
mainf_list = nodef.mainf_list
anx = nodef.accum_nx_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'], nodef.ns, pt0, pt1):
common.check_value('pt0 %s' % axis, p0, range(n))
common.check_value('pt1 %s' % axis, p1, range(n))
# allocation
dummied_shape = common.shape_two_points(pt0, pt1, is_dummy=True)
setf_list = []
slices_list = []
self.gpu, self.cpu = gpu, cpu
for i, mainf in enumerate(mainf_list):
nx0 = anx[i]
nx1 = anx[i+1]
overlap = common.intersection_two_lines((nx0, nx1), (pt0[0], pt1[0]))
if overlap != None:
x0, y0, z0 = pt0
x1, y1, z1 = pt1
slice_pt0 = (overlap[0]-x0, 0, 0)
slice_pt1 = (overlap[1]-x0, y1-y0, z1-z0)
slices = []
for j, p0, p1 in zip([0, 1, 2], slice_pt0, slice_pt1):
if dummied_shape[j] != 1:
slices.append( slice(p0, p1+1) )
slices_list.append(slices if slices!=[] else [slice(0, 1)] )
local_pt0 = (overlap[0]-nx0, y0, z0)
local_pt1 = (overlap[1]-nx0, y1, z1)
setf_list.append( getattr(self, mainf.device_type). \
SetFields(mainf, str_fs, local_pt0, local_pt1, \
is_array, is_overwrite) )
# global variables and functions
self.str_fs = str_fs
self.dtype = nodef.dtype
self.shape = common.shape_two_points(pt0, pt1, len(str_fs))
self.setf_list = setf_list
self.slices_list = slices_list
if is_array:
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):
"""
"""
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
nodef = fields
str_fs = common.convert_to_tuple(str_f)
mainf_list = nodef.mainf_list
anx = nodef.accum_nx_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'], nodef.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))
dummied_shape = common.shape_two_points(pt0, pt1, is_dummy=True)
host_array = np.zeros(shape, dtype=nodef.dtype)
split_host_array = np.split(host_array, len(str_fs))
split_host_array_dict = dict( zip(str_fs, split_host_array) )
getf_list = []
slices_list = []
self.gpu, self.cpu = gpu, cpu
for i, mainf in enumerate(mainf_list):
nx0 = anx[i]
nx1 = anx[i+1]-1 if i < len(mainf_list)-1 else anx[i+1]
overlap = common.intersection_two_lines((nx0, nx1), (pt0[0], pt1[0]))
if overlap != None:
x0, y0, z0 = pt0
x1, y1, z1 = pt1
slice_pt0 = (overlap[0]-x0, 0, 0)
slice_pt1 = (overlap[1]-x0, y1-y0, z1-z0)
slices = []
for j, p0, p1 in zip([0, 1, 2], slice_pt0, slice_pt1):
if dummied_shape[j] != 1:
slices.append( slice(p0, p1+1) )
slices_list.append(slices if slices!=[] else [slice(0, 1)] )
local_pt0 = (overlap[0]-nx0, y0, z0)
local_pt1 = (overlap[1]-nx0, y1, z1)
getf_list.append( getattr(self, mainf.device_type). \
GetFields(mainf, str_fs, local_pt0, local_pt1) )
# global variables
self.str_fs = str_fs
self.host_array = host_array
self.split_host_array_dict = split_host_array_dict
self.getf_list = getf_list
self.slices_list = slices_list
def set_savefields(self, str_f, pt0, pt1, tstep_range, dir_path, tag=0):
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('tstep_range', tstep_range, (list, tuple), int)
common.check_type('dir_path', dir_path, str)
common.check_type('tag', tag, int)
pt0 = common.convert_indices(self.ns, pt0)
pt1 = common.convert_indices(self.ns, pt1)
t0 = common.convert_index(self.max_tstep, tstep_range[0])
t1 = common.convert_index(self.max_tstep, tstep_range[1]) + 1
tgap = tstep_range[2]
tmax = self.max_tstep
tag = tag if tag not in self.savef_tag_list else max(
self.savef_tag_list) + 1
self.savef_tag_list.append(tag)
if is_mpi:
overlap = common.overlap_two_regions(self.node_pt0, self.node_pt1,
pt0, pt1)
if overlap != None:
sx0, sy0, sz0 = self.node_pt0
ox0, oy0, oz0 = overlap[0]
ox1, oy1, oz1 = overlap[1]
local_pt0 = (ox0 - sx0, oy0 - sy0, oz0 - sz0)
local_pt1 = (ox1 - sx0, oy1 - sy0, oz1 - sz0)
savef = node.SaveFields(self.nodef, str_f, local_pt0,
local_pt1, dir_path, tag, tmax, True,
rank)
self.savef_list.append(savef)
else:
savef = node.SaveFields(self.nodef, str_f, pt0, pt1, dir_path, tag,
tmax)
self.savef_list.append(savef)
# save the subdomain informations as pickle
if self.is_master:
if is_mpi:
divide_info_dict = common_mpi.divide_info_dict(
size, self.mpi_shape, pt0, pt1, self.asn_dict)
else:
divide_info_dict = { \
'shape': common.shape_two_points(pt0, pt1), \
'pt0': pt0, \
'pt1': pt1, \
'anx_list': self.nodef.accum_nx_list }
divide_info_dict['str_fs'] = common.convert_to_tuple(str_f)
divide_info_dict['tmax'] = tmax
divide_info_dict['tgap'] = tgap
pkl_file = open(dir_path + 'divide_info.pkl', 'wb')
pickle.dump(divide_info_dict, pkl_file)
pkl_file.close()
self.t0 = t0
self.t1 = t1
self.tgap = tgap
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
