def _arg_names_and_decls(self, codegen_state): implemented_data_info = codegen_state.implemented_data_info arg_names = [iai.name for iai in implemented_data_info] arg_decls = [ self.idi_to_cgen_declarator(codegen_state.kernel, idi) for idi in implemented_data_info ] # {{{ occa compatibility hackery from cgen import Value if self.target.occa_mode: from cgen import ArrayOf, Const from cgen.ispc import ISPCUniform arg_decls = [ Const(ISPCUniform(ArrayOf(Value("int", "loopy_dims")))), Const(ISPCUniform(Value("int", "o1"))), Const(ISPCUniform(Value("int", "o2"))), Const(ISPCUniform(Value("int", "o3"))), ] + arg_decls arg_names = ["loopy_dims", "o1", "o2", "o3"] + arg_names # }}} return arg_names, arg_decls
def make_codepy_module(self, toolchain, dtype): from codepy.libraries import add_codepy toolchain = toolchain.copy() add_codepy(toolchain) from cgen import (Value, Include, Statement, Typedef, FunctionBody, FunctionDeclaration, Block, Const, Line, POD, Initializer, CustomLoop) S = Statement from codepy.bpl import BoostPythonModule mod = BoostPythonModule() mod.add_to_preamble([ Include("vector"), Include("algorithm"), Include("hedge/base.hpp"), Include("hedge/volume_operators.hpp"), Include("boost/foreach.hpp"), Include("boost/numeric/ublas/io.hpp"), ]+self.get_cpu_extra_includes()) mod.add_to_module([ S("namespace ublas = boost::numeric::ublas"), S("using namespace hedge"), S("using namespace pyublas"), Line(), Typedef(POD(dtype, "value_type")), Line(), ]) mod.add_function(FunctionBody( FunctionDeclaration(Value("void", "process_elements"), [ Const(Value("uniform_element_ranges", "ers")), Const(Value("numpy_vector<value_type>", "field")), Value("numpy_vector<value_type>", "result"), ]+self.get_cpu_extra_parameter_declarators()), Block([ Typedef(Value("numpy_vector<value_type>::iterator", "it_type")), Typedef(Value("numpy_vector<value_type>::const_iterator", "cit_type")), Line(), Initializer(Value("it_type", "result_it"), "result.begin()"), Initializer(Value("cit_type", "field_it"), "field.begin()"), Line() ]+self.get_cpu_extra_preamble()+[ Line(), CustomLoop( "BOOST_FOREACH(const element_range er, ers)", Block(self.get_cpu_per_element_code()) ) ]))) #print mod.generate() #toolchain = toolchain.copy() #toolchain.enable_debugging return mod.compile(toolchain)
def emit_initializer(self, codegen_state, dtype, name, val_str, is_const): decl = POD(self, dtype, name) from cgen import Initializer, Const if is_const: decl = Const(decl) return Initializer(decl, val_str)
def get_global_arg_decl(self, name, shape, dtype, is_written): from cgen import RestrictPointer, Const arg_decl = RestrictPointer(POD(self, dtype, name)) if not is_written: arg_decl = Const(arg_decl) return arg_decl
def get_array_arg_decl(self, name, mem_address_space, shape, dtype, is_written): from cgen import RestrictPointer, Const arg_decl = RestrictPointer(POD(self, dtype, name)) if not is_written: arg_decl = Const(arg_decl) return arg_decl
def get_constant_arg_decl(self, name, shape, dtype, is_written): from loopy.target.c import POD # uses the correct complex type from cgen import RestrictPointer, Const arg_decl = RestrictPointer(POD(self, dtype, name)) if not is_written: arg_decl = Const(arg_decl) return arg_decl
def get_global_arg_decl(self, name, shape, dtype, is_written): from loopy.target.c import POD # uses the correct complex type from cgen import Const from cgen.cuda import CudaRestrictPointer arg_decl = CudaRestrictPointer(POD(self, dtype, name)) if not is_written: arg_decl = Const(arg_decl) return arg_decl
def generate_top_of_body(self, codegen_state): from loopy.kernel.data import ImageArg if any(isinstance(arg, ImageArg) for arg in codegen_state.kernel.args): from cgen import Value, Const, Initializer return [ Initializer(Const(Value("sampler_t", "loopy_sampler")), "CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP " "| CLK_FILTER_NEAREST") ] return []
def get_global_arg_decl(self, name, shape, dtype, is_written): from loopy.target.c import POD # uses the correct complex type from cgen import Const from cgen.ispc import ISPCUniformPointer, ISPCUniform arg_decl = ISPCUniformPointer(POD(self, dtype, name)) if not is_written: arg_decl = Const(arg_decl) arg_decl = ISPCUniform(arg_decl) return arg_decl
def get_temporary_decl(self, knl, schedule_index, temp_var, decl_info): temp_var_decl = POD(self, decl_info.dtype, decl_info.name) if temp_var.read_only: from cgen import Const temp_var_decl = Const(temp_var_decl) if decl_info.shape: from cgen import ArrayOf temp_var_decl = ArrayOf( temp_var_decl, " * ".join(str(s) for s in decl_info.shape)) return temp_var_decl
def get_value_arg_decl(self, name, shape, dtype, is_written): assert shape == () result = POD(self, dtype, name) if not is_written: from cgen import Const result = Const(result) if self.target.fortran_abi: from cgen import Pointer result = Pointer(result) return result
def make_greet_mod(greeting): from cgen import FunctionBody, FunctionDeclaration, Block, \ Const, Pointer, Value, Statement from codepy.bpl import BoostPythonModule mod = BoostPythonModule() mod.add_function( FunctionBody( FunctionDeclaration(Const(Pointer(Value("char", "greet"))), []), Block([Statement('return "%s"' % greeting)]))) from codepy.toolchain import guess_toolchain return mod.compile(guess_toolchain(), wait_on_error=True)
def get_temporary_decl(self, codegen_state, schedule_index, temp_var, decl_info): temp_var_decl = POD(self, decl_info.dtype, decl_info.name) if temp_var.read_only: from cgen import Const temp_var_decl = Const(temp_var_decl) if decl_info.shape: from cgen import ArrayOf ecm = self.get_expression_to_code_mapper(codegen_state) temp_var_decl = ArrayOf(temp_var_decl, ecm(p.flattened_product(decl_info.shape), prec=PREC_NONE, type_context="i")) return temp_var_decl
def get_value_arg_decl(self, name, shape, dtype, is_written): result = super().get_value_arg_decl(name, shape, dtype, is_written) from cgen import Reference, Const was_const = isinstance(result, Const) if was_const: result = result.subdecl if self.target.occa_mode: result = Reference(result) if was_const: result = Const(result) from cgen.ispc import ISPCUniform return ISPCUniform(result)
def idi_to_cgen_declarator(self, kernel, idi): from loopy.kernel.data import InameArg if (idi.offset_for_name is not None or idi.stride_for_name_and_axis is not None): assert not idi.is_written from cgen import Const return Const(POD(self, idi.dtype, idi.name)) elif issubclass(idi.arg_class, InameArg): return InameArg(idi.name, idi.dtype).get_arg_decl(self) else: name = idi.base_name or idi.name var_descr = kernel.get_var_descriptor(name) from loopy.kernel.data import ArrayBase if isinstance(var_descr, ArrayBase): return var_descr.get_arg_decl(self, idi.name[len(name):], idi.shape, idi.dtype, idi.is_written) else: return var_descr.get_arg_decl(self)
a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) c_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes) from cgen import FunctionBody, \ FunctionDeclaration, Typedef, POD, Value, \ Pointer, Module, Block, Initializer, Assign, Const from cgen.opencl import CLKernel, CLGlobal, \ CLRequiredWorkGroupSize mod = Module([ FunctionBody( CLKernel(CLRequiredWorkGroupSize((local_size,), FunctionDeclaration( Value("void", "add"), arg_decls=[CLGlobal(Pointer(Const(POD(dtype, name)))) for name in ["tgt", "op1", "op2"]]))), Block([ Initializer(POD(numpy.int32, "idx"), "get_local_id(0) + %d * get_group_id(0)" % (local_size*thread_strides)) ]+[ Assign( "tgt[idx+%d]" % (o*local_size), "op1[idx+%d] + op2[idx+%d]" % ( o*local_size, o*local_size)) for o in range(thread_strides)]))]) knl = cl.Program(ctx, str(mod)).build().add
def get_kernel(self, diff_op, elgroup, for_benchmark=False): from cgen import \ Pointer, POD, Value, ArrayOf, Const, \ Module, FunctionDeclaration, FunctionBody, Block, \ Comment, Line, Define, Include, \ Initializer, If, For, Statement, Assign from pycuda.tools import dtype_to_ctype from cgen.cuda import CudaShared, CudaGlobal discr = self.discr d = discr.dimensions dims = range(d) plan = self.plan given = plan.given elgroup, = discr.element_groups float_type = given.float_type f_decl = CudaGlobal(FunctionDeclaration(Value("void", "apply_diff_mat_smem"), [Pointer(POD(float_type, "debugbuf")), Pointer(POD(float_type, "field")), ] + [Pointer(POD(float_type, "drst%d_global" % i)) for i in dims] )) par = plan.parallelism cmod = Module([ Include("pycuda-helpers.hpp"), ]) if float_type == numpy.float64: cmod.append(Value("texture<fp_tex_double, 1, cudaReadModeElementType>", "diff_rst_mat_tex")) elif float_type == numpy.float32: rst_channels = given.devdata.make_valid_tex_channel_count(d) cmod.append(Value("texture<float%d, 1, cudaReadModeElementType>" % rst_channels, "diff_rst_mat_tex")) else: raise ValueError("unsupported float type: %s" % float_type) # only preimage size variation is supported here assert plan.image_dofs_per_el == given.dofs_per_el() assert plan.aligned_image_dofs_per_microblock == given.microblock.aligned_floats # FIXME: aligned_image_dofs_per_microblock must be divisible # by this, therefore hardcoding for now. chunk_size = 16 cmod.extend([ Line(), Define("DIMENSIONS", discr.dimensions), Define("IMAGE_DOFS_PER_EL", plan.image_dofs_per_el), Define("PREIMAGE_DOFS_PER_EL", plan.preimage_dofs_per_el), Define("ALIGNED_IMAGE_DOFS_PER_MB", plan.aligned_image_dofs_per_microblock), Define("ALIGNED_PREIMAGE_DOFS_PER_MB", plan.aligned_preimage_dofs_per_microblock), Define("ELS_PER_MB", given.microblock.elements), Define("IMAGE_DOFS_PER_MB", "(IMAGE_DOFS_PER_EL*ELS_PER_MB)"), Line(), Define("CHUNK_SIZE", chunk_size), Define("CHUNK_DOF", "threadIdx.x"), Define("PAR_MB_NR", "threadIdx.y"), Define("CHUNK_NR", "threadIdx.z"), Define("IMAGE_MB_DOF", "(CHUNK_NR*CHUNK_SIZE+CHUNK_DOF)"), Define("IMAGE_EL_DOF", "(IMAGE_MB_DOF - mb_el*IMAGE_DOFS_PER_EL)"), Line(), Define("MACROBLOCK_NR", "blockIdx.x"), Line(), Define("PAR_MB_COUNT", par.parallel), Define("INLINE_MB_COUNT", par.inline), Define("SEQ_MB_COUNT", par.serial), Line(), Define("GLOBAL_MB_NR_BASE", "(MACROBLOCK_NR*PAR_MB_COUNT*INLINE_MB_COUNT*SEQ_MB_COUNT)"), Define("GLOBAL_MB_NR", "(GLOBAL_MB_NR_BASE" "+ (seq_mb_number*PAR_MB_COUNT + PAR_MB_NR)*INLINE_MB_COUNT)"), Define("GLOBAL_MB_IMAGE_DOF_BASE", "(GLOBAL_MB_NR*ALIGNED_IMAGE_DOFS_PER_MB)"), Define("GLOBAL_MB_PREIMAGE_DOF_BASE", "(GLOBAL_MB_NR*ALIGNED_PREIMAGE_DOFS_PER_MB)"), Line(), CudaShared( ArrayOf( ArrayOf( ArrayOf( POD(float_type, "smem_field"), "PAR_MB_COUNT"), "INLINE_MB_COUNT"), "ALIGNED_PREIMAGE_DOFS_PER_MB")), Line(), ]) S = Statement f_body = Block([ Initializer(Const(POD(numpy.uint16, "mb_el")), "IMAGE_MB_DOF / IMAGE_DOFS_PER_EL"), Line(), ]) # --------------------------------------------------------------------- def get_load_code(): mb_img_dofs = plan.aligned_image_dofs_per_microblock mb_preimg_dofs = plan.aligned_preimage_dofs_per_microblock preimg_dofs_over_dofs = (mb_preimg_dofs+mb_img_dofs-1) // mb_img_dofs load_code = [] store_code = [] var_num = 0 for load_block in range(preimg_dofs_over_dofs): for inl in range(par.inline): # load and store are split for better pipelining # compiler can't figure that out because of branch var = "tmp%d" % var_num var_num += 1 load_code.append(POD(float_type, var)) block_addr = "%d * ALIGNED_IMAGE_DOFS_PER_MB + IMAGE_MB_DOF" % load_block load_instr = Assign(var, "field[GLOBAL_MB_PREIMAGE_DOF_BASE" " + %d*ALIGNED_PREIMAGE_DOFS_PER_MB" " + %s]" % (inl, block_addr)) store_instr = Assign( "smem_field[PAR_MB_NR][%d][%s]" % (inl, block_addr), var ) if (load_block+1)*mb_img_dofs >= mb_preimg_dofs: cond = "%s < ALIGNED_PREIMAGE_DOFS_PER_MB" % block_addr load_instr = If(cond, load_instr) store_instr = If(cond, store_instr) load_code.append(load_instr) store_code.append(store_instr) return Block(load_code + [Line()] + store_code) def get_scalar_diff_code(): code = [] for inl in range(par.inline): for axis in dims: code.append( Initializer(POD(float_type, "d%drst%d" % (inl, axis)), 0)) code.append(Line()) tex_channels = ["x", "y", "z", "w"] store_code = Block() for inl in range(par.inline): for rst_axis in dims: store_code.append(Assign( "drst%d_global[GLOBAL_MB_IMAGE_DOF_BASE + " "%d*ALIGNED_IMAGE_DOFS_PER_MB + IMAGE_MB_DOF]" % (rst_axis, inl), "d%drst%d" % (inl, rst_axis) )) from hedge.backends.cuda.tools import unroll code.extend([ Comment("everybody needs to be done with the old data"), S("__syncthreads()"), Line(), get_load_code(), Line(), Comment("all the new data must be loaded"), S("__syncthreads()"), Line(), ]) if float_type == numpy.float32: code.append(Value("float%d" % rst_channels, "dmat_entries")) code.extend([ POD(float_type, "field_value%d" % inl) for inl in range(par.inline) ]+[Line()]) def unroll_body(j): result = [ Assign("field_value%d" % inl, "smem_field[PAR_MB_NR][%d][mb_el*PREIMAGE_DOFS_PER_EL+%s]" % (inl, j)) for inl in range(par.inline) ] if float_type == numpy.float32: result.append(Assign("dmat_entries", "tex1Dfetch(diff_rst_mat_tex, IMAGE_EL_DOF + %s*IMAGE_DOFS_PER_EL)" % j)) result.extend( S("d%drst%d += dmat_entries.%s * field_value%d" % (inl, axis, tex_channels[axis], inl)) for inl in range(par.inline) for axis in dims) elif float_type == numpy.float64: result.extend( S("d%(inl)drst%(axis)d += " "fp_tex1Dfetch(diff_rst_mat_tex, %(axis)d " "+ DIMENSIONS*(IMAGE_EL_DOF + %(j)d*IMAGE_DOFS_PER_EL))" "* field_value%(inl)d" % { "inl": inl, "axis": axis, "j": j }) for inl in range(par.inline) for axis in dims) else: assert False return result code.append(If("IMAGE_MB_DOF < IMAGE_DOFS_PER_MB", Block(unroll(unroll_body, total_number=plan.preimage_dofs_per_el) +[store_code]))) return code f_body.extend([ For("unsigned short seq_mb_number = 0", "seq_mb_number < SEQ_MB_COUNT", "++seq_mb_number", Block(get_scalar_diff_code()) ) ]) # finish off ---------------------------------------------------------- cmod.append(FunctionBody(f_decl, f_body)) if not for_benchmark and "cuda_dump_kernels" in discr.debug: from hedge.tools import open_unique_debug_file open_unique_debug_file("diff", ".cu").write(str(cmod)) mod = SourceModule(cmod, keep="cuda_keep_kernels" in discr.debug, #options=["--maxrregcount=16"] ) func = mod.get_function("apply_diff_mat_smem") if "cuda_diff" in discr.debug: print "diff: lmem=%d smem=%d regs=%d" % ( func.local_size_bytes, func.shared_size_bytes, func.registers) diff_rst_mat_texref = mod.get_texref("diff_rst_mat_tex") gpu_diffmats = self.gpu_diffmats(diff_op, elgroup) if given.float_type == numpy.float32: gpu_diffmats.bind_to_texref_ext(diff_rst_mat_texref, rst_channels) elif given.float_type == numpy.float64: gpu_diffmats.bind_to_texref_ext(diff_rst_mat_texref, allow_double_hack=True) else: assert False assert given.microblock.aligned_floats % chunk_size == 0 block = ( chunk_size, plan.parallelism.parallel, given.microblock.aligned_floats//chunk_size) func.prepare( ["PP"] + discr.dimensions*["P"], texrefs=[diff_rst_mat_texref]) return block, func
from cgen import FunctionBody, \ FunctionDeclaration, POD, Value, \ Pointer, Module, Block, Initializer, Assign, Const from cgen.opencl import CLKernel, CLGlobal, \ CLRequiredWorkGroupSize mod = Module([ FunctionBody( CLKernel( CLRequiredWorkGroupSize( (local_size, ), FunctionDeclaration(Value("void", "add"), arg_decls=[ CLGlobal( Pointer(Const(POD(dtype, name)))) for name in ["tgt", "op1", "op2"] ]))), Block([ Initializer( POD(numpy.int32, "idx"), "get_local_id(0) + %d * get_group_id(0)" % (local_size * thread_strides)) ] + [ Assign( "tgt[idx+%d]" % (o * local_size), "op1[idx+%d] + op2[idx+%d]" % (o * local_size, o * local_size)) for o in range(thread_strides) ])) ])
def make_lift(self, fgroup, with_scale, dtype): discr = self.discr from cgen import (FunctionDeclaration, FunctionBody, Typedef, Const, Reference, Value, POD, Statement, Include, Line, Block, Initializer, Assign, For, If, Define) from pytools import to_uncomplex_dtype from codepy.bpl import BoostPythonModule mod = BoostPythonModule() S = Statement mod.add_to_preamble([ Include("hedge/face_operators.hpp"), Include("hedge/volume_operators.hpp"), Include("boost/foreach.hpp"), ]) mod.add_to_module([ S("namespace ublas = boost::numeric::ublas"), S("using namespace hedge"), S("using namespace pyublas"), Line(), Define("DOFS_PER_EL", fgroup.ldis_loc.node_count()), Define("FACES_PER_EL", fgroup.ldis_loc.face_count()), Define("DIMENSIONS", discr.dimensions), Line(), Typedef(POD(dtype, "value_type")), Typedef(POD(to_uncomplex_dtype(dtype), "uncomplex_type")), ]) def if_(cond, result, else_=None): if cond: return [result] else: if else_ is None: return [] else: return [else_] fdecl = FunctionDeclaration(Value("void", "lift"), [ Const( Reference(Value("face_group<face_pair<straight_face> >", "fg"))), Value("ublas::matrix<uncomplex_type>", "matrix"), Value("numpy_array<value_type>", "field"), Value("numpy_array<value_type>", "result") ] + if_( with_scale, Const( Reference(Value("numpy_array<double>", "elwise_post_scaling"))))) def make_it(name, is_const=True, tpname="value_type"): if is_const: const = "const_" else: const = "" return Initializer( Value("numpy_array<%s>::%siterator" % (tpname, const), name + "_it"), "%s.begin()" % name) fbody = Block([ make_it("field"), make_it("result", is_const=False), ] + if_(with_scale, make_it("elwise_post_scaling", tpname="double")) + [ Line(), For( "unsigned fg_el_nr = 0", "fg_el_nr < fg.element_count()", "++fg_el_nr", Block([ Initializer(Value("node_number_t", "dest_el_base"), "fg.local_el_write_base[fg_el_nr]"), Initializer(Value("node_number_t", "src_el_base"), "FACES_PER_EL*fg.face_length()*fg_el_nr"), Line(), For( "unsigned i = 0", "i < DOFS_PER_EL", "++i", Block([ Initializer(Value("value_type", "tmp"), 0), Line(), For( "unsigned j = 0", "j < FACES_PER_EL*fg.face_length()", "++j", S("tmp += matrix(i, j)*field_it[src_el_base+j]" )), Line(), ] + if_( with_scale, Assign( "result_it[dest_el_base+i]", "tmp * value_type(*elwise_post_scaling_it)"), Assign("result_it[dest_el_base+i]", "tmp")))), ] + if_(with_scale, S("elwise_post_scaling_it++")))) ]) mod.add_function(FunctionBody(fdecl, fbody)) #print "----------------------------------------------------------------" #print FunctionBody(fdecl, fbody) #raw_input() return mod.compile(self.discr.toolchain).lift
def get_boundary_flux_mod(fluxes, fvi, discr, dtype): from cgen import \ FunctionDeclaration, FunctionBody, Typedef, Struct, \ Const, Reference, Value, POD, MaybeUnused, \ Statement, Include, Line, Block, Initializer, Assign, \ CustomLoop, For from pytools import to_uncomplex_dtype, flatten from codepy.bpl import BoostPythonModule mod = BoostPythonModule() mod.add_to_preamble([ Include("cstdlib"), Include("algorithm"), Line(), Include("boost/foreach.hpp"), Line(), Include("hedge/face_operators.hpp"), ]) S = Statement mod.add_to_module([ S("using namespace hedge"), S("using namespace pyublas"), Line(), Typedef(POD(dtype, "value_type")), Typedef(POD(to_uncomplex_dtype(dtype), "uncomplex_type")), ]) arg_struct = Struct("arg_struct", [ Value("numpy_array<value_type>", "flux%d_on_faces" % i) for i in range(len(fluxes)) ]+[ Value("numpy_array<value_type>", arg_name) for arg_name in fvi.arg_names ]) mod.add_struct(arg_struct, "ArgStruct") mod.add_to_module([Line()]) fdecl = FunctionDeclaration( Value("void", "gather_flux"), [ Const(Reference(Value("face_group<face_pair<straight_face> >" , "fg"))), Reference(Value("arg_struct", "args")) ]) from pymbolic.mapper.stringifier import PREC_PRODUCT def gen_flux_code(): f2cm = FluxToCodeMapper() result = [ Assign("fof%d_it[loc_fof_base+i]" % flux_idx, "uncomplex_type(fp.int_side.face_jacobian) * " + flux_to_code(f2cm, False, flux_idx, fvi, flux.op.flux, PREC_PRODUCT)) for flux_idx, flux in enumerate(fluxes) ] return [ Initializer(Value("value_type", cse_name), cse_str) for cse_name, cse_str in f2cm.cse_name_list] + result fbody = Block([ Initializer( Const(Value("numpy_array<value_type>::iterator", "fof%d_it" % i)), "args.flux%d_on_faces.begin()" % i) for i in range(len(fluxes)) ]+[ Initializer( Const(Value("numpy_array<value_type>::const_iterator", "%s_it" % arg_name)), "args.%s.begin()" % arg_name) for arg_name in fvi.arg_names ]+[ Line(), CustomLoop("BOOST_FOREACH(const face_pair<straight_face> &fp, fg.face_pairs)", Block( list(flatten([ Initializer(Value("node_number_t", "%s_ebi" % where), "fp.%s.el_base_index" % where), Initializer(Value("index_lists_t::const_iterator", "%s_idx_list" % where), "fg.index_list(fp.%s.face_index_list_number)" % where), Line(), ] for where in ["int_side", "ext_side"] ))+[ Line(), Initializer(Value("node_number_t", "loc_fof_base"), "fg.face_length()*(fp.%(where)s.local_el_number*fg.face_count" " + fp.%(where)s.face_id)" % {"where": "int_side"}), Line(), For( "unsigned i = 0", "i < fg.face_length()", "++i", Block( [ Initializer(MaybeUnused( Value("node_number_t", "%s_idx" % where)), "%(where)s_ebi + %(where)s_idx_list[i]" % {"where": where}) for where in ["int_side", "ext_side"] ]+gen_flux_code() ) ) ])) ]) mod.add_function(FunctionBody(fdecl, fbody)) #print "----------------------------------------------------------------" #print mod.generate() #raw_input("[Enter]") return mod.compile(get_flux_toolchain(discr, fluxes))
def make_cuda_kernel(self, discr, dtype, eg): given = discr.given ldis = eg.local_discretization microblocks_per_block = 1 from cgen.cuda import CudaGlobal from cgen import (Module, Value, Include, Typedef, FunctionBody, FunctionDeclaration, Const, Line, POD, LiteralBlock, Define, Pointer) cmod = Module([ Include("pycuda-helpers.hpp"), Line(), Typedef(POD(dtype, "value_type")), Line(), Define("DOFS_PER_EL", given.dofs_per_el()), Define("ALIGNED_DOFS_PER_MB", given.microblock.aligned_floats), Define("VERTICES_PER_EL", ldis.vertex_count()), Define("ELS_PER_MB", given.microblock.elements), Define("MBS_PER_BLOCK", microblocks_per_block), Line(), Define("DOF_IN_MB_IDX", "threadIdx.x"), Define("DOF_IN_EL_IDX", "(DOF_IN_MB_IDX-el_idx_in_mb*DOFS_PER_EL)"), Define("MB_IN_BLOCK_IDX", "threadIdx.y"), Define("BLOCK_IDX", "blockIdx.x"), Define("MB_NUMBER", "(BLOCK_IDX * MBS_PER_BLOCK + MB_IN_BLOCK_IDX)"), Define("BLOCK_DATA", "whole_block[MB_IN_BLOCK_IDX]")] + self.get_cuda_extra_preamble(discr, dtype, eg) + [FunctionBody( CudaGlobal(FunctionDeclaration( Value("void", "elwise_kernel"), [ Pointer(Const(POD(dtype, "field"))), Pointer(POD(dtype, "result")), POD(numpy.uint32, "mb_count"), ])), LiteralBlock(""" int el_idx_in_mb = DOF_IN_MB_IDX / DOFS_PER_EL; if (MB_NUMBER >= mb_count) return; int idx = MB_NUMBER * ALIGNED_DOFS_PER_MB + DOF_IN_MB_IDX; int element_base_idx = ALIGNED_DOFS_PER_MB * MB_IN_BLOCK_IDX + (DOF_IN_MB_IDX / DOFS_PER_EL) * DOFS_PER_EL; int dof_in_element = DOF_IN_MB_IDX-el_idx_in_mb*DOFS_PER_EL; __shared__ value_type whole_block[MBS_PER_BLOCK][ALIGNED_DOFS_PER_MB+1]; int idx_in_block = ALIGNED_DOFS_PER_MB * MB_IN_BLOCK_IDX + DOF_IN_MB_IDX; BLOCK_DATA[idx_in_block] = field[idx]; __syncthreads(); %s result[idx] = node_result; """ % self.get_cuda_code(discr, dtype, eg))) ]) if False: for i, l in enumerate(str(cmod).split("\n")): print i+1, l raw_input() from pycuda.compiler import SourceModule mod = SourceModule( cmod, keep="cuda_keep_kernels" in discr.debug, ) func = mod.get_function("elwise_kernel") func.prepare( "PPI", block=( given.microblock.aligned_floats, microblocks_per_block, 1)) mb_count = len(discr.blocks) * discr.given.microblocks_per_block grid_dim = (mb_count + microblocks_per_block - 1) \ // microblocks_per_block from pytools import Record class KernelInfo(Record): pass return KernelInfo( func=func, grid_dim=grid_dim, mb_count=mb_count)
def make_diff(self, elgroup, dtype, shape): """ :param shape: If non-square, the resulting code takes two element_ranges arguments and supports non-square matrices. """ from hedge._internal import UniformElementRanges assert isinstance(elgroup.ranges, UniformElementRanges) ldis = elgroup.local_discretization discr = self.discr from cgen import ( FunctionDeclaration, FunctionBody, Typedef, Const, Reference, Value, POD, Statement, Include, Line, Block, Initializer, Assign, For, If, Define) from pytools import to_uncomplex_dtype from codepy.bpl import BoostPythonModule mod = BoostPythonModule() # {{{ preamble S = Statement mod.add_to_preamble([ Include("hedge/volume_operators.hpp"), Include("boost/foreach.hpp"), ]) mod.add_to_module([ S("namespace ublas = boost::numeric::ublas"), S("using namespace hedge"), S("using namespace pyublas"), Line(), Define("ROW_COUNT", shape[0]), Define("COL_COUNT", shape[1]), Define("DIMENSIONS", discr.dimensions), Line(), Typedef(POD(dtype, "value_type")), Typedef(POD(to_uncomplex_dtype(dtype), "uncomplex_type")), ]) fdecl = FunctionDeclaration( Value("void", "diff"), [ Const(Reference(Value("uniform_element_ranges", "from_ers"))), Const(Reference(Value("uniform_element_ranges", "to_ers"))), Value("numpy_array<value_type>", "field") ]+[ Value("ublas::matrix<uncomplex_type>", "diffmat_rst%d" % rst) for rst in range(discr.dimensions) ]+[ Value("numpy_array<value_type>", "result%d" % i) for i in range(discr.dimensions) ] ) # }}} # {{{ set-up def make_it(name, is_const=True, tpname="value_type"): if is_const: const = "const_" else: const = "" return Initializer( Value("numpy_array<%s>::%siterator" % (tpname, const), name+"_it"), "%s.begin()" % name) fbody = Block([ If("ROW_COUNT != diffmat_rst%d.size1()" % i, S('throw(std::runtime_error("unexpected matrix size"))')) for i in range(discr.dimensions) ] + [ If("COL_COUNT != diffmat_rst%d.size2()" % i, S('throw(std::runtime_error("unexpected matrix size"))')) for i in range(discr.dimensions) ]+[ If("ROW_COUNT != to_ers.el_size()", S('throw(std::runtime_error("unsupported image element size"))')), If("COL_COUNT != from_ers.el_size()", S('throw(std::runtime_error("unsupported preimage element size"))')), If("from_ers.size() != to_ers.size()", S('throw(std::runtime_error("image and preimage element groups ' 'do nothave the same element count"))')), Line(), make_it("field"), ]+[ make_it("result%d" % i, is_const=False) for i in range(discr.dimensions) ]+[ Line(), # }}} # {{{ computation For("element_number_t eg_el_nr = 0", "eg_el_nr < to_ers.size()", "++eg_el_nr", Block([ Initializer( Value("node_number_t", "from_el_base"), "from_ers.start() + eg_el_nr*COL_COUNT"), Initializer( Value("node_number_t", "to_el_base"), "to_ers.start() + eg_el_nr*ROW_COUNT"), Line(), For("unsigned i = 0", "i < ROW_COUNT", "++i", Block([ Initializer(Value("value_type", "drst_%d" % rst), 0) for rst in range(discr.dimensions) ]+[ Line(), ]+[ For("unsigned j = 0", "j < COL_COUNT", "++j", Block([ S("drst_%(rst)d += " "diffmat_rst%(rst)d(i, j)*field_it[from_el_base+j]" % {"rst":rst}) for rst in range(discr.dimensions) ]) ), Line(), ]+[ Assign("result%d_it[to_el_base+i]" % rst, "drst_%d" % rst) for rst in range(discr.dimensions) ]) ) ]) ) ]) # }}} # {{{ compilation mod.add_function(FunctionBody(fdecl, fbody)) #print "----------------------------------------------------------------" #print mod.generate() #raw_input() compiled_func = mod.compile(self.discr.toolchain).diff if self.discr.instrumented: from hedge.tools import time_count_flop compiled_func = time_count_flop(compiled_func, discr.diff_timer, discr.diff_counter, discr.diff_flop_counter, flops=discr.dimensions*( 2 # mul+add * ldis.node_count() * len(elgroup.members) * ldis.node_count() + 2 * discr.dimensions * len(elgroup.members) * ldis.node_count()), increment=discr.dimensions) return compiled_func