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): """ """ 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