def __init__(self, fields, axes): """ """ common.check_type('fields', fields, Fields) common.check_type('axes', axes, str) assert len(set(axes).intersection( set('xyz'))) > 0, 'axes option is wrong: %s is given' % repr(axes) # local variables nx, ny, nz = fields.ns dtype_str_list = fields.dtype_str_list # program macros = ['NMAX', 'IDX0', 'IDX1', 'DTYPE', 'PRAGMA_fp64'] program_dict = {} gs_dict = {} for axis in list(axes): program_dict[axis] = {} for e_or_h in ['e', 'h']: pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['get'] pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['get'] nmaxi_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list( pt0, pt1) idx0_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str, zid_str) pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['set'] pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['set'] nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list( pt0, pt1) idx1_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str, zid_str) values = [nmax_str, idx0_str, idx1_str] + dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy_self.cl').read(), \ macros, values) program = cl.Program(fields.context, ksrc).build() program_dict[axis][e_or_h] = program nmax = int(nmax_str) remainder = nmax % fields.ls gs_dict[ axis] = nmax if remainder == 0 else nmax - remainder + fields.ls # global variables self.mainf = fields self.axes = axes self.program_dict = program_dict self.gs_dict = gs_dict # append to the update list self.priority_type = 'pbc' self.mainf.append_instance(self)
def __init__(self, fields, axes): """ """ common.check_type('fields', fields, Fields) common.check_type('axes', axes, str) assert len( set(axes).intersection(set('xyz')) ) > 0, 'axes option is wrong: %s is given' % repr(axes) # local variables nx, ny, nz = fields.ns dtype_str_list = fields.dtype_str_list # program macros = ['NMAX', 'IDX0', 'IDX1', 'DTYPE', 'PRAGMA_fp64'] program_dict = {} gs_dict = {} for axis in list(axes): program_dict[axis] = {} for e_or_h in ['e', 'h']: pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['get'] pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['get'] nmaxi_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1) idx0_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str, zid_str) pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]['set'] pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]['set'] nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1) idx1_str = '%s*ny*nz + %s*nz + %s' % (xid_str, yid_str, zid_str) values = [nmax_str, idx0_str, idx1_str] + dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy_self.cl').read(), \ macros, values) program = cl.Program(fields.context, ksrc).build() program_dict[axis][e_or_h] = program nmax = int(nmax_str) remainder = nmax % fields.ls gs_dict[axis] = nmax if remainder == 0 else nmax - remainder + fields.ls # global variables self.mainf = fields self.axes = axes self.program_dict = program_dict self.gs_dict = gs_dict # append to the update list self.priority_type = 'pbc' self.mainf.append_instance(self)
def __init__(self, fields, axes): """ """ common.check_type("fields", fields, Fields) common.check_type("axes", axes, str) assert len(set(axes).intersection(set("xyz"))) > 0, "axes option is wrong: %s is given" % repr(axes) # local variables nx, ny, nz = fields.ns dtype_str_list = fields.dtype_str_list # program macros = ["NMAX", "IDX0", "IDX1", "DTYPE", "PRAGMA_fp64"] program_dict = {} for axis in list(axes): program_dict[axis] = {} for e_or_h in ["e", "h"]: pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]["get"] pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]["get"] nmax, xid, yid, zid = common_gpu.macro_replace_list(pt0, pt1) idx0_str = "%s*ny*nz + %s*nz + %s" % (xid, yid, zid) pt0 = common_exchange.pt0_dict(nx, ny, nz)[axis][e_or_h]["set"] pt1 = common_exchange.pt1_dict(nx, ny, nz)[axis][e_or_h]["set"] nmax, xid, yid, zid = common_gpu.macro_replace_list(pt0, pt1) idx1_str = "%s*ny*nz + %s*nz + %s" % (xid, yid, zid) values = [nmax, idx0_str, idx1_str] + dtype_str_list ksrc = common.replace_template_code(open(common_gpu.src_path + "copy_self.cl").read(), macros, values) program = cl.Program(fields.context, ksrc).build() program_dict[axis][e_or_h] = program # global variables self.mainf = fields self.axes = axes self.program_dict = program_dict # append to the update list self.priority_type = "pbc" self.mainf.append_instance(self)
def __init__(self, fields, str_f, pt0, pt1): """ """ common.check_type('fields', fields, Fields) common.check_type('str_f', str_f, (str, list, tuple), str) common.check_type('pt0', pt0, (list, tuple), (int, float)) common.check_type('pt1', pt1, (list, tuple), (int, float)) pt0 = list(common.convert_indices(fields.ns, pt0)) pt1 = list(common.convert_indices(fields.ns, pt1)) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list for strf in str_fs: strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz'] common.check_value('str_f', strf, strf_list) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) # program macros = ['NMAX', 'XID', 'YID', 'ZID', \ 'ARGS', \ 'TARGET', 'SOURCE', 'OVERWRITE', \ 'DTYPE'] values = common_gpu.macro_replace_list(pt0, pt1) + \ ['DTYPE *source', \ 'target[sub_idx]', 'source[idx]', '='] + \ dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy.cu').read(), macros, values) program = SourceModule(ksrc) kernel_copy = program.get_function('copy') # allocation source_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) host_array = np.zeros(shape, fields.dtype) split_host_array = np.split(host_array, len(str_fs)) split_host_array_dict = dict(zip(str_fs, split_host_array)) target_buf = cuda.to_device(host_array) # global variables self.mainf = fields self.kernel_copy = kernel_copy self.source_bufs = source_bufs self.target_buf = target_buf self.host_array = host_array self.split_host_array_dict = split_host_array_dict
def __init__(self, fields, str_f, pt0, pt1): """ """ common.check_type('fields', fields, Fields) common.check_type('str_f', str_f, (str, list, tuple), str) common.check_type('pt0', pt0, (list, tuple), (int, float)) common.check_type('pt1', pt1, (list, tuple), (int, float)) pt0 = list( common.convert_indices(fields.ns, pt0) ) pt1 = list( common.convert_indices(fields.ns, pt1) ) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list for strf in str_fs: strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz'] common.check_value('str_f', strf, strf_list) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) # program macros = ['NMAX', 'XID', 'YID', 'ZID', \ 'ARGS', \ 'TARGET', 'SOURCE', 'OVERWRITE', \ 'DTYPE'] values = common_gpu.macro_replace_list(pt0, pt1) + \ ['DTYPE *source', \ 'target[sub_idx]', 'source[idx]', '='] + \ dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy.cu').read(), macros, values) program = SourceModule(ksrc) kernel_copy = program.get_function('copy') # allocation source_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) host_array = np.zeros(shape, fields.dtype) split_host_array = np.split(host_array, len(str_fs)) split_host_array_dict = dict( zip(str_fs, split_host_array) ) target_buf = cuda.to_device(host_array) # global variables self.mainf = fields self.kernel_copy = kernel_copy self.source_bufs = source_bufs self.target_buf = target_buf self.host_array = host_array self.split_host_array_dict = split_host_array_dict
def __init__(self, fields, str_f, pt0, pt1): """ """ common.check_type('fields', fields, Fields) common.check_type('str_f', str_f, (str, list, tuple), str) common.check_type('pt0', pt0, (list, tuple), int) common.check_type('pt1', pt1, (list, tuple), int) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list for strf in str_fs: strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz'] common.check_value('str_f', strf, strf_list) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) # program macros = ['NMAX', 'XID', 'YID', 'ZID', \ 'ARGS', \ 'TARGET', 'SOURCE', 'OVERWRITE', \ 'DTYPE', 'PRAGMA_fp64'] values = common_gpu.macro_replace_list(pt0, pt1) + \ ['__global DTYPE *source', \ 'target[sub_idx]', 'source[idx]', '='] + \ dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy.cl').read(), macros, values) program = cl.Program(fields.context, ksrc).build() # allocation source_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) host_array = np.zeros(shape, dtype=fields.dtype) split_host_array = np.split(host_array, len(str_fs)) split_host_array_dict = dict(zip(str_fs, split_host_array)) target_buf = cl.Buffer( \ fields.context, \ cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, \ hostbuf=host_array) # global variables self.mainf = fields self.program = program self.source_bufs = source_bufs self.target_buf = target_buf self.host_array = host_array self.split_host_array_dict = split_host_array_dict
def __init__(self, fields, axis): """ """ common.check_type('fields', fields, Fields) common.check_value('axis', axis, ['x', 'y', 'z']) # local variables nx, ny, nz = fields.ns dtype_str_list = fields.dtype_str_list # program replace = lambda lst, idx, val: lst[:idx] + [val] + lst[idx+1:] base0 = {'e': [0, 0, 0], 'h': [1, 1, 1]} base1 = {'e': [nx-2, ny-2, nz-2], 'h': [nx-1, ny-1, nz-1]} axis_id = {'x':0, 'y':1, 'z':2}[axis] nn = fields.ns[axis_id] value_dict = {'e': [], 'h': []} for eh in ['e', 'h']: for idx in {'e': [0, nn-1], 'h':[nn-1, 0]}[eh]: pt0 = replace(base0[eh], axis_id, idx) pt1 = replace(base1[eh], axis_id, idx) nmax, xid, yid, zid = \ common_gpu.macro_replace_list(pt0, pt1) value_dict[eh].append( \ '%s*ny*nz + %s*nz + %s' % (xid, yid, zid) ) macros = ['NMAX', 'IDX0', 'IDX1', 'DTYPE', 'PRAGMA_fp64'] values_e = [nmax] + value_dict['e'] + dtype_str_list values_h = [nmax] + value_dict['h'] + dtype_str_list ksrc_e = common.replace_template_code( \ open(common_gpu.src_path + 'copy_self.cl').read(), \ macros, values_e) ksrc_h = common.replace_template_code( \ open(common_gpu.src_path + 'copy_self.cl').read(), \ macros, values_h) program_e = cl.Program(fields.context, ksrc_e).build() program_h = cl.Program(fields.context, ksrc_h).build() # global variables self.mainf = fields self.program_e = program_e self.program_h = program_h self.strfs_e = {'x':['ey','ez'], 'y':['ex','ez'], 'z':['ex','ey']}[axis] self.strfs_h = {'x':['hy','hz'], 'y':['hx','hz'], 'z':['hx','hy']}[axis] # append to the update list self.priority_type = 'pbc' self.mainf.append_instance(self)
def __init__(self, fields, str_f, pt0, pt1): """ """ common.check_type("fields", fields, Fields) common.check_type("str_f", str_f, (str, list, tuple), str) common.check_type("pt0", pt0, (list, tuple), int) common.check_type("pt1", pt1, (list, tuple), int) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list for strf in str_fs: strf_list = ["ex", "ey", "ez", "hx", "hy", "hz"] common.check_value("str_f", strf, strf_list) for axis, n, p0, p1 in zip(["x", "y", "z"], fields.ns, pt0, pt1): common.check_value("pt0 %s" % axis, p0, range(n)) common.check_value("pt1 %s" % axis, p1, range(n)) # program macros = ["NMAX", "XID", "YID", "ZID", "ARGS", "TARGET", "SOURCE", "OVERWRITE", "DTYPE", "PRAGMA_fp64"] values = ( common_gpu.macro_replace_list(pt0, pt1) + ["__global DTYPE *source", "target[sub_idx]", "source[idx]", "="] + dtype_str_list ) ksrc = common.replace_template_code(open(common_gpu.src_path + "copy.cl").read(), macros, values) program = cl.Program(fields.context, ksrc).build() # allocation source_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) host_array = np.zeros(shape, dtype=fields.dtype) split_host_array = np.split(host_array, len(str_fs)) split_host_array_dict = dict(zip(str_fs, split_host_array)) target_buf = cl.Buffer(fields.context, cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, hostbuf=host_array) # global variables self.mainf = fields self.program = program self.source_bufs = source_bufs self.target_buf = target_buf self.host_array = host_array self.split_host_array_dict = split_host_array_dict
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True): """ """ common.check_type('fields', fields, Fields) common.check_type('str_f', str_f, (str, list, tuple), str) common.check_type('pt0', pt0, (list, tuple), (int, float)) common.check_type('pt1', pt1, (list, tuple), (int, float)) common.check_type('is_array', is_array, bool) common.check_type('is_overwrite', is_overwrite, bool) pt0 = list(common.convert_indices(fields.ns, pt0)) pt1 = list(common.convert_indices(fields.ns, pt1)) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list overwrite_str = {True: '=', False: '+='}[is_overwrite] for strf in str_fs: strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz'] common.check_value('str_f', strf, strf_list) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) # program macros = ['NMAX', 'XID', 'YID', 'ZID', \ 'ARGS', \ 'TARGET', 'SOURCE', 'OVERWRITE', \ 'DTYPE'] if is_array: values = common_gpu.macro_replace_list(pt0, pt1) + \ ['DTYPE *source', \ 'target[idx]', 'source[sub_idx]', overwrite_str] + \ dtype_str_list else: values = common_gpu.macro_replace_list(pt0, pt1) + \ ['DTYPE source', \ 'target[idx]', 'source', overwrite_str] + \ dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy.cu').read(), macros, values) program = SourceModule(ksrc) kernel_copy = program.get_function('copy') # allocation target_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) if is_array: tmp_array = np.zeros(shape, fields.dtype) source_buf = cuda.to_device(tmp_array) # global variabels and functions self.mainf = fields self.kernel_copy = kernel_copy self.target_bufs = target_bufs self.shape = shape if is_array: self.source_buf = source_buf self.set_fields = self.set_fields_spatial_value else: self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True): """ """ common.check_type('fields', fields, Fields) common.check_type('str_f', str_f, (str, list, tuple), str) common.check_type('pt0', pt0, (list, tuple), (int, float)) common.check_type('pt1', pt1, (list, tuple), (int, float)) common.check_type('is_array', is_array, bool) common.check_type('is_overwrite', is_overwrite, bool) pt0 = list( common.convert_indices(fields.ns, pt0) ) pt1 = list( common.convert_indices(fields.ns, pt1) ) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list overwrite_str = {True: '=', False: '+='}[is_overwrite] for strf in str_fs: strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz'] common.check_value('str_f', strf, strf_list) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) # program macros = ['NMAX', 'XID', 'YID', 'ZID', \ 'ARGS', \ 'TARGET', 'SOURCE', 'OVERWRITE', \ 'DTYPE'] if is_array: values = common_gpu.macro_replace_list(pt0, pt1) + \ ['DTYPE *source', \ 'target[idx]', 'source[sub_idx]', overwrite_str] + \ dtype_str_list else: values = common_gpu.macro_replace_list(pt0, pt1) + \ ['DTYPE source', \ 'target[idx]', 'source', overwrite_str] + \ dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy.cu').read(), macros, values) program = SourceModule(ksrc) kernel_copy = program.get_function('copy') # allocation target_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) if is_array: tmp_array = np.zeros(shape, fields.dtype) source_buf = cuda.to_device(tmp_array) # global variabels and functions self.mainf = fields self.kernel_copy = kernel_copy self.target_bufs = target_bufs self.shape = shape if is_array: self.source_buf = source_buf self.set_fields = self.set_fields_spatial_value else: self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True): """ """ common.check_type('fields', fields, Fields) common.check_type('str_f', str_f, (str, list, tuple), str) common.check_type('pt0', pt0, (list, tuple), (int, float)) common.check_type('pt1', pt1, (list, tuple), (int, float)) common.check_type('is_array', is_array, bool) common.check_type('is_overwrite', is_overwrite, bool) pt0 = list( common.convert_indices(fields.ns, pt0) ) pt1 = list( common.convert_indices(fields.ns, pt1) ) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list overwrite_str = {True: '=', False: '+='}[is_overwrite] for strf in str_fs: strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz'] common.check_value('str_f', strf, strf_list) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) # program macros = ['NMAX', 'XID', 'YID', 'ZID', \ 'ARGS', \ 'TARGET', 'SOURCE', 'OVERWRITE', \ 'DTYPE', 'PRAGMA_fp64'] nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1) if is_array: values = [nmax_str, xid_str, yid_str, zid_str, \ '__global DTYPE *source', \ 'target[idx]', 'source[sub_idx]', overwrite_str] + \ dtype_str_list else: values = [nmax_str, xid_str, yid_str, zid_str, \ 'DTYPE source', \ 'target[idx]', 'source', overwrite_str] + \ dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy.cl').read(), macros, values) program = cl.Program(fields.context, ksrc).build() # allocation target_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) if is_array: tmp_array = np.zeros(shape, dtype=fields.dtype) source_buf = cl.Buffer( \ fields.context, \ cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, \ hostbuf=tmp_array) # global variabels and functions self.mainf = fields self.program = program self.target_bufs = target_bufs self.shape = shape nmax = int(nmax_str) remainder = nmax % fields.ls self.gs = nmax if remainder == 0 else nmax - remainder + fields.ls if is_array: self.source_buf = source_buf self.set_fields = self.set_fields_spatial_value else: self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True): """ """ common.check_type('fields', fields, Fields) common.check_type('str_f', str_f, (str, list, tuple), str) common.check_type('pt0', pt0, (list, tuple), (int, float)) common.check_type('pt1', pt1, (list, tuple), (int, float)) common.check_type('is_array', is_array, bool) common.check_type('is_overwrite', is_overwrite, bool) pt0 = list( common.convert_indices(fields.ns, pt0) ) pt1 = list( common.convert_indices(fields.ns, pt1) ) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list overwrite_str = {True: '=', False: '+='}[is_overwrite] for strf in str_fs: strf_list = ['ex', 'ey', 'ez', 'hx', 'hy', 'hz'] common.check_value('str_f', strf, strf_list) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) # program macros = ['NMAX', 'XID', 'YID', 'ZID', \ 'ARGS', \ 'TARGET', 'SOURCE', 'OVERWRITE', \ 'DTYPE', 'PRAGMA_fp64'] nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list(pt0, pt1) if is_array: values = [nmax_str, xid_str, yid_str, zid_str, \ '__global DTYPE *source', \ 'target[idx]', 'source[sub_idx]', overwrite_str] + \ dtype_str_list else: values = [nmax_str, xid_str, yid_str, zid_str, \ 'DTYPE source', \ 'target[idx]', 'source', overwrite_str] + \ dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'copy.cl').read(), macros, values) program = cl.Program(fields.context, ksrc).build() # allocation target_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) if is_array: tmp_array = np.zeros(shape, dtype=fields.dtype) source_buf = cl.Buffer( \ fields.context, \ cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, \ hostbuf=tmp_array) # global variabels and functions self.mainf = fields self.program = program self.target_bufs = target_bufs self.shape = shape nmax = int(nmax_str) remainder = nmax % fields.ls self.gs = nmax if remainder == 0 else nmax - remainder + fields.ls if is_array: self.source_buf = source_buf self.set_fields = self.set_fields_spatial_value else: self.set_fields = self.set_fields_single_value
def __init__(self, fields, str_f, pt0, pt1, is_array=False, is_overwrite=True): """ """ common.check_type("fields", fields, Fields) common.check_type("str_f", str_f, (str, list, tuple), str) common.check_type("pt0", pt0, (list, tuple), int) common.check_type("pt1", pt1, (list, tuple), int) common.check_type("is_array", is_array, bool) common.check_type("is_overwrite", is_overwrite, bool) # local variables str_fs = common.convert_to_tuple(str_f) dtype_str_list = fields.dtype_str_list overwrite_str = {True: "=", False: "+="}[is_overwrite] for strf in str_fs: strf_list = ["ex", "ey", "ez", "hx", "hy", "hz"] common.check_value("str_f", strf, strf_list) for axis, n, p0, p1 in zip(["x", "y", "z"], fields.ns, pt0, pt1): common.check_value("pt0 %s" % axis, p0, range(n)) common.check_value("pt1 %s" % axis, p1, range(n)) # program macros = ["NMAX", "XID", "YID", "ZID", "ARGS", "TARGET", "SOURCE", "OVERWRITE", "DTYPE", "PRAGMA_fp64"] if is_array: values = ( common_gpu.macro_replace_list(pt0, pt1) + ["__global DTYPE *source", "target[idx]", "source[sub_idx]", overwrite_str] + dtype_str_list ) else: values = ( common_gpu.macro_replace_list(pt0, pt1) + ["DTYPE source", "target[idx]", "source", overwrite_str] + dtype_str_list ) ksrc = common.replace_template_code(open(common_gpu.src_path + "copy.cl").read(), macros, values) program = cl.Program(fields.context, ksrc).build() # allocation target_bufs = [fields.get_buf(str_f) for str_f in str_fs] shape = common.shape_two_points(pt0, pt1, len(str_fs)) if is_array: tmp_array = np.zeros(shape, dtype=fields.dtype) source_buf = cl.Buffer( fields.context, cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, hostbuf=tmp_array ) # global variabels and functions self.mainf = fields self.program = program self.target_bufs = target_bufs self.shape = shape if is_array: self.source_buf = source_buf self.set_fields = self.set_fields_spatial_value else: self.set_fields = self.set_fields_single_value
def __init__(self, fields, pt0, pt1, ep_inf, drude_freq, gamma, mask_arrays=(1, 1, 1)): common.check_type('fields', fields, Fields) common.check_type('pt0', pt0, (list, tuple), (int, float)) common.check_type('pt1', pt1, (list, tuple), (int, float)) common.check_type('ep_inf', ep_inf, (int, float)) common.check_type('drude_freq', drude_freq, (int, float)) common.check_type('gamma', gamma, (int, float)) common.check_type('mask_arrays', mask_arrays, (list, tuple), (np.ndarray, int)) # local variables pt0 = common.convert_indices(fields.ns, pt0) pt1 = common.convert_indices(fields.ns, pt1) context = fields.context queue = fields.queue dtype = fields.dtype shape = common.shape_two_points(pt0, pt1, is_dummy=True) for axis, n, p0, p1 in zip(['x', 'y', 'z'], fields.ns, pt0, pt1): common.check_value('pt0 %s' % axis, p0, range(n)) common.check_value('pt1 %s' % axis, p1, range(n)) for mask_array in mask_arrays: if isinstance(mask_array, np.ndarray): assert common.shape_two_points(pt0, pt1) == mask_array.shape, \ 'shape mismatch : %s, %s' % (shape, mask_array.shape) # allocations psis = [np.zeros(shape, dtype) for i in range(3)] psi_bufs = [ cl.Buffer(context, cl.mem_flags.READ_WRITE, psi.nbytes) for psi in psis ] for psi_buf, psi in zip(psi_bufs, psis): cl.enqueue_copy(queue, psi_buf, psi) dt = fields.dt aa = (2 - gamma * dt) / (2 + gamma * dt) bb = drude_freq**2 * dt / (2 + gamma * dt) comm = 2 * ep_inf + bb * dt ca = 2 * dt / comm cb = -(aa + 3) * bb * dt / comm cc = -(aa + 1) * dt / comm cas = [ca * mask for mask in mask_arrays] shape = common.shape_two_points(pt0, pt1, is_dummy=True) f = np.zeros(shape, dtype) psi_bufs = [ cl.Buffer(context, cl.mem_flags.READ_WRITE, f.nbytes) for i in range(3) ] for psi_buf in psi_bufs: cl.enqueue_copy(queue, psi_buf, f) cf = np.ones(shape, dtype) mask_bufs = [ cl.Buffer(context, cl.mem_flags.READ_ONLY, cf.nbytes) for i in range(3) ] for mask_buf, mask in zip(mask_bufs, mask_arrays): cl.enqueue_copy(queue, mask_buf, cf * mask) # modify ce arrays slices = common.slices_two_points(pt0, pt1) for ce, ca in zip(fields.get_ces(), cas): ce[slices] = ca * mask + ce[slices] * mask.__invert__() # program nmax_str, xid_str, yid_str, zid_str = common_gpu.macro_replace_list( pt0, pt1) macros = ['NMAX', 'XID', 'YID', 'ZID', 'DX', 'DTYPE', 'PRAGMA_fp64'] values = [nmax_str, xid_str, yid_str, zid_str, str(fields.ls)] + fields.dtype_str_list ksrc = common.replace_template_code( \ open(common_gpu.src_path + 'drude.cl').read(), macros, values) program = cl.Program(fields.context, ksrc).build() # arguments pca = aa pcb = (aa + 1) * bb args = fields.ns + [dtype(cb), dtype(cc), dtype(pca), dtype(pcb)] \ + fields.eh_bufs[:3] + psi_bufs + mask_bufs # global variables self.mainf = fields self.program = program self.args = args nx, ny, nz = fields.ns nmax = int(nmax_str) remainder = nmax % fields.ls self.gs = nmax if remainder == 0 else nmax - remainder + fields.ls # append to the update list self.priority_type = 'material' fields.append_instance(self)