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_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 _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"), 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):
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 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 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 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)))]
args = [c.Pointer(c.Value(self.ptype.name, "particle_backup")),
c.Pointer(c.Value(self.ptype.name, "particle"))]
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), ("particle->%s" % 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(self.ptype.name, "particle"))]
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(("particle->%s" % 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]
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'] + list(field_args.keys())
+ list(const_args.keys()))
# Inner loop nest for forward runs
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[p].time > 0 ? 1 : -1")
dt_pos = c.Assign("__dt", "fmin(fabs(particles[p].dt), fabs(endtime - particles[p].time))")
pdt_eq_dt_pos = c.Assign("particles[p].dt", "__dt * sign_dt")
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.Statement("set_particle_backup(&particle_backup, &(particles[p]))")]
body += [pdt_eq_dt_pos]
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([c.Statement("get_particle_backup(&particle_backup, &(particles[p]))"),
dt_pos, c.Statement("break")]), 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, particle_backup, part_loop])
fdecl = c.FunctionDeclaration(c.Value("void", "particle_loop"), args)
ccode += [str(c.FunctionBody(fdecl, fbody))]
return "\n\n".join(ccode)
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)
def _generate_kernel_func(self):
IX = self._components['LIB_PAIR_INDEX_0']
IY = self._components['LIB_PAIR_INDEX_1']
CX = '_CX'
CY = '_CY'
if_block = cgen.If(
IX + '<_D_N_LOCAL',
cgen.Block([
self._components['KERNEL_GATHER'],
cgen.Initializer(cgen.Const(cgen.Value(host.int32_str, CX)), '_D_CRL[' + IX +']'),
cgen.For('int _jk=0','_jk<_D_CCC['+CX+']', '_jk++',
cgen.Block([
cgen.Initializer(
cgen.Const(cgen.Value(host.int32_str, IY)),
'_D_L_MATRIX[' + CX+'*_D_N_LAYERS' + '+_jk]'
),
cgen.If(
IX+'!='+IY,
cgen.Block([
self._components['KERNEL_MAPPING'],
cgen.Line(self._kernel.code)
])
),
])
),
cgen.For('int _k=0','_k<_D_N_OFFSETS', '_k++',
cgen.Block([
cgen.Initializer(cgen.Const(cgen.Value(host.int32_str, CY)), CX + '+ _D_OFFSETS[_k]'),
cgen.For('int _jk=0','_jk<_D_CCC['+CY+']', '_jk++',
cgen.Block([
cgen.Initializer(
cgen.Const(cgen.Value(host.int32_str, IY)),
'_D_L_MATRIX[' + CY+'*_D_N_LAYERS' + '+_jk]'
),
#cgen.If(IX+'!='+IY,
#cgen.Block([
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
)
