def visit_Operator(self, o, mode='all'):
# Kernel signature and body
body = flatten(self._visit(i) for i in o.children)
decls = self._args_decl(o.parameters)
signature = c.FunctionDeclaration(c.Value(o.retval, o.name), decls)
retval = [c.Line(), c.Statement("return 0")]
kernel = c.FunctionBody(signature, c.Block(body + retval))
# Elemental functions
esigns = []
efuncs = [blankline]
for i in o._func_table.values():
if i.local:
prefix = ' '.join(i.root.prefix + (i.root.retval, ))
esigns.append(
c.FunctionDeclaration(c.Value(prefix, i.root.name),
self._args_decl(i.root.parameters)))
efuncs.extend([self._visit(i.root), blankline])
# Definitions
headers = [c.Define(*i) for i in o._headers] + [blankline]
# Header files
includes = self._operator_includes(o) + [blankline]
# Type declarations
typedecls = self._operator_typedecls(o, mode)
if mode in ('all', 'public') and o._compiler.src_ext in ('cpp', 'cu'):
typedecls.append(c.Extern('C', signature))
typedecls = [i for j in typedecls for i in (j, blankline)]
return c.Module(headers + includes + typedecls + esigns +
[blankline, kernel] + efuncs)
def visit_Operator(self, o):
blankline = c.Line("")
# Kernel signature and body
body = flatten(self._visit(i) for i in o.children)
decls = self._args_decl(o.parameters)
signature = c.FunctionDeclaration(c.Value(o.retval, o.name), decls)
retval = [c.Statement("return 0")]
kernel = c.FunctionBody(signature, c.Block(body + retval))
# Elemental functions
esigns = []
efuncs = [blankline]
for i in o._func_table.values():
if i.local:
esigns.append(
c.FunctionDeclaration(c.Value(i.root.retval, i.root.name),
self._args_decl(i.root.parameters)))
efuncs.extend([i.root.ccode, blankline])
# Header files, extra definitions, ...
header = [c.Line(i) for i in o._headers]
includes = [c.Include(i, system=False) for i in o._includes]
includes += [blankline]
cdefs = [
i._C_typedecl for i in o.parameters if i._C_typedecl is not None
]
cdefs = filter_sorted(cdefs, key=lambda i: i.tpname)
if o._compiler.src_ext == 'cpp':
cdefs += [c.Extern('C', signature)]
cdefs = [i for j in cdefs for i in (j, blankline)]
return c.Module(header + includes + cdefs + esigns +
[blankline, kernel] + efuncs)
def visit_FunctionDef(self, node):
# Generate "ccode" attribute by traversing the Python AST
for stmt in node.body:
if not (hasattr(stmt, 'value') and type(stmt.value) is ast.Str): # ignore docstrings
self.visit(stmt)
# Create function declaration and argument list
decl = c.Static(c.DeclSpecifier(c.Value("ErrorCode", node.name), spec='inline'))
args = [c.Pointer(c.Value(self.ptype.name, "particle")),
c.Value("double", "time"), c.Value("float", "dt")]
for field_name, field in self.field_args.items():
args += [c.Pointer(c.Value("CField", "%s" % field_name))]
for field_name, field in self.vector_field_args.items():
fieldset = field.fieldset
Wname = field.W.name if field.W else 'not_defined'
for f in [field.U.name, field.V.name, Wname]:
try:
# Next line will break for example if field.U was created but not added to the fieldset
getattr(fieldset, f)
if f not in self.field_args:
args += [c.Pointer(c.Value("CField", "%s" % f))]
except:
if f != Wname:
raise RuntimeError("Field %s needed by a VectorField but it does not exist" % f)
else:
pass
for const, _ in self.const_args.items():
args += [c.Value("float", const)]
# Create function body as C-code object
body = [stmt.ccode for stmt in node.body if not (hasattr(stmt, 'value') and type(stmt.value) is ast.Str)]
body += [c.Statement("return SUCCESS")]
node.ccode = c.FunctionBody(c.FunctionDeclaration(decl, args), c.Block(body))
def _generate_kernel_func(self):
self._components['KERNEL_FUNC'] = cgen.FunctionBody(
cgen.FunctionDeclaration(
cgen.DeclSpecifier(
cgen.Value("void", 'k_' + self._kernel.name), 'inline'),
self._components['KERNEL_ARG_DECLS']),
cgen.Block([cgen.Line(self._kernel.code)]))
def visit_FunctionDef(self, node):
# Generate "ccode" attribute by traversing the Python AST
for stmt in node.body:
if not (hasattr(stmt, 'value') and type(stmt.value) is ast.Str): # ignore docstrings
self.visit(stmt)
# Create function declaration and argument list
decl = c.Static(c.DeclSpecifier(c.Value("ErrorCode", node.name), spec='inline'))
args = [c.Pointer(c.Value(self.ptype.name, "particle")),
c.Value("double", "time")]
for field in self.field_args.values():
args += [c.Pointer(c.Value("CField", "%s" % field.ccode_name))]
for field in self.vector_field_args.values():
for fcomponent in ['U', 'V', 'W']:
try:
f = getattr(field, fcomponent)
if f.ccode_name not in self.field_args:
args += [c.Pointer(c.Value("CField", "%s" % f.ccode_name))]
self.field_args[f.ccode_name] = f
except:
pass # field.W does not always exist
for const, _ in self.const_args.items():
args += [c.Value("float", const)]
# Create function body as C-code object
body = [stmt.ccode for stmt in node.body if not (hasattr(stmt, 'value') and type(stmt.value) is ast.Str)]
body += [c.Statement("return SUCCESS")]
node.ccode = c.FunctionBody(c.FunctionDeclaration(decl, args), c.Block(body))
def visit_FunctionDef(self, node):
# Generate "ccode" attribute by traversing the Python AST
for stmt in node.body:
if not (hasattr(stmt, 'value') and type(stmt.value) is ast.Str): # ignore docstrings
self.visit(stmt)
# Create function declaration and argument list
decl = c.Static(c.DeclSpecifier(c.Value("ErrorCode", node.name), spec='inline'))
args = [c.Pointer(c.Value(self.ptype.name, "particle")),
c.Value("double", "time"), c.Value("float", "dt")]
for field_name, field in self.field_args.items():
if field_name != 'UV':
args += [c.Pointer(c.Value("CField", "%s" % field_name))]
for field_name, field in self.field_args.items():
if field_name == 'UV':
fieldset = field.fieldset
for f in ['U', 'V', 'cosU', 'sinU', 'cosV', 'sinV']:
try:
getattr(fieldset, f)
if f not in self.field_args:
args += [c.Pointer(c.Value("CField", "%s" % f))]
except:
if fieldset.U.grid.gtype in [GridCode.CurvilinearZGrid, GridCode.CurvilinearSGrid]:
raise RuntimeError("cosU, sinU, cosV and sinV fields must be defined for a proper rotation of U, V fields in curvilinear grids")
for const, _ in self.const_args.items():
args += [c.Value("float", const)]
# Create function body as C-code object
body = [stmt.ccode for stmt in node.body if not (hasattr(stmt, 'value') and type(stmt.value) is ast.Str)]
body += [c.Statement("return SUCCESS")]
node.ccode = c.FunctionBody(c.FunctionDeclaration(decl, args), c.Block(body))
def visit_Callable(self, o):
body = flatten(self.visit(i) for i in o.children)
params = runtime_arguments(o.parameters)
decls = self._args_decl(params)
casts = self._args_cast(params)
signature = c.FunctionDeclaration(c.Value(o.retval, o.name), decls)
return c.FunctionBody(signature, c.Block(casts + body))
def _generate_kernel_func(self):
if_block = cgen.If(
self._components['LIB_PAIR_INDEX_0'] + '<_D_N_LOCAL',
cgen.Block([
self._components['KERNEL_GATHER'],
self._components['KERNEL_MAPPING'],
cgen.Line(self._kernel.code),
self._components['KERNEL_SCATTER']
]))
func = cgen.Block([
cgen.Initializer(
cgen.Const(
cgen.Value(host.int32_str,
self._components['LIB_PAIR_INDEX_0'])),
'threadIdx.x + blockIdx.x*blockDim.x'),
self._components['IF_GATHER'], if_block,
self._components['IF_SCATTER']
])
self._components['KERNEL_FUNC'] = cgen.FunctionBody(
cgen.FunctionDeclaration(
cgen.DeclSpecifier(
cgen.Value("void", 'k_' + self._kernel.name),
'__global__'), self._components['KERNEL_ARG_DECLS']), func)
def main(self):
statements = []
statements.append(cgen.Value(self._grid_structure_name, "grid"))
statements.append(cgen.Value(self.__convergence_structure_name,
"conv"))
statements.append(
cgen.Value(self._profiling_structure_name, "profiling"))
statements.append(cgen.Statement("opesci_execute(&grid, &profiling)"))
statements.append(cgen.Statement("opesci_convergence(&grid, &conv)"))
statements.append(cgen.Statement("opesci_free(&grid)"))
statements.append(self.grid.print_convergence)
if self.profiling:
statements.append(
cgen.Statement(
'printf("PAPI:: Max real_time: %f (sec)\\n", profiling.g_rtime)'
))
statements.append(
cgen.Statement(
'printf("PAPI:: Max proc_time: %f (sec)\\n", profiling.g_ptime)'
))
statements.append(
cgen.Statement(
'printf("PAPI:: Total MFlops/s: %f\\n", profiling.g_mflops)'
))
statements.append(cgen.Statement('return 0'))
return cgen.FunctionBody(
cgen.FunctionDeclaration(cgen.Value('int', 'main'), []),
cgen.Block(statements))
def visit_FunctionDef(self, node):
# Generate "ccode" attribute by traversing the Python AST
for stmt in node.body:
if not (hasattr(stmt, 'value')
and type(stmt.value) is ast.Str): # ignore docstrings
self.visit(stmt)
# Create function declaration and argument list
decl = c.Static(
c.DeclSpecifier(c.Value("ErrorCode", node.name), spec='inline'))
args = [
c.Pointer(c.Value(self.ptype.name, "particle")),
c.Value("double", "time")
]
for field in self.field_args.values():
args += [c.Pointer(c.Value("CField", "%s" % field.ccode_name))]
for field in self.vector_field_args.values():
Wname = field.W.ccode_name if field.W else 'not_defined'
for fname in [field.U.ccode_name, field.V.ccode_name, Wname]:
if fname not in self.field_args and fname != 'not_defined':
args += [c.Pointer(c.Value("CField", "%s" % fname))]
for const, _ in self.const_args.items():
args += [c.Value("float", const)]
# Create function body as C-code object
body = [
stmt.ccode for stmt in node.body
if not (hasattr(stmt, 'value') and type(stmt.value) is ast.Str)
]
body += [c.Statement("return SUCCESS")]
node.ccode = c.FunctionBody(c.FunctionDeclaration(decl, args),
c.Block(body))
def execute_function_signature(self):
return cgen.Extern(
"C",
cgen.FunctionDeclaration(cgen.Value('int', 'opesci_execute'), [
cgen.Pointer(cgen.Value(self._grid_structure_name, "grid")),
cgen.Pointer(
cgen.Value(self._profiling_structure_name, "profiling"))
]))
def copy_initialise_function(self):
statements = cgen.Block(self.grid.copy_memory)
return cgen.FunctionBody(
cgen.Extern(
"C",
cgen.FunctionDeclaration(
cgen.Value('void', 'copy_initialise'),
[cgen.Pointer(cgen.Value(self.grid.real_t, 'data'))])),
statements)
def generate(self, funcname, field_args, const_args, kernel_ast, c_include):
ccode = []
# Add include for Parcels and math header
ccode += [str(c.Include("parcels.h", system=False))]
ccode += [str(c.Include("math.h", system=False))]
# Generate type definition for particle type
vdecl = []
for v in self.ptype.variables:
if v.dtype == np.uint64:
vdecl.append(c.Pointer(c.POD(np.void, v.name)))
else:
vdecl.append(c.POD(v.dtype, v.name))
ccode += [str(c.Typedef(c.GenerableStruct("", vdecl, declname=self.ptype.name)))]
if c_include:
ccode += [c_include]
# Insert kernel code
ccode += [str(kernel_ast)]
# Generate outer loop for repeated kernel invocation
args = [c.Value("int", "num_particles"),
c.Pointer(c.Value(self.ptype.name, "particles")),
c.Value("double", "endtime"), c.Value("float", "dt")]
for field, _ in field_args.items():
args += [c.Pointer(c.Value("CField", "%s" % field))]
for const, _ in const_args.items():
args += [c.Value("float", const)]
fargs_str = ", ".join(['particles[p].time', 'sign_dt * __dt'] + list(field_args.keys())
+ list(const_args.keys()))
# Inner loop nest for forward runs
sign_dt = c.Assign("sign_dt", "dt > 0 ? 1 : -1")
sign_end_part = c.Assign("sign_end_part", "endtime - particles[p].time > 0 ? 1 : -1")
dt_pos = c.Assign("__dt", "fmin(fabs(particles[p].dt), fabs(endtime - particles[p].time))")
dt_0_break = c.If("particles[p].dt == 0", c.Statement("break"))
notstarted_continue = c.If("(sign_end_part != sign_dt) && (particles[p].dt != 0)",
c.Statement("continue"))
body = [c.Assign("res", "%s(&(particles[p]), %s)" % (funcname, fargs_str))]
body += [c.Assign("particles[p].state", "res")] # Store return code on particle
body += [c.If("res == SUCCESS", c.Block([c.Statement("particles[p].time += sign_dt * __dt"),
dt_pos, dt_0_break, c.Statement("continue")]))]
body += [c.If("res == REPEAT", c.Block([dt_pos, c.Statement("continue")]),
c.Statement("break"))]
time_loop = c.While("__dt > __tol || particles[p].dt == 0", c.Block(body))
part_loop = c.For("p = 0", "p < num_particles", "++p",
c.Block([sign_end_part, notstarted_continue, dt_pos, time_loop]))
fbody = c.Block([c.Value("int", "p, sign_dt, sign_end_part"), c.Value("ErrorCode", "res"),
c.Value("double", "__dt, __tol"), c.Assign("__tol", "1.e-6"),
sign_dt, part_loop])
fdecl = c.FunctionDeclaration(c.Value("void", "particle_loop"), args)
ccode += [str(c.FunctionBody(fdecl, fbody))]
return "\n\n".join(ccode)
def freemem(self):
statements = []
statements.append(self.grid.free_memory)
statements.append(cgen.Statement("return 0"))
return cgen.FunctionBody(
cgen.Extern(
"C",
cgen.FunctionDeclaration(cgen.Value('int', 'opesci_free'), [
cgen.Pointer(cgen.Value(self._grid_structure_name, "grid"))
])), cgen.Block(statements))
def generate_optimmat_code(self, pos, name=None):
"""Generates the code for computing the local optimization matrix
for the optimization over site nr. `pos`
The function has the following signature:
DTYPE const *const A,
DTYPE const *const X_0,
...,
DTYPE const *const X_N,
DTYPE *const result
:param pos: The local tensor to copy (should be `< len(X)`)
:param name: Name of the C function (default: get_optimmat_%(pos))
:returns: cgen.FunctionBody with given name
"""
name = 'get_optimmat_%i' % pos if name is None else name
finalization_src = '''
if (mid < {nr_meas:d}) {{
for (uint i = 0; i < {pdim:d}; ++i) {{
for (uint k_l = 0; k_l < {rank_l:d}; ++k_l) {{
for (uint k_r = 0; k_r < {rank_r:d}; ++k_r) {{
result[mid * {rank_l:d} * {pdim:d} * {rank_r:d}
+ k_l * {pdim:d} * {rank_r:d}
+ i * {rank_r:d}
+ k_r]
= left_c[k_l] * current_row[{offset:d} + i] * right_c[k_r];
}}
}}
}}
}}
'''.format(nr_meas=self._meas,
pdim=self._dims[pos],
rank_l=1 if pos == 0 else self._ranks[pos - 1],
rank_r=1 if pos == self._sites - 1 else self._ranks[pos],
offset=sum(self._dims[:pos]))
finalization = c.LiteralLines(finalization_src)
arg_decls = [ConstPointerToConstDecl(self._dtype, 'A')]
arg_decls += [
ConstPointerToConstDecl(self._dtype, 'X%i' % i)
for i in range(self._sites)
]
arg_decls += [c.Pointer(c.Const(c.POD(self._dtype, 'result')))]
return c.FunctionBody(
ccu.CudaGlobal(
c.FunctionDeclaration(c.Value('void', 'get_optimmat_%i' % pos),
arg_decls=arg_decls)),
c.Block(
self.declaration(pos) + self.left_contractions(pos) +
self.right_contractions(pos) + [finalization]))
def _generate_lib_func(self):
block = cgen.Block([
self.loop_timer.get_cpp_pre_loop_code_ast(),
self._components['LIB_OUTER_LOOP'],
self.loop_timer.get_cpp_post_loop_code_ast()
])
self._components['LIB_FUNC'] = cgen.FunctionBody(
cgen.FunctionDeclaration(
cgen.Value("void", self._components['LIB_NAME']),
self._components['LIB_ARG_DECLS'] +
self._components['KERNEL_LIB_ARG_DECLS']), block)
def __generate_string_methods(self):
declaration = c.FunctionDeclaration(c.Value('void', 'pack_string'), [
self._data_object_ref(self._packet_type(), 'packet'),
self._data_cref(self.get_dtype('string'), 'str')
])
body = c.FunctionBody(
declaration,
c.Block([
c.Statement('uint8_t size = str.size()'),
c.Statement(f'*packet << static_cast<uint8_t>(size)'),
c.For(
'uint8_t i = 0', 'i < size', '++i',
c.Block([
c.Statement(f'*packet << static_cast<uint8_t>(str[i])')
]))
]))
self.marshal.header.append(declaration)
self.marshal.source.append(body)
declaration = c.FunctionDeclaration(c.Value('bool', 'unpack_string'), [
self._pointer_type('PacketReader', 'packet'),
self._data_object_ref(self.get_dtype('string'), 'str')
])
body = c.FunctionBody(
declaration,
c.Block([
WrapUnpack._guard('sizeof(uint8_t)'),
c.Statement('uint8_t size = packet->peek<uint8_t>()'),
WrapUnpack._guard('size'),
c.Statement(f'str = packet->read<std::string>()'),
c.Statement('return true')
]))
self.marshal.header.append(declaration)
self.marshal.source.append(body)
def visit_FunctionDef(self, node):
# Generate "ccode" attribute by traversing the Python AST
for stmt in node.body:
self.visit(stmt)
# Create function declaration and argument list
decl = c.Static(c.DeclSpecifier(c.Value("KernelOp", node.name), spec='inline'))
args = [c.Pointer(c.Value(self.ptype.name, "particle")),
c.Value("double", "time"), c.Value("float", "dt")]
for field, _ in self.field_args.items():
args += [c.Pointer(c.Value("CField", "%s" % field))]
# Create function body as C-code object
body = [stmt.ccode for stmt in node.body]
body += [c.Statement("return SUCCESS")]
node.ccode = c.FunctionBody(c.FunctionDeclaration(decl, args), c.Block(body))
def _cgen(self):
arg_decls = []
for arg in self.func.arg_dict:
# print the variable type of input and output arrays (CPointer) as
# 'void *', so as to handle it using ctypes.c_void_p()
if isinstance(arg.typ, CPointer) and self.are_io_void_ptrs:
arg_decls.append(cgen.Value('void *', arg.__str__()))
else:
arg_decls.append(
cgen.Value(self.func.arg_dict[arg].__str__(),
arg.__str__()))
type_str = self.func.ret_typ.__str__()
if self.is_extern_c_func:
type_str = "extern \"C\" " + type_str
return cgen.FunctionDeclaration(cgen.Value(type_str, self.func.name),
arg_decls)
def convergence_function(self):
statements = []
statements.append(self.grid.define_constants)
statements.append(self.grid.load_fields)
statements.append(self.grid.converge_test)
statements.append(cgen.Statement("return 0"))
return cgen.FunctionBody(
cgen.Extern(
"C",
cgen.FunctionDeclaration(
cgen.Value('int', 'opesci_convergence'), [
cgen.Pointer(
cgen.Value(self._grid_structure_name, "grid")),
cgen.Pointer(
cgen.Value(self.__convergence_structure_name,
"conv"))
])), cgen.Block(statements))
def __init__(self, dtype, name, args, body, template: tuple = None):
self.name = name
self.is_template = template is not None
if self.is_template:
self.name = f'{self.name}<{template[1]}>'
self.value = c.Value(dtype, self.name)
self.decl = c.FunctionDeclaration(self.value, args)
if body is None:
self.fnc = None
else:
self.fnc = c.FunctionBody(self.decl, body)
self.modifiers = []
self.constraints = []
def visit_Operator(self, o):
# Generate the code for the cfile
ccode = super().visit_Operator(o, mode='private')
# Generate the code for the hfile
typedecls = self._operator_typedecls(o, mode='public')
guarded_typedecls = []
for i in typedecls:
guard = "DEVITO_%s" % i.tpname.upper()
iflines = [c.Define(guard, ""), blankline, i, blankline]
guarded_typedecl = c.IfNDef(guard, iflines, [])
guarded_typedecls.extend([guarded_typedecl, blankline])
decls = self._args_decl(o.parameters)
signature = c.FunctionDeclaration(c.Value(o.retval, o.name), decls)
hcode = c.Module(guarded_typedecls + [blankline, signature, blankline])
return ccode, hcode
def __init__(self,
dtype,
name,
args,
body,
template: tuple = None,
modifiers=''):
self.dtype = dtype
self.name = name
self.value = c.Value(dtype, name)
self.modifiers = modifiers
self.is_template = template is not None
decl = c.FunctionDeclaration(self.value, args)
if self.is_template:
decl = c.Template(template[0], decl)
self.fnc = c.FunctionBody(decl, body)
def map_Subroutine(self, node):
assert not node.prefix
assert not hasattr(node, "suffix")
scope = Scope(node.name, list(node.args))
self.scope_stack.append(scope)
body = self.map_statement_list(node.content)
pre_func_decl, in_func_decl = self.get_declarations()
body = in_func_decl + [cgen.Line()] + body
if isinstance(body[-1], cgen.Statement) and body[-1].text == "return":
body.pop()
def get_arg_decl(arg_idx, arg_name):
decl = self.get_declarator(arg_name)
if self.arg_needs_pointer(node.name, arg_idx):
hint = self.addr_space_hints.get((node.name, arg_name))
if hint:
decl = hint(cgen.Pointer(decl))
else:
if self.use_restrict_pointers:
decl = cgen.RestrictPointer(decl)
else:
decl = cgen.Pointer(decl)
return decl
result = cgen.FunctionBody(
cgen.FunctionDeclaration(
cgen.Value("void", node.name),
[get_arg_decl(i, arg) for i, arg in enumerate(node.args)]),
cgen.Block(body))
self.scope_stack.pop()
if pre_func_decl:
return pre_func_decl + [cgen.Line(), result]
else:
return result
def generate(self, funcname, field_args, kernel_ast, adaptive=False):
ccode = []
# Add include for Parcels and math header
ccode += [str(c.Include("parcels.h", system=False))]
ccode += [str(c.Include("math.h", system=False))]
# Generate type definition for particle type
vdecl = [c.POD(dtype, var) for var, dtype in self.ptype.var_types.items()]
ccode += [str(c.Typedef(c.GenerableStruct("", vdecl, declname=self.ptype.name)))]
# Insert kernel code
ccode += [str(kernel_ast)]
# Generate outer loop for repeated kernel invocation
args = [c.Value("int", "num_particles"),
c.Pointer(c.Value(self.ptype.name, "particles")),
c.Value("double", "endtime"), c.Value("float", "dt")]
for field, _ in field_args.items():
args += [c.Pointer(c.Value("CField", "%s" % field))]
fargs_str = ", ".join(['particles[p].time', 'particles[p].dt'] + list(field_args.keys()))
# Inner loop nest for forward runs
dt_fwd = c.Statement("__dt = fmin(particles[p].dt, endtime - particles[p].time)")
body_fwd = [c.Statement("res = %s(&(particles[p]), %s)" % (funcname, fargs_str)),
c.If("res == SUCCESS", c.Statement("particles[p].time += __dt")), dt_fwd]
time_fwd = c.While("__dt > __tol", c.Block(body_fwd))
part_fwd = c.For("p = 0", "p < num_particles", "++p", c.Block([dt_fwd, time_fwd]))
# Inner loop nest for backward runs
dt_bwd = c.Statement("__dt = fmax(particles[p].dt, endtime - particles[p].time)")
body_bwd = [c.Statement("res = %s(&(particles[p]), %s)" % (funcname, fargs_str)),
c.If("res == SUCCESS", c.Statement("particles[p].time += __dt")), dt_bwd]
time_bwd = c.While("__dt < -1. * __tol", c.Block(body_bwd))
part_bwd = c.For("p = 0", "p < num_particles", "++p", c.Block([dt_bwd, time_bwd]))
time_if = c.If("dt > 0.0", c.Block([part_fwd]), c.Block([part_bwd]))
fbody = c.Block([c.Value("int", "p"), c.Value("KernelOp", "res"),
c.Value("double", "__dt, __tol"), c.Assign("__tol", "1.e-6"),
time_if])
fdecl = c.FunctionDeclaration(c.Value("void", "particle_loop"), args)
ccode += [str(c.FunctionBody(fdecl, fbody))]
return "\n\n".join(ccode)
def visit_Operator(self, o):
# Kernel signature and body
body = flatten(self.visit(i) for i in o.children)
params = o.parameters
decls = self._args_decl(params)
signature = c.FunctionDeclaration(c.Value(o.retval, o.name), decls)
retval = [c.Statement("return 0")]
kernel = c.FunctionBody(signature, c.Block(body + retval))
# Elemental functions
efuncs = [i.root.ccode
for i in o.func_table.values() if i.local] + [blankline]
# Header files, extra definitions, ...
header = [c.Line(i) for i in o._headers]
includes = [c.Include(i, system=False) for i in o._includes]
includes += [blankline]
cglobals = list(o._globals)
if o._compiler.src_ext == 'cpp':
cglobals += [c.Extern('C', signature)]
cglobals = [i for j in cglobals for i in (j, blankline)]
return c.Module(header + includes + cglobals + efuncs + [kernel])
def ccode(self):
"""Returns the C code generated by this kernel.
This function generates the internal code block from Iteration
and Expression objects, and adds the necessary template code
around it.
"""
header_vars = [
c.Pointer(c.POD(v.dtype, '%s_vec' % v.name))
for v in self.signature
]
header = c.Extern(
"C", c.FunctionDeclaration(c.Value('int', self.name), header_vars))
cast_shapes = [(v, ''.join(['[%d]' % d for d in v.shape[1:]]))
for v in self.signature]
casts = [
c.Initializer(
c.POD(v.dtype, '(*%s)%s' % (v.name, shape)), '(%s (*)%s) %s' %
(c.dtype_to_ctype(v.dtype), shape, '%s_vec' % v.name))
for v, shape in cast_shapes
]
body = [e.ccode for e in self.expressions]
ret = [c.Statement("return 0")]
return c.FunctionBody(header, c.Block(casts + body + ret))
def visit_Callable(self, o):
body = flatten(self._visit(i) for i in o.children)
decls = self._args_decl(o.parameters)
signature = c.FunctionDeclaration(c.Value(o.retval, o.name), decls)
return c.FunctionBody(signature, c.Block(body))
def generate(self, funcname, field_args, const_args, kernel_ast,
c_include):
ccode = []
pname = self.ptype.name + 'p'
# ==== Add include for Parcels and math header ==== #
ccode += [str(c.Include("parcels.h", system=False))]
#ccode += [str(c.Include("math.h", system=False))] # removed by Lyc because it is already in parcels.h ???
#ccode += [str(c.Include("stdbool.h", system=False))] # added by Luc to accomodate crossdike.h booleans
ccode += [str(c.Assign('double _next_dt', '0'))]
ccode += [str(c.Assign('size_t _next_dt_set', '0'))]
ccode += [
str(
c.Assign(
'const int ngrid',
str(self.fieldset.gridset.size if self.
fieldset is not None else 1)))
]
# ==== Generate type definition for particle type ==== #
vdeclp = [
c.Pointer(c.POD(v.dtype, v.name)) for v in self.ptype.variables
]
ccode += [
str(c.Typedef(c.GenerableStruct("", vdeclp, declname=pname)))
]
# Generate type definition for single particle type
vdecl = [
c.POD(v.dtype, v.name) for v in self.ptype.variables
if v.dtype != np.uint64
]
ccode += [
str(
c.Typedef(
c.GenerableStruct("", vdecl, declname=self.ptype.name)))
]
args = [
c.Pointer(c.Value(self.ptype.name, "particle_backup")),
c.Pointer(c.Value(pname, "particles")),
c.Value("int", "pnum")
]
p_back_set_decl = c.FunctionDeclaration(
c.Static(
c.DeclSpecifier(c.Value("void", "set_particle_backup"),
spec='inline')), args)
body = []
for v in self.ptype.variables:
if v.dtype != np.uint64 and v.name not in ['dt', 'state']:
body += [
c.Assign(("particle_backup->%s" % v.name),
("particles->%s[pnum]" % v.name))
]
p_back_set_body = c.Block(body)
p_back_set = str(c.FunctionBody(p_back_set_decl, p_back_set_body))
ccode += [p_back_set]
args = [
c.Pointer(c.Value(self.ptype.name, "particle_backup")),
c.Pointer(c.Value(pname, "particles")),
c.Value("int", "pnum")
]
p_back_get_decl = c.FunctionDeclaration(
c.Static(
c.DeclSpecifier(c.Value("void", "get_particle_backup"),
spec='inline')), args)
body = []
for v in self.ptype.variables:
if v.dtype != np.uint64 and v.name not in ['dt', 'state']:
body += [
c.Assign(("particles->%s[pnum]" % v.name),
("particle_backup->%s" % v.name))
]
p_back_get_body = c.Block(body)
p_back_get = str(c.FunctionBody(p_back_get_decl, p_back_get_body))
ccode += [p_back_get]
update_next_dt_decl = c.FunctionDeclaration(
c.Static(
c.DeclSpecifier(c.Value("void", "update_next_dt"),
spec='inline')), [c.Value('double', 'dt')])
if 'update_next_dt' in str(kernel_ast):
body = []
body += [c.Assign("_next_dt", "dt")]
body += [c.Assign("_next_dt_set", "1")]
update_next_dt_body = c.Block(body)
update_next_dt = str(
c.FunctionBody(update_next_dt_decl, update_next_dt_body))
ccode += [update_next_dt]
if c_include:
ccode += [c_include]
# ==== Insert kernel code ==== #
ccode += [str(kernel_ast)]
# Generate outer loop for repeated kernel invocation
args = [
c.Value("int", "num_particles"),
c.Pointer(c.Value(pname, "particles")),
c.Value("double", "endtime"),
c.Value("double", "dt")
]
for field, _ in field_args.items():
args += [c.Pointer(c.Value("CField", "%s" % field))]
for const, _ in const_args.items():
args += [c.Value("double", const)]
fargs_str = ", ".join(['particles->time[pnum]'] +
list(field_args.keys()) +
list(const_args.keys()))
# ==== statement clusters use to compose 'body' variable and variables 'time_loop' and 'part_loop' ==== ##
sign_dt = c.Assign("sign_dt", "dt > 0 ? 1 : -1")
particle_backup = c.Statement("%s particle_backup" % self.ptype.name)
sign_end_part = c.Assign(
"sign_end_part", "(endtime - particles->time[pnum]) > 0 ? 1 : -1")
reset_res_state = c.Assign("res", "particles->state[pnum]")
update_state = c.Assign("particles->state[pnum]", "res")
update_pdt = c.If(
"_next_dt_set == 1",
c.Block([
c.Assign("_next_dt_set", "0"),
c.Assign("particles->dt[pnum]", "_next_dt")
]))
dt_pos = c.Assign(
"__dt",
"fmin(fabs(particles->dt[pnum]), fabs(endtime - particles->time[pnum]))"
) # original
pdt_eq_dt_pos = c.Assign("__pdt_prekernels", "__dt * sign_dt")
partdt = c.Assign("particles->dt[pnum]", "__pdt_prekernels")
check_pdt = c.If(
"(res == SUCCESS) & !is_equal_dbl(__pdt_prekernels, particles->dt[pnum])",
c.Assign("res", "REPEAT"))
dt_0_break = c.If("is_zero_dbl(particles->dt[pnum])",
c.Statement("break"))
notstarted_continue = c.If(
"(( sign_end_part != sign_dt) || is_close_dbl(__dt, 0) ) && !is_zero_dbl(particles->dt[pnum])",
c.Block([
c.If("fabs(particles->time[pnum]) >= fabs(endtime)",
c.Assign("particles->state[pnum]", "SUCCESS")),
c.Statement("continue")
]))
# ==== main computation body ==== #
body = [
c.Statement(
"set_particle_backup(&particle_backup, particles, pnum)")
]
body += [pdt_eq_dt_pos]
body += [partdt]
body += [
c.Value("StatusCode", "state_prev"),
c.Assign("state_prev", "particles->state[pnum]")
]
body += [
c.Assign("res", "%s(particles, pnum, %s)" % (funcname, fargs_str))
]
body += [
c.If("(res==SUCCESS) && (particles->state[pnum] != state_prev)",
c.Assign("res", "particles->state[pnum]"))
]
body += [check_pdt]
body += [
c.If(
"res == SUCCESS || res == DELETE",
c.Block([
c.Statement(
"particles->time[pnum] += particles->dt[pnum]"),
update_pdt, dt_pos, sign_end_part,
c.If(
"(res != DELETE) && !is_close_dbl(__dt, 0) && (sign_dt == sign_end_part)",
c.Assign("res", "EVALUATE")),
c.If("sign_dt != sign_end_part",
c.Assign("__dt", "0")), update_state, dt_0_break
]),
c.Block([
c.Statement(
"get_particle_backup(&particle_backup, particles, pnum)"
), dt_pos, sign_end_part,
c.If("sign_dt != sign_end_part", c.Assign("__dt", "0")),
update_state,
c.Statement("break")
]))
]
time_loop = c.While(
"(particles->state[pnum] == EVALUATE || particles->state[pnum] == REPEAT) || is_zero_dbl(particles->dt[pnum])",
c.Block(body))
part_loop = c.For(
"pnum = 0", "pnum < num_particles", "++pnum",
c.Block([
sign_end_part, reset_res_state, dt_pos, notstarted_continue,
time_loop
]))
fbody = c.Block([
c.Value("int", "pnum, sign_dt, sign_end_part"),
c.Value("StatusCode", "res"),
c.Value("double", "__pdt_prekernels"),
c.Value("double",
"__dt"), # 1e-8 = built-in tolerance for np.isclose()
sign_dt,
particle_backup,
part_loop
])
fdecl = c.FunctionDeclaration(c.Value("void", "particle_loop"), args)
ccode += [str(c.FunctionBody(fdecl, fbody))]
return "\n\n".join(ccode)
import cgen as c
func = c.FunctionBody(
c.FunctionDeclaration(c.Const(c.Pointer(c.Value("char", "greet"))), []),
c.Block([c.Statement('return "hello world"')]))
code = c.Module([])
code.append(c.Value('int', 'cont'))
code.append(c.Assign('cont', '0'))
code.append(c.Increment('cont', '5'))
print(code)
