def mat_mul_outer_loop(fetch_count):
if with_scaling:
inv_jac_multiplier = (
"fp_tex1Dfetch(scaling_tex,"
"(GLOBAL_MB_NR + %(inl)d)*MB_EL_COUNT + mb_el)")
else:
inv_jac_multiplier = "1"
write_condition = "MB_DOF < DOFS_PER_EL*MB_EL_COUNT"
if self.with_index_check:
write_condition += " && GLOBAL_MB_NR < microblock_count"
return For(
"unsigned short seq_mb_number = 0",
"seq_mb_number < SEQ_MB_COUNT", "++seq_mb_number",
Block([
Initializer(POD(float_type, "result%d" % inl), 0)
for inl in range(par.inline)
] + [Line()] + get_mat_mul_code(fetch_count) + [
If(
write_condition,
Block([
Assign(
"out_vector[GLOBAL_MB_DOF_BASE"
" + %d*ALIGNED_DOFS_PER_MB"
" + MB_DOF]" % inl, "result%d * %s" %
(inl, (inv_jac_multiplier % {
"inl": inl
}))) for inl in range(par.inline)
]))
]))
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 get_elwise_module_descriptor(arguments, operation, name="kernel"):
from codepy.bpl import BoostPythonModule
from cgen import FunctionBody, FunctionDeclaration, \
Value, POD, Struct, For, Initializer, Include, Statement, \
Line, Block
S = Statement
mod = BoostPythonModule()
mod.add_to_preamble([
Include("pyublas/numpy.hpp"),
])
mod.add_to_module([
S("namespace ublas = boost::numeric::ublas"),
S("using namespace pyublas"),
Line(),
])
body = Block([
Initializer(
Value("numpy_array<%s >::iterator"
% dtype_to_ctype(varg.dtype),
varg.name),
"args.%s_ary.begin()" % varg.name)
for varg in arguments if isinstance(varg, VectorArg)]
+[Initializer(
sarg.declarator(), "args." + sarg.name)
for sarg in arguments if isinstance(sarg, ScalarArg)]
)
body.extend([
Line(),
For("unsigned i = 0",
"i < codepy_length",
"++i",
Block([S(operation)])
)
])
arg_struct = Struct("arg_struct",
[arg.declarator() for arg in arguments])
mod.add_struct(arg_struct, "ArgStruct")
mod.add_to_module([Line()])
mod.add_function(
FunctionBody(
FunctionDeclaration(
Value("void", name),
[POD(numpy.uintp, "codepy_length"),
Value("arg_struct", "args")]),
body))
return mod
def generate_body(self, kernel, codegen_state):
from cgen import Block
body = Block()
# {{{ declare temporaries
body.extend(
idi.cgen_declarator
for tv in six.itervalues(kernel.temporary_variables)
for idi in tv.decl_info(kernel.target, is_written=True, index_dtype=kernel.index_dtype)
)
# }}}
from loopy.codegen.loop import set_up_hw_parallel_loops
gen_code = set_up_hw_parallel_loops(kernel, 0, codegen_state)
from cgen import Line
body.append(Line())
if isinstance(gen_code.ast, Block):
body.extend(gen_code.ast.contents)
else:
body.append(gen_code.ast)
return body, gen_code.implemented_domains
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())
def get_mat_entry(row, col, axis):
return ("smem_diff_rst_mat["
"%(row)s*DIFFMAT_COLUMNS + %(axis)s*DOFS_PER_EL"
" + %(col)s"
"]" % {
"row": row,
"col": col,
"axis": axis
})
tex_channels = ["x", "y", "z", "w"]
from hedge.backends.cuda.tools import unroll
code.extend([
POD(float_type, "field_value%d" % inl)
for inl in range(par.inline)
] + [Line()] + unroll(
lambda j: [
Assign(
"field_value%d" % inl,
"fp_tex1Dfetch(field_tex, GLOBAL_MB_DOF_BASE + %d*ALIGNED_DOFS_PER_MB "
"+ mb_el*DOFS_PER_EL + %s)" % (inl, j))
for inl in range(par.inline)
] + [Line()] + [
S("d%drst%d += %s * field_value%d" %
(inl, axis, get_mat_entry("SEGMENT_DOF", j, axis), inl))
for axis in dims for inl in range(par.inline)
] + [Line()], given.dofs_per_el(), self.plan.max_unroll))
store_code = Block()
for inl in range(par.inline):
for rst_axis in dims:
store_code.append(
Assign(
"drst%d_global[GLOBAL_MB_DOF_BASE"
" + %d*ALIGNED_DOFS_PER_MB + MB_DOF]" %
(rst_axis, inl),
"d%drst%d" % (inl, rst_axis),
))
code.append(If("MB_DOF < DOFS_PER_EL*ELS_PER_MB", store_code))
return code
def test_cgen():
s = Struct(
"yuck",
[
POD(
np.float32,
"h",
),
POD(np.float32, "order"),
POD(np.float32, "face_jacobian"),
ArrayOf(POD(np.float32, "normal"), 17),
POD(np.uint16, "a_base"),
POD(np.uint16, "b_base"),
#CudaGlobal(POD(np.uint8, "a_ilist_number")),
POD(np.uint8, "b_ilist_number"),
POD(np.uint8, "bdry_flux_number"), # 0 if not on boundary
POD(np.uint8, "reserved"),
POD(np.uint32, "b_global_base"),
])
f_decl = FunctionDeclaration(POD(np.uint16, "get_num"), [
POD(np.uint8, "reserved"),
POD(np.uint32, "b_global_base"),
])
f_body = FunctionBody(
f_decl,
Block([
POD(np.uint32, "i"),
For(
"i = 0",
"i < 17",
"++i",
If(
"a > b",
Assign("a", "b"),
Block([
Assign("a", "b-1"),
#Break(),
])),
),
#BlankLine(),
Comment("all done"),
]))
t_decl = Template(
'typename T',
FunctionDeclaration(
Value('CUdeviceptr', 'scan'),
[Value('CUdeviceptr', 'inputPtr'),
Value('int', 'length')]))
print(s)
print(f_body)
print(t_decl)
def get_elwise_module_descriptor(arguments, operation, name="kernel"):
from codepy.bpl import BoostPythonModule
from cgen import FunctionBody, FunctionDeclaration, \
Value, POD, Struct, For, Initializer, Include, Statement, \
Line, Block
S = Statement # noqa: N806
mod = BoostPythonModule()
mod.add_to_preamble([
Include("pyublas/numpy.hpp"),
])
mod.add_to_module([
S("namespace ublas = boost::numeric::ublas"),
S("using namespace pyublas"),
Line(),
])
body = Block([
Initializer(
Value(
"numpy_array<{} >::iterator".format(dtype_to_ctype(
varg.dtype)), varg.name), f"args.{varg.name}_ary.begin()")
for varg in arguments if isinstance(varg, VectorArg)
] + [
Initializer(sarg.declarator(), f"args.{sarg.name}")
for sarg in arguments if isinstance(sarg, ScalarArg)
])
body.extend([
Line(),
For("unsigned i = 0", "i < codepy_length", "++i",
Block([S(operation)]))
])
arg_struct = Struct("arg_struct", [arg.declarator() for arg in arguments])
mod.add_struct(arg_struct, "ArgStruct")
mod.add_to_module([Line()])
mod.add_function(
FunctionBody(
FunctionDeclaration(Value("void", name), [
POD(numpy.uintp, "codepy_length"),
Value("arg_struct", "args")
]), body))
return mod
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())
def get_mat_entry(row, col, axis):
return ("smem_diff_rst_mat["
"%(row)s*DIFFMAT_COLUMNS + %(axis)s*DOFS_PER_EL"
" + %(col)s"
"]" % {"row":row, "col":col, "axis":axis}
)
tex_channels = ["x", "y", "z", "w"]
from hedge.backends.cuda.tools import unroll
code.extend(
[POD(float_type, "field_value%d" % inl)
for inl in range(par.inline)]
+[Line()]
+unroll(lambda j: [
Assign("field_value%d" % inl,
"fp_tex1Dfetch(field_tex, GLOBAL_MB_DOF_BASE + %d*ALIGNED_DOFS_PER_MB "
"+ mb_el*DOFS_PER_EL + %s)" % (inl, j)
)
for inl in range(par.inline)]
+[Line()]
+[S("d%drst%d += %s * field_value%d"
% (inl, axis, get_mat_entry("SEGMENT_DOF", j, axis), inl))
for axis in dims
for inl in range(par.inline)]
+[Line()],
given.dofs_per_el(), self.plan.max_unroll)
)
store_code = Block()
for inl in range(par.inline):
for rst_axis in dims:
store_code.append(Assign(
"drst%d_global[GLOBAL_MB_DOF_BASE"
" + %d*ALIGNED_DOFS_PER_MB + MB_DOF]" % (rst_axis, inl),
"d%drst%d" % (inl, rst_axis),
))
code.append(If("MB_DOF < DOFS_PER_EL*ELS_PER_MB", store_code))
return code
def get_matmul_code():
from hedge.backends.cuda.tools import unroll
index_check_condition = "GLOBAL_MB_NR < microblock_count"
def if_(conditions, then):
final_cond = " && ".join(cond for cond in conditions if cond)
if final_cond:
return If(final_cond, then)
else:
return then
result = Block([
Comment("everybody needs to be done with the old data"),
S("__syncthreads()"), Line(),
]+[If(index_check_condition, get_load_code())]+[
Line(),
Comment("all the new data must be loaded"),
S("__syncthreads()"),
Line(),
]+[
Initializer(POD(float_type, "result%d" % inl), 0)
for inl in range(par.inline)
]+[
Line(),
POD(float_type, "mat_entry"),
Line(),
])
result.append(if_(["IMAGE_MB_DOF < IMAGE_DOFS_PER_MB", index_check_condition],
Block(unroll(lambda j:
[Assign("mat_entry", "fp_tex2D(mat_tex, IMAGE_EL_DOF, %s)" % j)]
+[
S("result%d += mat_entry "
"* smem_in_vector[PAR_MB_NR][%d][mb_el*PREIMAGE_DOFS_PER_EL + %s]"
% (inl, inl, j))
for inl in range(par.inline)
],
total_number=plan.preimage_dofs_per_el)
+[Line()]
+[Assign(
"out_vector[GLOBAL_MB_IMAGE_DOF_BASE + "
"%d*ALIGNED_IMAGE_DOFS_PER_MB + IMAGE_MB_DOF]" % inl,
"result%d" % inl)
for inl in range(par.inline)]
)))
return result
def get_kernel_call(self, codegen_state, name, gsize, lsize, extra_args):
ecm = self.get_expression_to_code_mapper(codegen_state)
from pymbolic.mapper.stringifier import PREC_COMPARISON, PREC_NONE
result = []
from cgen import Statement as S, Block
if lsize:
result.append(
S("assert(programCount == %s)"
% ecm(lsize[0], PREC_COMPARISON)))
if gsize:
launch_spec = "[%s]" % ", ".join(
ecm(gs_i, PREC_NONE)
for gs_i in gsize)
else:
launch_spec = ""
arg_names, arg_decls = self._arg_names_and_decls(codegen_state)
result.append(S(
"launch%s %s(%s)" % (
launch_spec,
name,
", ".join(arg_names)
)))
return Block(result)
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_load_code(dest, base, bytes, word_type=numpy.uint32,
descr=None):
from cgen import (
Pointer, POD,
Comment, Block, Line, \
Constant, For, Statement)
from cgen import dtype_to_ctype
copy_dtype = numpy.dtype(word_type)
copy_dtype_str = dtype_to_ctype(copy_dtype)
code = []
if descr is not None:
code.append(Comment(descr))
code.extend([
Block([
Constant(Pointer(POD(copy_dtype, "load_base")),
("(%s *) (%s)" % (copy_dtype_str, base))),
For("unsigned word_nr = THREAD_NUM",
"word_nr*sizeof(int) < (%s)" % bytes,
"word_nr += COALESCING_THREAD_COUNT",
Statement("((%s *) (%s))[word_nr] = load_base[word_nr]"
% (copy_dtype_str, dest))
),
]),
Line(),
])
return code
def unroll(body_gen, total_number, max_unroll=None, start=0):
from cgen import For, Line, Block
from pytools import flatten
if max_unroll is None:
max_unroll = total_number
result = []
if total_number > max_unroll:
loop_items = (total_number // max_unroll) * max_unroll
result.extend([
For("unsigned j = 0",
"j < %d" % loop_items,
"j += %d" % max_unroll,
Block(list(flatten(
body_gen("(j+%d)" % i)
for i in range(max_unroll))))
),
Line()
])
start += loop_items
result.extend(flatten(
body_gen(i) for i in range(start, total_number)))
return result
def add_struct(self,
struct, py_name=None, py_member_name_transform=lambda x: x,
by_value_members=None):
if by_value_members is None:
by_value_members = set()
from cgen import Block, Line, Statement, Typedef, Value
if py_name is None:
py_name = struct.tpname
self.mod_body.append(struct)
member_defs = []
for f in struct.fields:
py_f_name = py_member_name_transform(f.name)
tp_lines, declarator = f.get_decl_pair()
if f.name in by_value_members or tp_lines[0].startswith("numpy_"):
member_defs.append(
".def(pyublas::by_value_rw_member"
f'("{py_f_name}", &cl::{f.name}))')
else:
member_defs.append(
f'.def_readwrite("{py_f_name}", &cl::{f.name})'
)
self.init_body.append(
Block([
Typedef(Value(struct.tpname, "cl")),
Line(),
Statement(
'boost::python::class_<cl>("{}"){}'.format(
py_name, "".join(member_defs))),
]))
def generate(self):
"""Generate (i.e. yield) the source code of the
module line-by-line.
"""
from cgen import Block, Module, Include, Line, Define, \
PrivateNamespace
body = []
if self.max_arity is not None:
body.append(Define("BOOST_PYTHON_MAX_ARITY", self.max_arity))
if self.use_private_namespace:
mod_body = [PrivateNamespace(self.mod_body)]
else:
mod_body = self.mod_body
body += ([Include("boost/python.hpp")]
+ self.preamble + [Line()]
+ mod_body
+ [Line(), Line(f"BOOST_PYTHON_MODULE({self.name})")]
+ [Block(self.init_body)])
return Module(body)
def generate_c_instruction_code(codegen_state, insn):
kernel = codegen_state.kernel
if codegen_state.vectorization_info is not None:
raise Unvectorizable("C instructions cannot be vectorized")
body = []
from loopy.target.c import POD
from cgen import Initializer, Block, Line
from pymbolic.primitives import Variable
for name, iname_expr in insn.iname_exprs:
if (isinstance(iname_expr, Variable)
and name not in codegen_state.var_subst_map):
# No need, the bare symbol will work
continue
body.append(
Initializer(
POD(codegen_state.ast_builder, kernel.index_dtype, name),
codegen_state.expression_to_code_mapper(iname_expr,
prec=PREC_NONE,
type_context="i")))
if body:
body.append(Line())
body.extend(Line(line) for line in insn.code.split("\n"))
return Block(body)
def get_flux_code(flux_writer):
flux_code = Block([])
flux_code.extend([
Initializer(Pointer(Value("face_pair", "fpair")),
"data.facepairs+fpair_nr"),
Initializer(
MaybeUnused(POD(numpy.uint32, "a_index")),
"fpair->a_base + tex1Dfetch(tex_index_lists, "
"fpair->a_ilist_index + FACEDOF_NR)"),
Initializer(
MaybeUnused(POD(numpy.uint32, "b_index")),
"fpair->b_base + tex1Dfetch(tex_index_lists, "
"fpair->b_ilist_index + FACEDOF_NR)"),
Line(),
flux_writer(),
Line(),
S("fpair_nr += CONCURRENT_FACES")
])
return flux_code
def get_flux_code(flux_writer):
flux_code = Block([])
flux_code.extend([
Initializer(Pointer(
Value("face_pair", "fpair")),
"data.facepairs+fpair_nr"),
Initializer(
MaybeUnused(POD(numpy.uint32, "a_index")),
"fpair->a_base + tex1Dfetch(tex_index_lists, "
"fpair->a_ilist_index + FACEDOF_NR)"),
Initializer(
MaybeUnused(POD(numpy.uint32, "b_index")),
"fpair->b_base + tex1Dfetch(tex_index_lists, "
"fpair->b_ilist_index + FACEDOF_NR)"),
Line(),
flux_writer(),
Line(),
S("fpair_nr += CONCURRENT_FACES")
])
return flux_code
def emit_sequential_loop(self, codegen_state, iname, iname_dtype, lbound,
ubound, inner):
from cgen import Pragma, Block
loop = super().emit_sequential_loop(codegen_state, iname, iname_dtype,
lbound, ubound, inner)
pragma = self.target.iname_pragma_map.get(iname)
if pragma:
return Block(contents=[
Pragma(pragma),
loop,
])
return loop
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 expose_vector_type(self, name, py_name=None):
self.add_codepy_include()
if py_name is None:
py_name = name
from cgen import (Block, Typedef, Line, Statement, Value)
self.init_body.append(
Block([
Typedef(Value(name, "cl")),
Line(),
Statement(
f'boost::python::class_<cl>("{py_name}")'
".def(codepy::no_compare_indexing_suite<cl>())"),
]))
def get_cpu_per_element_code(self):
from cgen import (Value, Statement, Initializer, While,
Comment, Block, For, Line, Pointer)
S = Statement
return [
# assumes there is more than one coefficient
Initializer(Value("cit_type", "el_modes"), "field_it+er.first"),
Line(),
Comment("zero out reduced_modes"),
For("npy_uint32 mode_idx = 0",
"mode_idx < max_degree+1",
"++mode_idx",
S("reduced_modes[mode_idx] = 0")),
Line(),
Comment("gather modes by degree"),
For("npy_uint32 mode_idx = 0",
"mode_idx < mode_count",
"++mode_idx",
S("reduced_modes[mode_degrees_iterator[mode_idx]]"
" += el_modes[mode_idx]")),
Line(),
Comment("perform skyline procedure"),
Initializer(Pointer(Value("value_type", "start")),
"reduced_modes.get()"),
Initializer(Pointer(Value("value_type", "end")),
"start+max_degree+1"),
Initializer(Value("value_type", "cur_max"),
"std::max(*(end-1), *(end-2))"),
Line(),
While("end != start", Block([
S("--end"),
S("*end = std::max(cur_max, *end)"),
])),
Line(),
Comment("scatter modes by degree"),
Initializer(Value("it_type", "tgt_base"), "result_it+er.first"),
For("npy_uint32 mode_idx = 0",
"mode_idx < mode_count",
"++mode_idx",
S("tgt_base[mode_idx] = "
"reduced_modes[mode_degrees_iterator[mode_idx]]")),
]
def get_kernel_call(self, codegen_state, name, gsize, lsize, extra_args):
ecm = self.get_expression_to_code_mapper(codegen_state)
from pymbolic.mapper.stringifier import PREC_NONE
result = []
from cgen import Statement as S, Block
if lsize:
result.append(
S("assert(programCount == (%s))" % ecm(lsize[0], PREC_NONE)))
arg_names, arg_decls = self._arg_names_and_decls(codegen_state)
from cgen.ispc import ISPCLaunch
result.append(
ISPCLaunch(tuple(ecm(gs_i, PREC_NONE) for gs_i in gsize),
"%s(%s)" % (name, ", ".join(arg_names))))
return Block(result)
def get_cpu_per_element_code(self):
from cgen import (Value, Statement, Initializer, While, Block)
S = Statement
return [
# assumes there is more than one coefficient
Initializer(Value("cit_type", "start"), "field_it+er.first"),
Initializer(Value("cit_type", "end"), "field_it+er.second"),
Initializer(Value("it_type", "tgt"), "result_it+er.first"),
Initializer(Value("cit_type", "cur"), "start"),
While("cur != end",
Block([
Initializer(Value("cit_type", "avg_start"),
"std::max(start, cur-1)"),
Initializer(Value("cit_type", "avg_end"),
"std::min(end, cur+2)"),
S("*tgt++ = std::accumulate(avg_start, avg_end, value_type(0))"
"/std::distance(avg_start, avg_end)"),
S("++cur"),
])
)
]
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
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
Pointer, Module, Block, Initializer, Assign
from cgen.cuda import CudaGlobal
mod = Module([
FunctionBody(
CudaGlobal(
FunctionDeclaration(Value("void", "add"),
arg_decls=[
Pointer(POD(dtype, name))
for name in ["tgt", "op1", "op2"]
])),
Block([
Initializer(
POD(numpy.int32, "idx"), "threadIdx.x + %d*blockIdx.x" %
(block_size * thread_strides)),
] + [
Assign(
"tgt[idx+%d]" % (o * block_size), "op1[idx+%d] + op2[idx+%d]" %
(o * block_size, o * block_size))
for o in range(thread_strides)
]))
])
mod = SourceModule(mod)
func = mod.get_function("add")
func(c_gpu, a_gpu, b_gpu, block=(block_size, 1, 1), grid=(macroblock_count, 1))
c = cuda.from_device_like(c_gpu, a)
assert la.norm(c - (a + b)) == 0
def get_kernel(self, diff_op_cls, elgroup, for_benchmark=False):
from cgen import \
Pointer, POD, Value, ArrayOf, \
Module, FunctionDeclaration, FunctionBody, Block, \
Line, Define, Include, \
Initializer, If, For, Statement, Assign
from cgen import dtype_to_ctype
from cgen.cuda import CudaShared, CudaGlobal
discr = self.discr
d = discr.dimensions
dims = range(d)
given = self.plan.given
par = self.plan.parallelism
diffmat_data = self.gpu_diffmats(diff_op_cls, elgroup)
elgroup, = discr.element_groups
float_type = given.float_type
f_decl = CudaGlobal(FunctionDeclaration(Value("void", "apply_diff_mat"),
[Pointer(POD(numpy.uint8, "gmem_diff_rst_mat")),
#Pointer(POD(float_type, "debugbuf")),
] + [Pointer(POD(float_type, "drst%d_global" % i)) for i in dims]
))
rst_channels = given.devdata.make_valid_tex_channel_count(d)
cmod = Module([
Include("pycuda-helpers.hpp"),
Line(),
Value("texture<fp_tex_%s, 1, cudaReadModeElementType>"
% dtype_to_ctype(float_type),
"field_tex"),
Line(),
Define("DIMENSIONS", discr.dimensions),
Define("DOFS_PER_EL", given.dofs_per_el()),
Line(),
Define("SEGMENT_DOF", "threadIdx.x"),
Define("PAR_MB_NR", "threadIdx.y"),
Line(),
Define("MB_SEGMENT", "blockIdx.x"),
Define("MACROBLOCK_NR", "blockIdx.y"),
Line(),
Define("DOFS_PER_SEGMENT", self.plan.segment_size),
Define("SEGMENTS_PER_MB", self.plan.segments_per_microblock()),
Define("ALIGNED_DOFS_PER_MB", given.microblock.aligned_floats),
Define("ELS_PER_MB", given.microblock.elements),
Line(),
Define("PAR_MB_COUNT", par.parallel),
Define("INLINE_MB_COUNT", par.inline),
Define("SEQ_MB_COUNT", par.serial),
Line(),
Define("THREAD_NUM", "(SEGMENT_DOF+PAR_MB_NR*DOFS_PER_SEGMENT)"),
Define("COALESCING_THREAD_COUNT", "(PAR_MB_COUNT*DOFS_PER_SEGMENT)"),
Line(),
Define("MB_DOF_BASE", "(MB_SEGMENT*DOFS_PER_SEGMENT)"),
Define("MB_DOF", "(MB_DOF_BASE+SEGMENT_DOF)"),
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_DOF_BASE", "(GLOBAL_MB_NR*ALIGNED_DOFS_PER_MB)"),
Line(),
Define("DIFFMAT_SEGMENT_FLOATS", diffmat_data.block_floats),
Define("DIFFMAT_SEGMENT_BYTES", "(DIFFMAT_SEGMENT_FLOATS*%d)"
% given.float_size()),
Define("DIFFMAT_COLUMNS", diffmat_data.matrix_columns),
Line(),
CudaShared(ArrayOf(POD(float_type, "smem_diff_rst_mat"),
"DIFFMAT_COLUMNS*DOFS_PER_SEGMENT")),
Line(),
])
S = Statement
f_body = Block()
f_body.extend_log_block("calculate responsibility data", [
Initializer(POD(numpy.uint16, "mb_el"),
"MB_DOF/DOFS_PER_EL"),
])
from hedge.backends.cuda.tools import get_load_code
f_body.extend(
get_load_code(
dest="smem_diff_rst_mat",
base="gmem_diff_rst_mat + MB_SEGMENT*DIFFMAT_SEGMENT_BYTES",
bytes="DIFFMAT_SEGMENT_BYTES",
descr="load diff mat segment")
+[S("__syncthreads()"), Line()])
# ---------------------------------------------------------------------
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())
def get_mat_entry(row, col, axis):
return ("smem_diff_rst_mat["
"%(row)s*DIFFMAT_COLUMNS + %(axis)s*DOFS_PER_EL"
" + %(col)s"
"]" % {"row":row, "col":col, "axis":axis}
)
tex_channels = ["x", "y", "z", "w"]
from hedge.backends.cuda.tools import unroll
code.extend(
[POD(float_type, "field_value%d" % inl)
for inl in range(par.inline)]
+[Line()]
+unroll(lambda j: [
Assign("field_value%d" % inl,
"fp_tex1Dfetch(field_tex, GLOBAL_MB_DOF_BASE + %d*ALIGNED_DOFS_PER_MB "
"+ mb_el*DOFS_PER_EL + %s)" % (inl, j)
)
for inl in range(par.inline)]
+[Line()]
+[S("d%drst%d += %s * field_value%d"
% (inl, axis, get_mat_entry("SEGMENT_DOF", j, axis), inl))
for axis in dims
for inl in range(par.inline)]
+[Line()],
given.dofs_per_el(), self.plan.max_unroll)
)
store_code = Block()
for inl in range(par.inline):
for rst_axis in dims:
store_code.append(Assign(
"drst%d_global[GLOBAL_MB_DOF_BASE"
" + %d*ALIGNED_DOFS_PER_MB + MB_DOF]" % (rst_axis, inl),
"d%drst%d" % (inl, rst_axis),
))
code.append(If("MB_DOF < DOFS_PER_EL*ELS_PER_MB", 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=10"]
)
field_texref = mod.get_texref("field_tex")
func = mod.get_function("apply_diff_mat")
func.prepare(
discr.dimensions*[float_type] + ["P"],
block=(self.plan.segment_size, par.parallel, 1),
texrefs=[field_texref])
if "cuda_diff" in discr.debug:
print "diff: lmem=%d smem=%d regs=%d" % (
func.local_size_bytes,
func.shared_size_bytes,
func.num_regs)
return func, field_texref
def get_kernel(self, fdata, ilist_data, for_benchmark):
from cgen.cuda import CudaShared, CudaGlobal
from pycuda.tools import dtype_to_ctype
discr = self.discr
given = self.plan.given
fplan = self.plan
d = discr.dimensions
dims = range(d)
elgroup, = discr.element_groups
float_type = given.float_type
f_decl = CudaGlobal(
FunctionDeclaration(Value("void", "apply_flux"), [
Pointer(POD(float_type, "debugbuf")),
Pointer(POD(numpy.uint8, "gmem_facedata")),
] + [
Pointer(POD(float_type, "gmem_fluxes_on_faces%d" % flux_nr))
for flux_nr in range(len(self.fluxes))
]))
cmod = Module()
cmod.append(Include("pycuda-helpers.hpp"))
for dep_expr in self.all_deps:
cmod.extend([
Value(
"texture<%s, 1, cudaReadModeElementType>" %
dtype_to_ctype(float_type, with_fp_tex_hack=True),
"field%d_tex" % self.dep_to_index[dep_expr])
])
if fplan.flux_count != len(self.fluxes):
from warnings import warn
warn(
"Flux count in flux execution plan different from actual flux count.\n"
"You may want to specify the tune_for= kwarg in the Discretization\n"
"constructor.")
cmod.extend([
Line(),
Typedef(POD(float_type, "value_type")),
Line(),
flux_header_struct(float_type, discr.dimensions),
Line(),
face_pair_struct(float_type, discr.dimensions),
Line(),
Define("DIMENSIONS", discr.dimensions),
Define("DOFS_PER_FACE", fplan.dofs_per_face),
Define("THREADS_PER_FACE", fplan.threads_per_face()),
Line(),
Define("CONCURRENT_FACES", fplan.parallel_faces),
Define("BLOCK_MB_COUNT", fplan.mbs_per_block),
Line(),
Define("FACEDOF_NR", "threadIdx.x"),
Define("BLOCK_FACE", "threadIdx.y"),
Line(),
Define("FLUX_COUNT", len(self.fluxes)),
Line(),
Define("THREAD_NUM", "(FACEDOF_NR + BLOCK_FACE*THREADS_PER_FACE)"),
Define("THREAD_COUNT", "(THREADS_PER_FACE*CONCURRENT_FACES)"),
Define(
"COALESCING_THREAD_COUNT",
"(THREAD_COUNT < 0x10 ? THREAD_COUNT : THREAD_COUNT & ~0xf)"),
Line(),
Define("DATA_BLOCK_SIZE", fdata.block_bytes),
Define("ALIGNED_FACE_DOFS_PER_MB",
fplan.aligned_face_dofs_per_microblock()),
Define("ALIGNED_FACE_DOFS_PER_BLOCK",
"(ALIGNED_FACE_DOFS_PER_MB*BLOCK_MB_COUNT)"),
Line(),
Define("FOF_BLOCK_BASE",
"(blockIdx.x*ALIGNED_FACE_DOFS_PER_BLOCK)"),
Line(),
] + ilist_data.code + [
Line(),
Value("texture<index_list_entry_t, 1, cudaReadModeElementType>",
"tex_index_lists"),
Line(),
fdata.struct,
Line(),
CudaShared(Value("flux_data", "data")),
])
if not fplan.direct_store:
cmod.extend([
CudaShared(
ArrayOf(
ArrayOf(POD(float_type, "smem_fluxes_on_faces"),
"FLUX_COUNT"),
"ALIGNED_FACE_DOFS_PER_MB*BLOCK_MB_COUNT")),
Line(),
])
S = Statement
f_body = Block()
from hedge.backends.cuda.tools import get_load_code
f_body.extend(
get_load_code(dest="&data",
base="gmem_facedata + blockIdx.x*DATA_BLOCK_SIZE",
bytes="sizeof(flux_data)",
descr="load face_pair data") +
[S("__syncthreads()"), Line()])
def get_flux_code(flux_writer):
flux_code = Block([])
flux_code.extend([
Initializer(Pointer(Value("face_pair", "fpair")),
"data.facepairs+fpair_nr"),
Initializer(
MaybeUnused(POD(numpy.uint32, "a_index")),
"fpair->a_base + tex1Dfetch(tex_index_lists, "
"fpair->a_ilist_index + FACEDOF_NR)"),
Initializer(
MaybeUnused(POD(numpy.uint32, "b_index")),
"fpair->b_base + tex1Dfetch(tex_index_lists, "
"fpair->b_ilist_index + FACEDOF_NR)"),
Line(),
flux_writer(),
Line(),
S("fpair_nr += CONCURRENT_FACES")
])
return flux_code
flux_computation = Block([
Comment("fluxes for dual-sided (intra-block) interior face pairs"),
While("fpair_nr < data.header.same_facepairs_end",
get_flux_code(lambda: self.write_interior_flux_code(True))),
Line(),
Comment("work around nvcc assertion failure"),
S("fpair_nr+=1"),
S("fpair_nr-=1"),
Line(),
Comment(
"fluxes for single-sided (inter-block) interior face pairs"),
While("fpair_nr < data.header.diff_facepairs_end",
get_flux_code(lambda: self.write_interior_flux_code(False))),
Line(),
Comment("fluxes for single-sided boundary face pairs"),
While(
"fpair_nr < data.header.bdry_facepairs_end",
get_flux_code(
lambda: self.write_boundary_flux_code(for_benchmark))),
])
f_body.extend_log_block("compute the fluxes", [
Initializer(POD(numpy.uint32, "fpair_nr"), "BLOCK_FACE"),
If("FACEDOF_NR < DOFS_PER_FACE", flux_computation)
])
if not fplan.direct_store:
f_body.extend([Line(), S("__syncthreads()"), Line()])
f_body.extend_log_block(
"store fluxes",
[
#Assign("debugbuf[blockIdx.x]", "FOF_BLOCK_BASE"),
#Assign("debugbuf[0]", "FOF_BLOCK_BASE"),
#Assign("debugbuf[0]", "sizeof(face_pair)"),
For(
"unsigned word_nr = THREAD_NUM",
"word_nr < ALIGNED_FACE_DOFS_PER_MB*BLOCK_MB_COUNT",
"word_nr += COALESCING_THREAD_COUNT",
Block([
Assign(
"gmem_fluxes_on_faces%d[FOF_BLOCK_BASE+word_nr]"
% flux_nr,
"smem_fluxes_on_faces[%d][word_nr]" % flux_nr)
for flux_nr in range(len(self.fluxes))
]
#+[If("isnan(smem_fluxes_on_faces[%d][word_nr])" % flux_nr,
#Block([
#Assign("debugbuf[blockIdx.x]", "word_nr"),
#])
#)
#for flux_nr in range(len(self.fluxes))]
))
])
if False:
f_body.extend([
Assign("debugbuf[blockIdx.x*96+32+BLOCK_FACE*32+threadIdx.x]",
"fpair_nr"),
Assign("debugbuf[blockIdx.x*96+16]",
"data.header.same_facepairs_end"),
Assign("debugbuf[blockIdx.x*96+17]",
"data.header.diff_facepairs_end"),
Assign("debugbuf[blockIdx.x*96+18]",
"data.header.bdry_facepairs_end"),
])
# 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("flux_gather", ".cu").write(str(cmod))
#from pycuda.tools import allow_user_edit
mod = SourceModule(
#allow_user_edit(cmod, "kernel.cu", "the flux kernel"),
cmod,
keep="cuda_keep_kernels" in discr.debug)
expr_to_texture_map = dict(
(dep_expr,
mod.get_texref("field%d_tex" % self.dep_to_index[dep_expr]))
for dep_expr in self.all_deps)
index_list_texref = mod.get_texref("tex_index_lists")
index_list_texref.set_address(ilist_data.device_memory,
ilist_data.bytes)
index_list_texref.set_format(
cuda.dtype_to_array_format(ilist_data.type), 1)
index_list_texref.set_flags(cuda.TRSF_READ_AS_INTEGER)
func = mod.get_function("apply_flux")
block = (fplan.threads_per_face(), fplan.parallel_faces, 1)
func.prepare(
(2 + len(self.fluxes)) * "P",
texrefs=expr_to_texture_map.values() + [index_list_texref])
if "cuda_flux" in discr.debug:
print "flux: lmem=%d smem=%d regs=%d" % (
func.local_size_bytes, func.shared_size_bytes, func.num_regs)
return block, func, expr_to_texture_map
def generate_assignment_instruction_code(codegen_state, insn):
kernel = codegen_state.kernel
ecm = codegen_state.expression_to_code_mapper
from loopy.expression import dtype_to_type_context, VectorizabilityChecker
# {{{ vectorization handling
if codegen_state.vectorization_info:
if insn.atomicity:
raise Unvectorizable("atomic operation")
vinfo = codegen_state.vectorization_info
vcheck = VectorizabilityChecker(kernel, vinfo.iname, vinfo.length)
lhs_is_vector = vcheck(insn.assignee)
rhs_is_vector = vcheck(insn.expression)
if not lhs_is_vector and rhs_is_vector:
raise Unvectorizable("LHS is scalar, RHS is vector, cannot assign")
is_vector = lhs_is_vector
del lhs_is_vector
del rhs_is_vector
# }}}
from pymbolic.primitives import Variable, Subscript
from loopy.symbolic import LinearSubscript
lhs = insn.assignee
if isinstance(lhs, Variable):
assignee_var_name = lhs.name
assignee_indices = ()
elif isinstance(lhs, Subscript):
assignee_var_name = lhs.aggregate.name
assignee_indices = lhs.index_tuple
elif isinstance(lhs, LinearSubscript):
assignee_var_name = lhs.aggregate.name
assignee_indices = (lhs.index, )
else:
raise RuntimeError("invalid lvalue '%s'" % lhs)
lhs_var = kernel.get_var_descriptor(assignee_var_name)
lhs_dtype = lhs_var.dtype
if insn.atomicity is not None:
lhs_atomicity = [
a for a in insn.atomicity if a.var_name == assignee_var_name
]
assert len(lhs_atomicity) <= 1
if lhs_atomicity:
lhs_atomicity, = lhs_atomicity
else:
lhs_atomicity = None
else:
lhs_atomicity = None
from loopy.kernel.data import AtomicInit, AtomicUpdate
lhs_code = ecm(insn.assignee, prec=PREC_NONE, type_context=None)
rhs_type_context = dtype_to_type_context(kernel.target, lhs_dtype)
if lhs_atomicity is None:
result = codegen_state.ast_builder.emit_assignment(
codegen_state, lhs_code,
ecm(insn.expression,
prec=PREC_NONE,
type_context=rhs_type_context,
needed_dtype=lhs_dtype))
elif isinstance(lhs_atomicity, AtomicInit):
raise NotImplementedError("atomic init")
elif isinstance(lhs_atomicity, AtomicUpdate):
codegen_state.seen_atomic_dtypes.add(lhs_dtype)
result = codegen_state.ast_builder.generate_atomic_update(
kernel, codegen_state, lhs_atomicity, lhs_var, insn.assignee,
insn.expression, lhs_dtype, rhs_type_context)
else:
raise ValueError("unexpected lhs atomicity type: %s" %
type(lhs_atomicity).__name__)
# {{{ tracing
if kernel.options.trace_assignments or kernel.options.trace_assignment_values:
if codegen_state.vectorization_info and is_vector:
raise Unvectorizable("tracing does not support vectorization")
from cgen import Statement as S # noqa
gs, ls = kernel.get_grid_size_upper_bounds()
printf_format = "%s.%s[%s][%s]: %s" % (kernel.name, insn.id, ", ".join(
"gid%d=%%d" % i for i in range(len(gs))), ", ".join(
"lid%d=%%d" % i for i in range(len(ls))), assignee_var_name)
printf_args = (["gid(%d)" % i for i in range(len(gs))] +
["lid(%d)" % i for i in range(len(ls))])
if assignee_indices:
printf_format += "[%s]" % ",".join(len(assignee_indices) * ["%d"])
printf_args.extend(
ecm(i, prec=PREC_NONE, type_context="i")
for i in assignee_indices)
if kernel.options.trace_assignment_values:
if lhs_dtype.numpy_dtype.kind == "i":
printf_format += " = %d"
printf_args.append(lhs_code)
elif lhs_dtype.numpy_dtype.kind == "f":
printf_format += " = %g"
printf_args.append(lhs_code)
elif lhs_dtype.numpy_dtype.kind == "c":
printf_format += " = %g + %gj"
printf_args.extend(["(%s).x" % lhs_code, "(%s).y" % lhs_code])
if printf_args:
printf_args_str = ", " + ", ".join(printf_args)
else:
printf_args_str = ""
printf_insn = S("printf(\"%s\\n\"%s)" %
(printf_format, printf_args_str))
from cgen import Block
if kernel.options.trace_assignment_values:
result = Block([result, printf_insn])
else:
# print first, execute later -> helps find segfaults
result = Block([printf_insn, result])
# }}}
return result
def generate_code(self, kernel, codegen_state, impl_arg_info):
from cgen import (FunctionBody, FunctionDeclaration, Value, Module,
Block, Line, Statement as S)
from cgen.ispc import ISPCExport, ISPCTask
knl_body, implemented_domains = kernel.target.generate_body(
kernel, codegen_state)
inner_name = "lp_ispc_inner_"+kernel.name
arg_decls = [iai.cgen_declarator for iai in impl_arg_info]
arg_names = [iai.name for iai in impl_arg_info]
# {{{ occa compatibility hackery
if self.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
# }}}
knl_fbody = FunctionBody(
ISPCTask(
FunctionDeclaration(
Value("void", inner_name),
arg_decls)),
knl_body)
# {{{ generate wrapper
wrapper_body = Block()
gsize, lsize = kernel.get_grid_sizes_as_exprs()
if len(lsize) > 1:
for i, ls_i in enumerate(lsize[1:]):
if ls_i != 1:
raise LoopyError("local axis %d (0-based) "
"has length > 1, which is unsupported "
"by ISPC" % ls_i)
from pymbolic.mapper.stringifier import PREC_COMPARISON, PREC_NONE
ccm = self.get_expression_to_code_mapper(codegen_state)
wrapper_body.extend([
S("assert(programCount == %s)"
% ccm(lsize[0], PREC_COMPARISON)),
S("launch[%s] %s(%s)"
% (
", ".join(
ccm(gs_i, PREC_NONE)
for gs_i in gsize),
inner_name,
", ".join(arg_names)
))
])
wrapper_fbody = FunctionBody(
ISPCExport(
FunctionDeclaration(
Value("void", kernel.name),
arg_decls)),
wrapper_body)
# }}}
mod = Module([
knl_fbody,
Line(),
wrapper_fbody,
])
return str(mod), implemented_domains
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
def emit_atomic_update(self, codegen_state, lhs_atomicity, lhs_var,
lhs_expr, rhs_expr, lhs_dtype, rhs_type_context):
from pymbolic.primitives import Sum
from cgen import Statement
from pymbolic.mapper.stringifier import PREC_NONE
if isinstance(lhs_dtype, NumpyType) and lhs_dtype.numpy_dtype in [
np.int32, np.int64, np.float32, np.float64
]:
# atomicAdd
if isinstance(rhs_expr, Sum):
ecm = self.get_expression_to_code_mapper(codegen_state)
new_rhs_expr = Sum(
tuple(c for c in rhs_expr.children if c != lhs_expr))
lhs_expr_code = ecm(lhs_expr)
rhs_expr_code = ecm(new_rhs_expr)
return Statement("atomicAdd(&{}, {})".format(
lhs_expr_code, rhs_expr_code))
else:
from cgen import Block, DoWhile, Assign
from loopy.target.c import POD
old_val_var = codegen_state.var_name_generator("loopy_old_val")
new_val_var = codegen_state.var_name_generator("loopy_new_val")
from loopy.kernel.data import TemporaryVariable
ecm = codegen_state.expression_to_code_mapper.with_assignments(
{
old_val_var: TemporaryVariable(old_val_var, lhs_dtype),
new_val_var: TemporaryVariable(new_val_var, lhs_dtype),
})
lhs_expr_code = ecm(lhs_expr,
prec=PREC_NONE,
type_context=None)
from pymbolic.mapper.substitutor import make_subst_func
from pymbolic import var
from loopy.symbolic import SubstitutionMapper
subst = SubstitutionMapper(
make_subst_func({lhs_expr: var(old_val_var)}))
rhs_expr_code = ecm(subst(rhs_expr),
prec=PREC_NONE,
type_context=rhs_type_context,
needed_dtype=lhs_dtype)
cast_str = ""
old_val = old_val_var
new_val = new_val_var
if lhs_dtype.numpy_dtype.kind == "f":
if lhs_dtype.numpy_dtype == np.float32:
ctype = "int"
elif lhs_dtype.numpy_dtype == np.float64:
ctype = "long"
else:
raise AssertionError()
old_val = "*(%s *) &" % ctype + old_val
new_val = "*(%s *) &" % ctype + new_val
cast_str = "(%s *) " % (ctype)
return Block([
POD(self, NumpyType(lhs_dtype.dtype, target=self.target),
old_val_var),
POD(self, NumpyType(lhs_dtype.dtype, target=self.target),
new_val_var),
DoWhile(
"atomicCAS("
"%(cast_str)s&(%(lhs_expr)s), "
"%(old_val)s, "
"%(new_val)s"
") != %(old_val)s" % {
"cast_str": cast_str,
"lhs_expr": lhs_expr_code,
"old_val": old_val,
"new_val": new_val,
},
Block([
Assign(old_val_var, lhs_expr_code),
Assign(new_val_var, rhs_expr_code),
]))
])
else:
raise NotImplementedError("atomic update for '%s'" % lhs_dtype)
def get_kernel(self, with_scaling, for_benchmark=False):
from cgen import \
Pointer, POD, Value, ArrayOf, \
Module, FunctionDeclaration, FunctionBody, Block, \
Line, Define, Include, \
Initializer, If, For, Statement, Assign, \
ArrayInitializer
from cgen import dtype_to_ctype
from cgen.cuda import CudaShared, CudaConstant, CudaGlobal
discr = self.discr
d = discr.dimensions
dims = range(d)
given = self.plan.given
float_type = given.float_type
f_decl = CudaGlobal(FunctionDeclaration(Value("void", "apply_el_local_mat_smem_mat"),
[
Pointer(POD(float_type, "out_vector")),
Pointer(POD(numpy.uint8, "gmem_matrix")),
Pointer(POD(float_type, "debugbuf")),
POD(numpy.uint32, "microblock_count"),
]
))
cmod = Module([
Include("pycuda-helpers.hpp"),
Line(),
Value("texture<fp_tex_%s, 1, cudaReadModeElementType>"
% dtype_to_ctype(float_type),
"in_vector_tex"),
])
if with_scaling:
cmod.append(
Value("texture<fp_tex_%s, 1, cudaReadModeElementType>"
% dtype_to_ctype(float_type),
"scaling_tex"),
)
par = self.plan.parallelism
cmod.extend([
Line(),
Define("DIMENSIONS", discr.dimensions),
Define("DOFS_PER_EL", given.dofs_per_el()),
Define("PREIMAGE_DOFS_PER_EL", self.plan.preimage_dofs_per_el),
Line(),
Define("SEGMENT_DOF", "threadIdx.x"),
Define("PAR_MB_NR", "threadIdx.y"),
Line(),
Define("MB_SEGMENT", "blockIdx.x"),
Define("MACROBLOCK_NR", "blockIdx.y"),
Line(),
Define("DOFS_PER_SEGMENT", self.plan.segment_size),
Define("SEGMENTS_PER_MB", self.plan.segments_per_microblock()),
Define("ALIGNED_DOFS_PER_MB", given.microblock.aligned_floats),
Define("ALIGNED_PREIMAGE_DOFS_PER_MB",
self.plan.aligned_preimage_dofs_per_microblock),
Define("MB_EL_COUNT", given.microblock.elements),
Line(),
Define("PAR_MB_COUNT", par.parallel),
Define("INLINE_MB_COUNT", par.inline),
Define("SEQ_MB_COUNT", par.serial),
Line(),
Define("THREAD_NUM", "(SEGMENT_DOF+PAR_MB_NR*DOFS_PER_SEGMENT)"),
Define("COALESCING_THREAD_COUNT", "(PAR_MB_COUNT*DOFS_PER_SEGMENT)"),
Line(),
Define("MB_DOF_BASE", "(MB_SEGMENT*DOFS_PER_SEGMENT)"),
Define("MB_DOF", "(MB_DOF_BASE+SEGMENT_DOF)"),
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_DOF_BASE", "(GLOBAL_MB_NR*ALIGNED_DOFS_PER_MB)"),
Define("GLOBAL_MB_PREIMG_DOF_BASE", "(GLOBAL_MB_NR*ALIGNED_PREIMAGE_DOFS_PER_MB)"),
Line(),
Define("MATRIX_COLUMNS", self.plan.gpu_matrix_columns()),
Define("MATRIX_SEGMENT_FLOATS", self.plan.gpu_matrix_block_floats()),
Define("MATRIX_SEGMENT_BYTES",
"(MATRIX_SEGMENT_FLOATS*%d)" % given.float_size()),
Line(),
CudaShared(ArrayOf(POD(float_type, "smem_matrix"),
"MATRIX_SEGMENT_FLOATS")),
CudaShared(
ArrayOf(
ArrayOf(
ArrayOf(
POD(float_type, "dof_buffer"),
"PAR_MB_COUNT"),
"INLINE_MB_COUNT"),
"DOFS_PER_SEGMENT"),
),
CudaShared(POD(numpy.uint16, "segment_start_el")),
CudaShared(POD(numpy.uint16, "segment_stop_el")),
CudaShared(POD(numpy.uint16, "segment_el_count")),
Line(),
ArrayInitializer(
CudaConstant(
ArrayOf(
POD(numpy.uint32, "segment_start_el_lookup"),
"SEGMENTS_PER_MB")),
[(chk*self.plan.segment_size)//given.dofs_per_el()
for chk in range(self.plan.segments_per_microblock())]
),
ArrayInitializer(
CudaConstant(
ArrayOf(
POD(numpy.uint32, "segment_stop_el_lookup"),
"SEGMENTS_PER_MB")),
[min(given.microblock.elements,
(chk*self.plan.segment_size+self.plan.segment_size-1)
//given.dofs_per_el()+1)
for chk in range(self.plan.segments_per_microblock())]
),
])
S = Statement
f_body = Block()
f_body.extend_log_block("calculate this dof's element", [
Initializer(POD(numpy.uint8, "mb_el"),
"MB_DOF/DOFS_PER_EL") ])
if self.plan.use_prefetch_branch:
f_body.extend_log_block("calculate segment responsibility data", [
If("THREAD_NUM==0",
Block([
Assign("segment_start_el", "segment_start_el_lookup[MB_SEGMENT]"),
Assign("segment_stop_el", "segment_stop_el_lookup[MB_SEGMENT]"),
Assign("segment_el_count", "segment_stop_el-segment_start_el"),
])
),
S("__syncthreads()")
])
from hedge.backends.cuda.tools import get_load_code
f_body.extend(
get_load_code(
dest="smem_matrix",
base=("gmem_matrix + MB_SEGMENT*MATRIX_SEGMENT_BYTES"),
bytes="MATRIX_SEGMENT_BYTES",
descr="load matrix segment")
+[S("__syncthreads()")]
)
# ---------------------------------------------------------------------
def get_batched_fetch_mat_mul_code(el_fetch_count):
result = []
dofs = range(self.plan.preimage_dofs_per_el)
for load_segment_start in range(0, self.plan.preimage_dofs_per_el,
self.plan.segment_size):
result.extend(
[S("__syncthreads()")]
+[Assign(
"dof_buffer[PAR_MB_NR][%d][SEGMENT_DOF]" % inl,
"fp_tex1Dfetch(in_vector_tex, "
"GLOBAL_MB_PREIMG_DOF_BASE"
" + %d*ALIGNED_PREIMAGE_DOFS_PER_MB"
" + (segment_start_el)*PREIMAGE_DOFS_PER_EL + %d + SEGMENT_DOF)"
% (inl, load_segment_start)
)
for inl in range(par.inline)]
+[S("__syncthreads()"),
Line(),
])
for dof in dofs[load_segment_start:load_segment_start+self.plan.segment_size]:
for inl in range(par.inline):
result.append(
S("result%d += "
"smem_matrix[SEGMENT_DOF*MATRIX_COLUMNS + %d]"
"*"
"dof_buffer[PAR_MB_NR][%d][%d]"
% (inl, dof, inl, dof-load_segment_start))
)
result.append(Line())
return result
from hedge.backends.cuda.tools import unroll
def get_direct_tex_mat_mul_code():
return (
[POD(float_type, "fof%d" % inl) for inl in range(par.inline)]
+ [POD(float_type, "lm"), Line()]
+ unroll(
lambda j: [
Assign("fof%d" % inl,
"fp_tex1Dfetch(in_vector_tex, "
"GLOBAL_MB_PREIMG_DOF_BASE"
" + %(inl)d * ALIGNED_PREIMAGE_DOFS_PER_MB"
" + mb_el*PREIMAGE_DOFS_PER_EL+%(j)s)"
% {"j":j, "inl":inl, "row": "SEGMENT_DOF"},)
for inl in range(par.inline)
]+[
Assign("lm",
"smem_matrix["
"%(row)s*MATRIX_COLUMNS + %(j)s]"
% {"j":j, "row": "SEGMENT_DOF"},
)
]+[
S("result%(inl)d += fof%(inl)d*lm" % {"inl":inl})
for inl in range(par.inline)
],
total_number=self.plan.preimage_dofs_per_el,
max_unroll=self.plan.max_unroll)
+ [Line()])
def get_mat_mul_code(el_fetch_count):
if el_fetch_count == 1:
return get_batched_fetch_mat_mul_code(el_fetch_count)
else:
return get_direct_tex_mat_mul_code()
def mat_mul_outer_loop(fetch_count):
if with_scaling:
inv_jac_multiplier = ("fp_tex1Dfetch(scaling_tex,"
"(GLOBAL_MB_NR + %(inl)d)*MB_EL_COUNT + mb_el)")
else:
inv_jac_multiplier = "1"
write_condition = "MB_DOF < DOFS_PER_EL*MB_EL_COUNT"
if self.with_index_check:
write_condition += " && GLOBAL_MB_NR < microblock_count"
return For("unsigned short seq_mb_number = 0",
"seq_mb_number < SEQ_MB_COUNT",
"++seq_mb_number",
Block([
Initializer(POD(float_type, "result%d" % inl), 0)
for inl in range(par.inline)
]+[Line()]
+get_mat_mul_code(fetch_count)
+[
If(write_condition,
Block([
Assign(
"out_vector[GLOBAL_MB_DOF_BASE"
" + %d*ALIGNED_DOFS_PER_MB"
" + MB_DOF]" % inl,
"result%d * %s" % (inl, (inv_jac_multiplier % {"inl":inl}))
)
for inl in range(par.inline)
])
)
])
)
if self.plan.use_prefetch_branch:
from cgen import make_multiple_ifs
f_body.append(make_multiple_ifs([
("segment_el_count == %d" % fetch_count,
mat_mul_outer_loop(fetch_count))
for fetch_count in
range(1, self.plan.max_elements_touched_by_segment()+1)]
))
else:
f_body.append(mat_mul_outer_loop(0))
# 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(self.plan.debug_name, ".cu").write(str(cmod))
mod = SourceModule(cmod,
keep="cuda_keep_kernels" in discr.debug,
#options=["--maxrregcount=12"]
)
func = mod.get_function("apply_el_local_mat_smem_mat")
if self.plan.debug_name in discr.debug:
print "%s: lmem=%d smem=%d regs=%d" % (
self.plan.debug_name,
func.local_size_bytes,
func.shared_size_bytes,
func.num_regs)
in_vector_texref = mod.get_texref("in_vector_tex")
texrefs = [in_vector_texref]
if with_scaling:
scaling_texref = mod.get_texref("scaling_tex")
texrefs.append(scaling_texref)
else:
scaling_texref = None
func.prepare(
"PPPI",
block=(self.plan.segment_size, self.plan.parallelism.parallel, 1),
texrefs=texrefs)
return func, in_vector_texref, scaling_texref
def get_kernel(self, for_benchmark=False):
from cgen import \
Pointer, POD, Value, ArrayOf, Const, \
Module, FunctionDeclaration, FunctionBody, Block, \
Comment, Line, Include, \
Define, \
Initializer, If, For, Statement, Assign
from cgen import dtype_to_ctype
from cgen.cuda import CudaShared, CudaGlobal
discr = self.discr
d = discr.dimensions
dims = range(d)
given = self.plan.given
float_type = given.float_type
f_decl = CudaGlobal(FunctionDeclaration(Value("void", "apply_el_local_mat_smem_field"),
[
Pointer(POD(float_type, "out_vector")),
Pointer(POD(float_type, "in_vector")),
Pointer(POD(float_type, "debugbuf")),
POD(numpy.uint32, "microblock_count"),
]
))
cmod = Module([
Include("pycuda-helpers.hpp"),
Line(),
Value("texture<fp_tex_%s, 2, cudaReadModeElementType>"
% dtype_to_ctype(float_type),
"mat_tex"),
])
plan = self.plan
par = plan.parallelism
# 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),
Line(),
Define("MB_EL_COUNT", plan.elements_per_microblock),
Line(),
Define("IMAGE_DOFS_PER_MB", "(IMAGE_DOFS_PER_EL*MB_EL_COUNT)"),
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_in_vector"),
"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,
"in_vector[GLOBAL_MB_PREIMAGE_DOF_BASE"
" + %d*ALIGNED_PREIMAGE_DOFS_PER_MB"
" + %s]" % (inl, block_addr))
store_instr = Assign(
"smem_in_vector[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_matmul_code():
from hedge.backends.cuda.tools import unroll
index_check_condition = "GLOBAL_MB_NR < microblock_count"
def if_(conditions, then):
final_cond = " && ".join(cond for cond in conditions if cond)
if final_cond:
return If(final_cond, then)
else:
return then
result = Block([
Comment("everybody needs to be done with the old data"),
S("__syncthreads()"), Line(),
]+[If(index_check_condition, get_load_code())]+[
Line(),
Comment("all the new data must be loaded"),
S("__syncthreads()"),
Line(),
]+[
Initializer(POD(float_type, "result%d" % inl), 0)
for inl in range(par.inline)
]+[
Line(),
POD(float_type, "mat_entry"),
Line(),
])
result.append(if_(["IMAGE_MB_DOF < IMAGE_DOFS_PER_MB", index_check_condition],
Block(unroll(lambda j:
[Assign("mat_entry", "fp_tex2D(mat_tex, IMAGE_EL_DOF, %s)" % j)]
+[
S("result%d += mat_entry "
"* smem_in_vector[PAR_MB_NR][%d][mb_el*PREIMAGE_DOFS_PER_EL + %s]"
% (inl, inl, j))
for inl in range(par.inline)
],
total_number=plan.preimage_dofs_per_el)
+[Line()]
+[Assign(
"out_vector[GLOBAL_MB_IMAGE_DOF_BASE + "
"%d*ALIGNED_IMAGE_DOFS_PER_MB + IMAGE_MB_DOF]" % inl,
"result%d" % inl)
for inl in range(par.inline)]
)))
return result
f_body.append(For("unsigned short seq_mb_number = 0",
"seq_mb_number < SEQ_MB_COUNT",
"++seq_mb_number", get_matmul_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(plan.debug_name, ".cu").write(str(cmod))
mod = SourceModule(cmod,
keep="cuda_keep_kernels" in discr.debug,
#options=["--maxrregcount=12"]
)
func = mod.get_function("apply_el_local_mat_smem_field")
if plan.debug_name in discr.debug:
print "%s: lmem=%d smem=%d regs=%d" % (
plan.debug_name,
func.local_size_bytes,
func.shared_size_bytes,
func.num_regs)
mat_texref = mod.get_texref("mat_tex")
texrefs = [mat_texref]
func.prepare(
"PPPI",
texrefs=texrefs)
assert plan.aligned_image_dofs_per_microblock % chunk_size == 0
block = (
chunk_size,
plan.parallelism.parallel,
plan.aligned_image_dofs_per_microblock
//chunk_size)
return func, block, mat_texref
def write_boundary_flux_code(self, for_benchmark):
given = self.plan.given
flux_write_code = Block()
fluxes_by_bdry_number = {}
for flux_nr, wdflux in enumerate(self.fluxes):
for bflux_info in wdflux.boundaries:
if for_benchmark:
bdry_number = 0
else:
bdry_number = self.executor.boundary_tag_to_number[
bflux_info.bpair.tag]
fluxes_by_bdry_number.setdefault(bdry_number, [])\
.append((flux_nr, bflux_info))
flux_write_code.extend([
Initializer(MaybeUnused(POD(given.float_type, "flux%d" % flux_nr)),
0) for flux_nr in range(len(self.fluxes))
])
for bdry_number, nrs_and_fluxes in fluxes_by_bdry_number.iteritems():
bblock = []
from pytools import set_sum
int_deps = set_sum(flux_rec.int_dependencies
for flux_nr, flux_rec in nrs_and_fluxes)
ext_deps = set_sum(flux_rec.ext_dependencies
for flux_nr, flux_rec in nrs_and_fluxes)
for dep in int_deps:
bblock.extend([
Comment(str(dep)),
Initializer(
MaybeUnused(
POD(given.float_type,
"val_a_field%d" % self.dep_to_index[dep])),
"fp_tex1Dfetch(field%d_tex, a_index)" %
self.dep_to_index[dep])
])
for dep in ext_deps:
bblock.extend([
Comment(str(dep)),
Initializer(
MaybeUnused(
POD(given.float_type,
"val_b_field%d" % self.dep_to_index[dep])),
"fp_tex1Dfetch(field%s_tex, b_index)" %
self.dep_to_index[dep])
])
f2cm = FluxToCodeMapper(given.float_type)
comp_code = [Line()]
for flux_nr, flux_rec in nrs_and_fluxes:
comp_code.append(
Statement(
("flux%d += " % flux_nr) +
flux_to_code(f2cm,
is_flipped=False,
int_field_expr=flux_rec.bpair.field,
ext_field_expr=flux_rec.bpair.bfield,
dep_to_index=self.dep_to_index,
flux=flux_rec.flux_expr,
prec=PREC_NONE)))
if f2cm.cse_name_list:
bblock.append(Line())
bblock.extend(
Initializer(Value("value_type", cse_name), cse_str)
for cse_name, cse_str in f2cm.cse_name_list)
flux_write_code.extend([
Line(),
Comment(nrs_and_fluxes[0][1].bpair.tag),
If("(fpair->boundary_bitmap) & (1 << %d)" % (bdry_number),
Block(bblock + comp_code)),
])
flux_write_code.extend(
[
Line(),
] + [
self.gen_store(flux_nr, "fpair->a_dest+FACEDOF_NR",
"fpair->face_jacobian * flux%d" % flux_nr)
for flux_nr in range(len(self.fluxes))
]
#Assign("debugbuf[blockIdx.x*96+fpair_nr]", "10000+fpair->a_dest"),
)
return flux_write_code
def get_kernel(self, fdata, ilist_data, for_benchmark):
from cgen.cuda import CudaShared, CudaGlobal
from pycuda.tools import dtype_to_ctype
discr = self.discr
given = self.plan.given
fplan = self.plan
d = discr.dimensions
dims = range(d)
elgroup, = discr.element_groups
float_type = given.float_type
f_decl = CudaGlobal(FunctionDeclaration(Value("void", "apply_flux"),
[
Pointer(POD(float_type, "debugbuf")),
Pointer(POD(numpy.uint8, "gmem_facedata")),
]+[
Pointer(POD(float_type, "gmem_fluxes_on_faces%d" % flux_nr))
for flux_nr in range(len(self.fluxes))
]
))
cmod = Module()
cmod.append(Include("pycuda-helpers.hpp"))
for dep_expr in self.all_deps:
cmod.extend([
Value("texture<%s, 1, cudaReadModeElementType>"
% dtype_to_ctype(float_type, with_fp_tex_hack=True),
"field%d_tex" % self.dep_to_index[dep_expr])
])
if fplan.flux_count != len(self.fluxes):
from warnings import warn
warn("Flux count in flux execution plan different from actual flux count.\n"
"You may want to specify the tune_for= kwarg in the Discretization\n"
"constructor.")
cmod.extend([
Line(),
Typedef(POD(float_type, "value_type")),
Line(),
flux_header_struct(float_type, discr.dimensions),
Line(),
face_pair_struct(float_type, discr.dimensions),
Line(),
Define("DIMENSIONS", discr.dimensions),
Define("DOFS_PER_FACE", fplan.dofs_per_face),
Define("THREADS_PER_FACE", fplan.threads_per_face()),
Line(),
Define("CONCURRENT_FACES", fplan.parallel_faces),
Define("BLOCK_MB_COUNT", fplan.mbs_per_block),
Line(),
Define("FACEDOF_NR", "threadIdx.x"),
Define("BLOCK_FACE", "threadIdx.y"),
Line(),
Define("FLUX_COUNT", len(self.fluxes)),
Line(),
Define("THREAD_NUM", "(FACEDOF_NR + BLOCK_FACE*THREADS_PER_FACE)"),
Define("THREAD_COUNT", "(THREADS_PER_FACE*CONCURRENT_FACES)"),
Define("COALESCING_THREAD_COUNT",
"(THREAD_COUNT < 0x10 ? THREAD_COUNT : THREAD_COUNT & ~0xf)"),
Line(),
Define("DATA_BLOCK_SIZE", fdata.block_bytes),
Define("ALIGNED_FACE_DOFS_PER_MB", fplan.aligned_face_dofs_per_microblock()),
Define("ALIGNED_FACE_DOFS_PER_BLOCK",
"(ALIGNED_FACE_DOFS_PER_MB*BLOCK_MB_COUNT)"),
Line(),
Define("FOF_BLOCK_BASE", "(blockIdx.x*ALIGNED_FACE_DOFS_PER_BLOCK)"),
Line(),
] + ilist_data.code + [
Line(),
Value("texture<index_list_entry_t, 1, cudaReadModeElementType>",
"tex_index_lists"),
Line(),
fdata.struct,
Line(),
CudaShared(Value("flux_data", "data")),
])
if not fplan.direct_store:
cmod.extend([
CudaShared(
ArrayOf(
ArrayOf(
POD(float_type, "smem_fluxes_on_faces"),
"FLUX_COUNT"),
"ALIGNED_FACE_DOFS_PER_MB*BLOCK_MB_COUNT")
),
Line(),
])
S = Statement
f_body = Block()
from hedge.backends.cuda.tools import get_load_code
f_body.extend(get_load_code(
dest="&data",
base="gmem_facedata + blockIdx.x*DATA_BLOCK_SIZE",
bytes="sizeof(flux_data)",
descr="load face_pair data")
+[S("__syncthreads()"), Line() ])
def get_flux_code(flux_writer):
flux_code = Block([])
flux_code.extend([
Initializer(Pointer(
Value("face_pair", "fpair")),
"data.facepairs+fpair_nr"),
Initializer(
MaybeUnused(POD(numpy.uint32, "a_index")),
"fpair->a_base + tex1Dfetch(tex_index_lists, "
"fpair->a_ilist_index + FACEDOF_NR)"),
Initializer(
MaybeUnused(POD(numpy.uint32, "b_index")),
"fpair->b_base + tex1Dfetch(tex_index_lists, "
"fpair->b_ilist_index + FACEDOF_NR)"),
Line(),
flux_writer(),
Line(),
S("fpair_nr += CONCURRENT_FACES")
])
return flux_code
flux_computation = Block([
Comment("fluxes for dual-sided (intra-block) interior face pairs"),
While("fpair_nr < data.header.same_facepairs_end",
get_flux_code(lambda:
self.write_interior_flux_code(True))
),
Line(),
Comment("work around nvcc assertion failure"),
S("fpair_nr+=1"),
S("fpair_nr-=1"),
Line(),
Comment("fluxes for single-sided (inter-block) interior face pairs"),
While("fpair_nr < data.header.diff_facepairs_end",
get_flux_code(lambda:
self.write_interior_flux_code(False))
),
Line(),
Comment("fluxes for single-sided boundary face pairs"),
While("fpair_nr < data.header.bdry_facepairs_end",
get_flux_code(
lambda: self.write_boundary_flux_code(for_benchmark))
),
])
f_body.extend_log_block("compute the fluxes", [
Initializer(POD(numpy.uint32, "fpair_nr"), "BLOCK_FACE"),
If("FACEDOF_NR < DOFS_PER_FACE", flux_computation)
])
if not fplan.direct_store:
f_body.extend([
Line(),
S("__syncthreads()"),
Line()
])
f_body.extend_log_block("store fluxes", [
#Assign("debugbuf[blockIdx.x]", "FOF_BLOCK_BASE"),
#Assign("debugbuf[0]", "FOF_BLOCK_BASE"),
#Assign("debugbuf[0]", "sizeof(face_pair)"),
For("unsigned word_nr = THREAD_NUM",
"word_nr < ALIGNED_FACE_DOFS_PER_MB*BLOCK_MB_COUNT",
"word_nr += COALESCING_THREAD_COUNT",
Block([Assign(
"gmem_fluxes_on_faces%d[FOF_BLOCK_BASE+word_nr]" % flux_nr,
"smem_fluxes_on_faces[%d][word_nr]" % flux_nr)
for flux_nr in range(len(self.fluxes))]
#+[If("isnan(smem_fluxes_on_faces[%d][word_nr])" % flux_nr,
#Block([
#Assign("debugbuf[blockIdx.x]", "word_nr"),
#])
#)
#for flux_nr in range(len(self.fluxes))]
)
)
])
if False:
f_body.extend([
Assign("debugbuf[blockIdx.x*96+32+BLOCK_FACE*32+threadIdx.x]", "fpair_nr"),
Assign("debugbuf[blockIdx.x*96+16]", "data.header.same_facepairs_end"),
Assign("debugbuf[blockIdx.x*96+17]", "data.header.diff_facepairs_end"),
Assign("debugbuf[blockIdx.x*96+18]", "data.header.bdry_facepairs_end"),
]
)
# 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("flux_gather", ".cu").write(str(cmod))
#from pycuda.tools import allow_user_edit
mod = SourceModule(
#allow_user_edit(cmod, "kernel.cu", "the flux kernel"),
cmod,
keep="cuda_keep_kernels" in discr.debug)
expr_to_texture_map = dict(
(dep_expr, mod.get_texref(
"field%d_tex" % self.dep_to_index[dep_expr]))
for dep_expr in self.all_deps)
index_list_texref = mod.get_texref("tex_index_lists")
index_list_texref.set_address(
ilist_data.device_memory,
ilist_data.bytes)
index_list_texref.set_format(
cuda.dtype_to_array_format(ilist_data.type), 1)
index_list_texref.set_flags(cuda.TRSF_READ_AS_INTEGER)
func = mod.get_function("apply_flux")
block = (fplan.threads_per_face(), fplan.parallel_faces, 1)
func.prepare(
(2+len(self.fluxes))*"P",
texrefs=expr_to_texture_map.values()
+ [index_list_texref])
if "cuda_flux" in discr.debug:
print "flux: lmem=%d smem=%d regs=%d" % (
func.local_size_bytes,
func.shared_size_bytes,
func.num_regs)
return block, func, expr_to_texture_map
def write_boundary_flux_code(self, for_benchmark):
given = self.plan.given
flux_write_code = Block()
fluxes_by_bdry_number = {}
for flux_nr, wdflux in enumerate(self.fluxes):
for bflux_info in wdflux.boundaries:
if for_benchmark:
bdry_number = 0
else:
bdry_number = self.executor.boundary_tag_to_number[
bflux_info.bpair.tag]
fluxes_by_bdry_number.setdefault(bdry_number, [])\
.append((flux_nr, bflux_info))
flux_write_code.extend([
Initializer(
MaybeUnused(POD(given.float_type, "flux%d" % flux_nr)),
0)
for flux_nr in range(len(self.fluxes))])
for bdry_number, nrs_and_fluxes in fluxes_by_bdry_number.iteritems():
bblock = []
from pytools import set_sum
int_deps = set_sum(flux_rec.int_dependencies
for flux_nr, flux_rec in nrs_and_fluxes)
ext_deps = set_sum(flux_rec.ext_dependencies
for flux_nr, flux_rec in nrs_and_fluxes)
for dep in int_deps:
bblock.extend([
Comment(str(dep)),
Initializer(
MaybeUnused(POD(given.float_type, "val_a_field%d"
% self.dep_to_index[dep])),
"fp_tex1Dfetch(field%d_tex, a_index)" % self.dep_to_index[dep])
])
for dep in ext_deps:
bblock.extend([
Comment(str(dep)),
Initializer(
MaybeUnused(POD(given.float_type, "val_b_field%d"
% self.dep_to_index[dep])),
"fp_tex1Dfetch(field%s_tex, b_index)" % self.dep_to_index[dep])
])
f2cm = FluxToCodeMapper(given.float_type)
comp_code = [Line()]
for flux_nr, flux_rec in nrs_and_fluxes:
comp_code.append(
Statement(("flux%d += " % flux_nr) +
flux_to_code(f2cm, is_flipped=False,
int_field_expr=flux_rec.bpair.field,
ext_field_expr=flux_rec.bpair.bfield,
dep_to_index=self.dep_to_index,
flux=flux_rec.flux_expr, prec=PREC_NONE)))
if f2cm.cse_name_list:
bblock.append(Line())
bblock.extend(
Initializer(
Value("value_type", cse_name), cse_str)
for cse_name, cse_str in f2cm.cse_name_list)
flux_write_code.extend([
Line(),
Comment(nrs_and_fluxes[0][1].bpair.tag),
If("(fpair->boundary_bitmap) & (1 << %d)" % (bdry_number),
Block(bblock+comp_code)),
])
flux_write_code.extend([Line(),]
+[
self.gen_store(flux_nr, "fpair->a_dest+FACEDOF_NR",
"fpair->face_jacobian * flux%d" % flux_nr)
for flux_nr in range(len(self.fluxes))
]
#Assign("debugbuf[blockIdx.x*96+fpair_nr]", "10000+fpair->a_dest"),
)
return flux_write_code
def write_interior_flux_code(self, is_twosided):
given = self.plan.given
def get_field(flux_rec, is_interior, flipped):
if is_interior ^ flipped:
prefix = "a"
else:
prefix = "b"
return ("val_%s_field%d" % (prefix, self.dep_to_index[flux_rec.field_expr]))
flux_write_code = Block([])
flux_var_decl = [Initializer(POD(given.float_type, "a_flux"), 0)]
if is_twosided:
flux_var_decl.append(Initializer(POD(given.float_type, "b_flux"), 0))
prefixes = ["a", "b"]
flip_values = [False, True]
else:
prefixes = ["a"]
flip_values = [False]
flux_write_code.append(Line())
for dep in self.interior_deps:
flux_write_code.append(Comment(str(dep)))
for side in ["a", "b"]:
flux_write_code.append(
Initializer(
MaybeUnused(POD(given.float_type, "val_%s_field%d"
% (side, self.dep_to_index[dep]))),
"fp_tex1Dfetch(field%d_tex, %s_index)"
% (self.dep_to_index[dep], side)))
f2cm = FluxToCodeMapper(given.float_type)
flux_sub_codes = []
for flux_nr, wdflux in enumerate(self.fluxes):
my_flux_block = Block(flux_var_decl)
for int_rec in wdflux.interiors:
for prefix, is_flipped in zip(prefixes, flip_values):
my_flux_block.append(
Statement("%s_flux += %s"
% (prefix,
flux_to_code(f2cm, is_flipped,
int_rec.field_expr,
int_rec.field_expr,
self.dep_to_index,
int_rec.flux_expr, PREC_NONE),
)))
my_flux_block.append(Line())
my_flux_block.append(
self.gen_store(flux_nr, "fpair->a_dest+FACEDOF_NR",
"fpair->face_jacobian*a_flux"))
#my_flux_block.append(
#Statement("if(isnan(val_b_field5)) debugbuf[blockIdx.x] = 1"),
#)
if is_twosided:
my_flux_block.append(
self.gen_store(flux_nr,
"fpair->b_dest+tex1Dfetch(tex_index_lists, "
"fpair->b_write_ilist_index + FACEDOF_NR)",
"fpair->face_jacobian*b_flux"))
#my_flux_block.append(
#Assign("debugbuf[blockIdx.x*96+fpair_nr+8]", "10000+fpair->b_dest"),
#)
flux_sub_codes.append(my_flux_block)
if f2cm.cse_name_list:
flux_write_code.append(Line())
flux_write_code.extend(
Initializer(
Value("value_type", cse_name), cse_str)
for cse_name, cse_str in f2cm.cse_name_list)
flux_write_code.extend(flux_sub_codes)
return flux_write_code
def emit_atomic_update(self, codegen_state, lhs_atomicity, lhs_var,
lhs_expr, rhs_expr, lhs_dtype, rhs_type_context):
from pymbolic.mapper.stringifier import PREC_NONE
# FIXME: Could detect operations, generate atomic_{add,...} when
# appropriate.
if isinstance(lhs_dtype, NumpyType) and lhs_dtype.numpy_dtype in [
np.int32, np.int64, np.float32, np.float64
]:
from cgen import Block, DoWhile, Assign
from loopy.target.c import POD
old_val_var = codegen_state.var_name_generator("loopy_old_val")
new_val_var = codegen_state.var_name_generator("loopy_new_val")
from loopy.kernel.data import TemporaryVariable, AddressSpace
ecm = codegen_state.expression_to_code_mapper.with_assignments({
old_val_var:
TemporaryVariable(old_val_var, lhs_dtype),
new_val_var:
TemporaryVariable(new_val_var, lhs_dtype),
})
lhs_expr_code = ecm(lhs_expr, prec=PREC_NONE, type_context=None)
from pymbolic.mapper.substitutor import make_subst_func
from pymbolic import var
from loopy.symbolic import SubstitutionMapper
subst = SubstitutionMapper(
make_subst_func({lhs_expr: var(old_val_var)}))
rhs_expr_code = ecm(subst(rhs_expr),
prec=PREC_NONE,
type_context=rhs_type_context,
needed_dtype=lhs_dtype)
if lhs_dtype.numpy_dtype.itemsize == 4:
func_name = "atomic_cmpxchg"
elif lhs_dtype.numpy_dtype.itemsize == 8:
func_name = "atom_cmpxchg"
else:
raise LoopyError("unexpected atomic size")
cast_str = ""
old_val = old_val_var
new_val = new_val_var
if lhs_dtype.numpy_dtype.kind == "f":
if lhs_dtype.numpy_dtype == np.float32:
ctype = "int"
elif lhs_dtype.numpy_dtype == np.float64:
ctype = "long"
else:
assert False
from loopy.kernel.data import (TemporaryVariable, ArrayArg)
if (isinstance(lhs_var, ArrayArg)
and lhs_var.address_space == AddressSpace.GLOBAL):
var_kind = "__global"
elif (isinstance(lhs_var, ArrayArg)
and lhs_var.address_space == AddressSpace.LOCAL):
var_kind = "__local"
elif (isinstance(lhs_var, TemporaryVariable)
and lhs_var.address_space == AddressSpace.LOCAL):
var_kind = "__local"
elif (isinstance(lhs_var, TemporaryVariable)
and lhs_var.address_space == AddressSpace.GLOBAL):
var_kind = "__global"
else:
raise LoopyError("unexpected kind of variable '%s' in "
"atomic operation: " %
(lhs_var.name, type(lhs_var).__name__))
old_val = "*(%s *) &" % ctype + old_val
new_val = "*(%s *) &" % ctype + new_val
cast_str = "(%s %s *) " % (var_kind, ctype)
return Block([
POD(self, NumpyType(lhs_dtype.dtype, target=self.target),
old_val_var),
POD(self, NumpyType(lhs_dtype.dtype, target=self.target),
new_val_var),
DoWhile(
"%(func_name)s("
"%(cast_str)s&(%(lhs_expr)s), "
"%(old_val)s, "
"%(new_val)s"
") != %(old_val)s" % {
"func_name": func_name,
"cast_str": cast_str,
"lhs_expr": lhs_expr_code,
"old_val": old_val,
"new_val": new_val,
},
Block([
Assign(old_val_var, lhs_expr_code),
Assign(new_val_var, rhs_expr_code),
]))
])
else:
raise NotImplementedError("atomic update for '%s'" % lhs_dtype)
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
def generate_body(self, kernel, codegen_state):
from cgen import Block
body = Block()
temp_decls = []
# {{{ declare temporaries
base_storage_sizes = {}
base_storage_to_is_local = {}
base_storage_to_align_bytes = {}
from cgen import ArrayOf, Pointer, Initializer, AlignedAttribute
from loopy.codegen import POD # uses the correct complex type
class ConstRestrictPointer(Pointer):
def get_decl_pair(self):
sub_tp, sub_decl = self.subdecl.get_decl_pair()
return sub_tp, ("*const restrict %s" % sub_decl)
for tv in sorted(
six.itervalues(kernel.temporary_variables),
key=lambda tv: tv.name):
decl_info = tv.decl_info(self, index_dtype=kernel.index_dtype)
if not tv.base_storage:
for idi in decl_info:
temp_var_decl = POD(self, idi.dtype, idi.name)
if idi.shape:
temp_var_decl = ArrayOf(temp_var_decl,
" * ".join(str(s) for s in idi.shape))
temp_decls.append(
self.wrap_temporary_decl(temp_var_decl, tv.is_local))
else:
offset = 0
base_storage_sizes.setdefault(tv.base_storage, []).append(
tv.nbytes)
base_storage_to_is_local.setdefault(tv.base_storage, []).append(
tv.is_local)
align_size = tv.dtype.itemsize
from loopy.kernel.array import VectorArrayDimTag
for dim_tag, axis_len in zip(tv.dim_tags, tv.shape):
if isinstance(dim_tag, VectorArrayDimTag):
align_size *= axis_len
base_storage_to_align_bytes.setdefault(tv.base_storage, []).append(
align_size)
for idi in decl_info:
cast_decl = POD(self, idi.dtype, "")
temp_var_decl = POD(self, idi.dtype, idi.name)
cast_decl = self.wrap_temporary_decl(cast_decl, tv.is_local)
temp_var_decl = self.wrap_temporary_decl(
temp_var_decl, tv.is_local)
# The 'restrict' part of this is a complete lie--of course
# all these temporaries are aliased. But we're promising to
# not use them to shovel data from one representation to the
# other. That counts, right?
cast_decl = ConstRestrictPointer(cast_decl)
temp_var_decl = ConstRestrictPointer(temp_var_decl)
cast_tp, cast_d = cast_decl.get_decl_pair()
temp_var_decl = Initializer(
temp_var_decl,
"(%s %s) (%s + %s)" % (
" ".join(cast_tp), cast_d,
tv.base_storage,
offset))
temp_decls.append(temp_var_decl)
from pytools import product
offset += (
idi.dtype.itemsize
* product(si for si in idi.shape))
for bs_name, bs_sizes in sorted(six.iteritems(base_storage_sizes)):
bs_var_decl = POD(self, np.int8, bs_name)
bs_var_decl = self.wrap_temporary_decl(
bs_var_decl, base_storage_to_is_local[bs_name])
bs_var_decl = ArrayOf(bs_var_decl, max(bs_sizes))
alignment = max(base_storage_to_align_bytes[bs_name])
bs_var_decl = AlignedAttribute(alignment, bs_var_decl)
body.append(bs_var_decl)
body.extend(temp_decls)
# }}}
from loopy.codegen.loop import set_up_hw_parallel_loops
gen_code = set_up_hw_parallel_loops(kernel, 0, codegen_state)
from cgen import Line
body.append(Line())
if isinstance(gen_code.ast, Block):
body.extend(gen_code.ast.contents)
else:
body.append(gen_code.ast)
return body, gen_code.implemented_domains
def _cusp_solver(M, parameters):
cache_key = lambda t, p: (t, p['ksp_type'], p['pc_type'], p['ksp_rtol'], p[
'ksp_atol'], p['ksp_max_it'], p['ksp_gmres_restart'], p['ksp_monitor'])
module = _cusp_cache.get(cache_key(M.ctype, parameters))
if module:
return module
import codepy.toolchain
from cgen import FunctionBody, FunctionDeclaration
from cgen import Block, Statement, Include, Value
from codepy.bpl import BoostPythonModule
from codepy.cuda import CudaModule
gcc_toolchain = codepy.toolchain.guess_toolchain()
nvcc_toolchain = codepy.toolchain.guess_nvcc_toolchain()
if 'CUSP_HOME' in os.environ:
nvcc_toolchain.add_library('cusp', [os.environ['CUSP_HOME']], [], [])
host_mod = BoostPythonModule()
nvcc_mod = CudaModule(host_mod)
nvcc_includes = [
'thrust/device_vector.h', 'thrust/fill.h', 'cusp/csr_matrix.h',
'cusp/krylov/cg.h', 'cusp/krylov/bicgstab.h', 'cusp/krylov/gmres.h',
'cusp/precond/diagonal.h', 'cusp/precond/smoothed_aggregation.h',
'cusp/precond/ainv.h', 'string'
]
nvcc_mod.add_to_preamble([Include(s) for s in nvcc_includes])
nvcc_mod.add_to_preamble([Statement('using namespace std')])
# We're translating PETSc preconditioner types to CUSP
diag = Statement(
'cusp::precond::diagonal< ValueType, cusp::device_memory >M(A)')
ainv = Statement(
'cusp::precond::scaled_bridson_ainv< ValueType, cusp::device_memory >M(A)'
)
amg = Statement(
'cusp::precond::smoothed_aggregation< IndexType, ValueType, cusp::device_memory >M(A)'
)
none = Statement(
'cusp::identity_operator< ValueType, cusp::device_memory >M(nrows, ncols)'
)
preconditioners = {
'diagonal': diag,
'jacobi': diag,
'ainv': ainv,
'ainvcusp': ainv,
'amg': amg,
'hypre': amg,
'none': none,
None: none
}
try:
precond_call = preconditioners[parameters['pc_type']]
except KeyError:
raise RuntimeError("Cusp does not support preconditioner type %s" %
parameters['pc_type'])
solvers = {
'cg':
Statement('cusp::krylov::cg(A, x, b, monitor, M)'),
'bicgstab':
Statement('cusp::krylov::bicgstab(A, x, b, monitor, M)'),
'gmres':
Statement(
'cusp::krylov::gmres(A, x, b, %(ksp_gmres_restart)d, monitor, M)' %
parameters)
}
try:
solve_call = solvers[parameters['ksp_type']]
except KeyError:
raise RuntimeError("Cusp does not support solver type %s" %
parameters['ksp_type'])
monitor = 'monitor(b, %(ksp_max_it)d, %(ksp_rtol)g, %(ksp_atol)g)' % parameters
nvcc_function = FunctionBody(
FunctionDeclaration(Value('void', '__cusp_solve'), [
Value('CUdeviceptr', '_rowptr'),
Value('CUdeviceptr', '_colidx'),
Value('CUdeviceptr', '_csrdata'),
Value('CUdeviceptr', '_b'),
Value('CUdeviceptr', '_x'),
Value('int', 'nrows'),
Value('int', 'ncols'),
Value('int', 'nnz')
]),
Block([
Statement('typedef int IndexType'),
Statement('typedef %s ValueType' % M.ctype),
Statement(
'typedef typename cusp::array1d_view< thrust::device_ptr<IndexType> > indices'
),
Statement(
'typedef typename cusp::array1d_view< thrust::device_ptr<ValueType> > values'
),
Statement(
'typedef cusp::csr_matrix_view< indices, indices, values, IndexType, ValueType, cusp::device_memory > matrix'
),
Statement(
'thrust::device_ptr< IndexType > rowptr((IndexType *)_rowptr)'
),
Statement(
'thrust::device_ptr< IndexType > colidx((IndexType *)_colidx)'
),
Statement(
'thrust::device_ptr< ValueType > csrdata((ValueType *)_csrdata)'
),
Statement('thrust::device_ptr< ValueType > d_b((ValueType *)_b)'),
Statement('thrust::device_ptr< ValueType > d_x((ValueType *)_x)'),
Statement('indices row_offsets(rowptr, rowptr + nrows + 1)'),
Statement('indices column_indices(colidx, colidx + nnz)'),
Statement('values matrix_values(csrdata, csrdata + nnz)'),
Statement('values b(d_b, d_b + nrows)'),
Statement('values x(d_x, d_x + ncols)'),
Statement('thrust::fill(x.begin(), x.end(), (ValueType)0)'),
Statement(
'matrix A(nrows, ncols, nnz, row_offsets, column_indices, matrix_values)'
),
Statement('cusp::%s_monitor< ValueType > %s' %
('verbose' if parameters['ksp_monitor'] else 'default',
monitor)), precond_call, solve_call
]))
host_mod.add_to_preamble(
[Include('boost/python/extract.hpp'),
Include('string')])
host_mod.add_to_preamble([Statement('using namespace boost::python')])
host_mod.add_to_preamble([Statement('using namespace std')])
nvcc_mod.add_function(nvcc_function)
host_mod.add_function(
FunctionBody(
FunctionDeclaration(Value('void', 'solve'), [
Value('object', '_rowptr'),
Value('object', '_colidx'),
Value('object', '_csrdata'),
Value('object', '_b'),
Value('object', '_x'),
Value('object', '_nrows'),
Value('object', '_ncols'),
Value('object', '_nnz')
]),
Block([
Statement(
'CUdeviceptr rowptr = extract<CUdeviceptr>(_rowptr.attr("gpudata"))'
),
Statement(
'CUdeviceptr colidx = extract<CUdeviceptr>(_colidx.attr("gpudata"))'
),
Statement(
'CUdeviceptr csrdata = extract<CUdeviceptr>(_csrdata.attr("gpudata"))'
),
Statement(
'CUdeviceptr b = extract<CUdeviceptr>(_b.attr("gpudata"))'
),
Statement(
'CUdeviceptr x = extract<CUdeviceptr>(_x.attr("gpudata"))'
),
Statement('int nrows = extract<int>(_nrows)'),
Statement('int ncols = extract<int>(_ncols)'),
Statement('int nnz = extract<int>(_nnz)'),
Statement(
'__cusp_solve(rowptr, colidx, csrdata, b, x, nrows, ncols, nnz)'
)
])))
nvcc_toolchain.cflags.append('-arch')
nvcc_toolchain.cflags.append('sm_20')
nvcc_toolchain.cflags.append('-O3')
module = nvcc_mod.compile(gcc_toolchain,
nvcc_toolchain,
debug=configuration["debug"])
_cusp_cache[cache_key(M.ctype, parameters)] = module
return module
def generate_assignment_instruction_code(codegen_state, insn):
kernel = codegen_state.kernel
ecm = codegen_state.expression_to_code_mapper
from loopy.expression import VectorizabilityChecker
# {{{ vectorization handling
if codegen_state.vectorization_info:
if insn.atomicity:
raise Unvectorizable("atomic operation")
vinfo = codegen_state.vectorization_info
vcheck = VectorizabilityChecker(kernel, vinfo.iname, vinfo.length)
lhs_is_vector = vcheck(insn.assignee)
rhs_is_vector = vcheck(insn.expression)
if not lhs_is_vector and rhs_is_vector:
raise Unvectorizable("LHS is scalar, RHS is vector, cannot assign")
is_vector = lhs_is_vector
del lhs_is_vector
del rhs_is_vector
# }}}
from pymbolic.primitives import Variable, Subscript, Lookup
from loopy.symbolic import LinearSubscript
lhs = insn.assignee
if isinstance(lhs, Lookup):
lhs = lhs.aggregate
if isinstance(lhs, Variable):
assignee_var_name = lhs.name
assignee_indices = ()
elif isinstance(lhs, Subscript):
assignee_var_name = lhs.aggregate.name
assignee_indices = lhs.index_tuple
elif isinstance(lhs, LinearSubscript):
assignee_var_name = lhs.aggregate.name
assignee_indices = (lhs.index, )
else:
raise RuntimeError("invalid lvalue '%s'" % lhs)
del lhs
result = codegen_state.ast_builder.emit_assignment(codegen_state, insn)
# {{{ tracing
lhs_dtype = codegen_state.kernel.get_var_descriptor(
assignee_var_name).dtype
if kernel.options.trace_assignments or kernel.options.trace_assignment_values:
if codegen_state.vectorization_info and is_vector:
raise Unvectorizable("tracing does not support vectorization")
from pymbolic.mapper.stringifier import PREC_NONE
lhs_code = codegen_state.expression_to_code_mapper(
insn.assignee, PREC_NONE)
from cgen import Statement as S # noqa
gs, ls = kernel.get_grid_size_upper_bounds()
printf_format = "{}.{}[{}][{}]: {}".format(
kernel.name, insn.id, ", ".join("gid%d=%%d" % i
for i in range(len(gs))),
", ".join("lid%d=%%d" % i for i in range(len(ls))),
assignee_var_name)
printf_args = (["gid(%d)" % i for i in range(len(gs))] +
["lid(%d)" % i for i in range(len(ls))])
if assignee_indices:
printf_format += "[%s]" % ",".join(len(assignee_indices) * ["%d"])
printf_args.extend(
ecm(i, prec=PREC_NONE, type_context="i")
for i in assignee_indices)
if kernel.options.trace_assignment_values:
if lhs_dtype.numpy_dtype.kind == "i":
printf_format += " = %d"
printf_args.append(lhs_code)
elif lhs_dtype.numpy_dtype.kind == "f":
printf_format += " = %g"
printf_args.append(lhs_code)
elif lhs_dtype.numpy_dtype.kind == "c":
printf_format += " = %g + %gj"
printf_args.extend(["(%s).x" % lhs_code, "(%s).y" % lhs_code])
if printf_args:
printf_args_str = ", " + ", ".join(str(v) for v in printf_args)
else:
printf_args_str = ""
printf_insn = S('printf("{}\\n"{})'.format(printf_format,
printf_args_str))
from cgen import Block
if kernel.options.trace_assignment_values:
result = Block([result, printf_insn])
else:
# print first, execute later -> helps find segfaults
result = Block([printf_insn, result])
# }}}
return result
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
knl(queue, (local_size * macroblock_count, ), (local_size, ), c_buf, a_buf,
b_buf)
c = numpy.empty_like(a)
cl.enqueue_copy(queue, c, c_buf).wait()
