def field_extraction_code(field,
is_temporary,
declaration_only=False,
no_declaration=False,
is_gpu=False):
"""Returns code string for getting a field pointer.
This can happen in two ways: either the field is extracted from a walberla block, or a temporary field to swap is
created.
Args:
field: the field for which the code should be created
is_temporary: new_filtered field from block (False) or create a temporary copy of an existing field (True)
declaration_only: only create declaration instead of the full code
no_declaration: create the extraction code, and assume that declarations are elsewhere
is_gpu: if the field is a GhostLayerField or a GpuField
"""
# Determine size of f coordinate which is a template parameter
f_size = get_field_fsize(field)
field_name = field.name
dtype = get_base_type(field.dtype)
field_type = make_field_type(dtype, f_size, is_gpu)
if not is_temporary:
dtype = get_base_type(field.dtype)
field_type = make_field_type(dtype, f_size, is_gpu)
if declaration_only:
return "%s * %s;" % (field_type, field_name)
else:
prefix = "" if no_declaration else "auto "
return "%s%s = block->uncheckedFastGetData< %s >(%sID);" % (
prefix, field_name, field_type, field_name)
else:
assert field_name.endswith('_tmp')
original_field_name = field_name[:-len('_tmp')]
if declaration_only:
return "%s * %s;" % (field_type, field_name)
else:
declaration = "{type} * {tmp_field_name};".format(
type=field_type, tmp_field_name=field_name)
tmp_field_str = temporary_fieldTemplate.format(
original_field_name=original_field_name,
tmp_field_name=field_name,
type=field_type)
return tmp_field_str if no_declaration else declaration + tmp_field_str
def __new_stage2__(cls, field_name, field_dtype, const):
name = f"_data_{field_name}"
dtype = PointerType(get_base_type(field_dtype),
const=const,
restrict=True)
obj = super(FieldPointerSymbol, cls).__xnew__(cls, name, dtype)
obj.field_name = field_name
return obj
def check_type(e):
if only_type is None:
return True
try:
base_type = get_base_type(get_type_of_expression(e))
except ValueError:
return False
if only_type == 'int' and (base_type.is_int() or base_type.is_uint()):
return True
if only_type == 'real' and (base_type.is_float()):
return True
else:
return base_type == only_type
def generate_benchmark(ast, likwid=False, openmp=False, timing=False):
"""Return C code of a benchmark program for the given kernel.
Args:
ast: the pystencils AST object as returned by create_kernel
likwid: if True likwid markers are added to the code
openmp: relevant only if likwid=True, to generated correct likwid initialization code
timing: add timing output to the code, prints time per iteration to stdout
Returns:
C code as string
"""
accessed_fields = {f.name: f for f in ast.fields_accessed}
constants = []
fields = []
call_parameters = []
for p in ast.get_parameters():
if not p.is_field_parameter:
constants.append((p.symbol.name, str(p.symbol.dtype)))
call_parameters.append(p.symbol.name)
else:
assert p.is_field_pointer, "Benchmark implemented only for kernels with fixed loop size"
field = accessed_fields[p.field_name]
dtype = str(get_base_type(p.symbol.dtype))
fields.append((p.field_name, dtype, prod(field.shape)))
call_parameters.append(p.field_name)
header_list = get_headers(ast)
includes = "\n".join(["#include %s" % (include_file,) for include_file in header_list])
# Strip "#pragma omp parallel" from within kernel, because main function takes care of that
# when likwid and openmp are enabled
if likwid and openmp:
if len(ast.body.args) > 0 and isinstance(ast.body.args[0], PragmaBlock):
ast.body.args[0].pragma_line = ''
args = {
'likwid': likwid,
'openmp': openmp,
'kernel_code': generate_c(ast, dialect='c'),
'kernelName': ast.function_name,
'fields': fields,
'constants': constants,
'call_argument_list': ",".join(call_parameters),
'includes': includes,
'timing': timing,
}
return benchmark_template.render(**args)
def get_field_stride(param):
field = param.fields[0]
type_str = get_base_type(param.symbol.dtype).base_name
stride_names = ['xStride()', 'yStride()', 'zStride()', 'fStride()']
stride_names = [
"%s(%s->%s)" % (type_str, param.field_name, e) for e in stride_names
]
strides = stride_names[:field.spatial_dimensions]
if field.index_dimensions > 0:
additional_strides = [1]
for shape in reversed(field.index_shape[1:]):
additional_strides.append(additional_strides[-1] * shape)
assert len(additional_strides) == field.index_dimensions
f_stride_name = stride_names[-1]
strides.extend([
"%s(%d * %s)" % (type_str, e, f_stride_name)
for e in reversed(additional_strides)
])
return strides[param.symbol.coordinate]
def generate_members(ctx,
kernel_info,
parameters_to_ignore=(),
only_fields=False):
ast = kernel_info.ast
fields = {f.name: f for f in ast.fields_accessed}
params_to_skip = tuple(parameters_to_ignore) + tuple(
kernel_info.temporary_fields)
params_to_skip += tuple(e[1] for e in kernel_info.varying_parameters)
is_gpu = ctx['target'] == 'gpu'
result = []
for param in kernel_info.parameters:
if only_fields and not param.is_field_parameter:
continue
if param.is_field_pointer and param.field_name not in params_to_skip:
result.append("BlockDataID %sID;" % (param.field_name, ))
elif not param.is_field_parameter and param.symbol.name not in params_to_skip:
result.append("%s %s;" % (
param.symbol.dtype,
param.symbol.name,
))
for field_name in kernel_info.temporary_fields:
f = fields[field_name]
if field_name in parameters_to_ignore:
continue
assert field_name.endswith('_tmp')
original_field_name = field_name[:-len('_tmp')]
f_size = get_field_fsize(f)
field_type = make_field_type(get_base_type(f.dtype), f_size, is_gpu)
result.append(
temporary_fieldMemberTemplate.format(
type=field_type, original_field_name=original_field_name))
if hasattr(kernel_info, 'varying_parameters'):
result.extend(["%s %s;" % e for e in kernel_info.varying_parameters])
return "\n".join(result)