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 )