Skip to content
Snippets Groups Projects
Commit 12a86893 authored by Jan Hönig's avatar Jan Hönig
Browse files

Progress

parent 49ebf273
No related branches found
No related tags found
No related merge requests found
Hide whitespace changes
Inline Side-by-side
pystencils_benchmark/benchmark.py
+ 81
− 15
View file @ 12a86893
from pathlib import Path
from typing import Union, List
from jinja2 import Environment, PackageLoader, StrictUndefined
from pystencils.backends.cbackend import generate_c, get_headers
from pystencils.astnodes import KernelFunction
from pystencils.astnodes import KernelFunction, PragmaBlock
from pystencils.enums import Backend
from pystencils.data_types import get_base_type
from pystencils.sympyextensions import prod
from pystencils.integer_functions import modulo_ceil
import numpy as np
def generate_benchmark(kernel_asts: Union[List[KernelFunction], KernelFunction],
def generate_benchmark(kernel_ast: KernelFunction,
path: Path = None,
*,
dialect: Backend = Backend.C) -> None:
......@@ -21,19 +24,82 @@ def generate_benchmark(kernel_asts: Union[List[KernelFunction], KernelFunction],
include_path = path / 'include'
include_path.mkdir(parents=True, exist_ok=True)
if isinstance(kernel_asts, KernelFunction):
kernel_asts = [kernel_asts]
for kernel_ast in kernel_asts:
kernel_name = kernel_ast.function_name
kernel_name = kernel_ast.function_name
header = kernel_header(kernel_ast, dialect)
with open(include_path / f'{kernel_name}.h', 'w+') as f:
f.write(header)
header = kernel_header(kernel_ast, dialect)
with open(include_path / f'{kernel_name}.h', 'w+') as f:
f.write(header)
source = kernel_source(kernel_ast, dialect)
with open(src_path / f'{kernel_name}.c', 'w+') as f:
f.write(source)
with open(src_path / 'main.c', 'w+') as f:
f.write(kernel_main(kernel_ast))
def kernel_main(kernel: KernelFunction, timing: bool = False):
"""
Return C code of a benchmark program for the given kernel.
Args:
kernel: the pystencils AST object as returned by create_kernel
timing: add timing output to the code, prints time per iteration to stdout
Returns:
C code as string
"""
kernel_name = kernel.function_name
accessed_fields = {f.name: f for f in kernel.fields_accessed}
constants = []
fields = []
call_parameters = []
for p in kernel.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))
np_dtype = get_base_type(p.symbol.dtype).numpy_dtype
size_data_type = np_dtype.itemsize
dim0_size = field.shape[-1]
dim1_size = np.prod(field.shape[:-1])
elements = prod(field.shape)
if kernel.instruction_set:
align = kernel.instruction_set['width'] * size_data_type
padding_elements = modulo_ceil(dim0_size, kernel.instruction_set['width']) - dim0_size
padding_bytes = padding_elements * size_data_type
ghost_layers = max(max(kernel.ghost_layers))
size = dim1_size * padding_bytes + np.prod(field.shape) * size_data_type
assert align % np_dtype.itemsize == 0
offset = ((dim0_size + padding_elements + ghost_layers) % kernel.instruction_set['width']) * size_data_type
fields.append((p.field_name, dtype, elements, size, offset, align))
call_parameters.append(p.field_name)
else:
size = elements * size_data_type
fields.append((p.field_name, dtype, elements, size, 0, 0))
call_parameters.append(p.field_name)
includes = f'#include "{kernel_name}"\n'
jinja_context = {
'kernel_code': generate_c(kernel, dialect=Backend.C),
'kernel_name': kernel_name,
'fields': fields,
'constants': constants,
'call_argument_list': ",".join(call_parameters),
'includes': includes,
'timing': timing,
}
source = kernel_source(kernel_ast, dialect)
with open(src_path / f'{kernel_name}.cpp', 'w+') as f:
f.write(source)
env = Environment(loader=PackageLoader('pystencils_benchmark'), undefined=StrictUndefined)
main = env.get_template('main.c').render(**jinja_context)
return main
def kernel_header(kernel_ast: KernelFunction, dialect: Backend = Backend.C) -> str:
......
pystencils_benchmark/templates/benchmark.c → pystencils_benchmark/templates/main.c
+ 6
− 38
View file @ 12a86893
......@@ -8,10 +8,6 @@
{{ includes }}
{%- if likwid %}
#include <likwid.h>
{%- endif %}
#define RESTRICT __restrict__
#define FUNC_PREFIX
void dummy(void *);
......@@ -56,27 +52,22 @@ void aligned_free( void *ptr )
}
{{kernel_code}}
int main(int argc, char **argv)
{
{%- if likwid %}
likwid_markerInit();
{%- endif %}
{%- for field_name, dataType, elements, size, offset, alignment in fields %}
// Initialization {{field_name}}
{%- if alignment > 0 %}
{%- if alignment > 0 %}
{{dataType}} * {{field_name}} = ({{dataType}} *) aligned_malloc_with_offset({{size}}, {{alignment}}, {{offset}});
{%- else %}
{%- else %}
{{dataType}} * {{field_name}} = new {{dataType}}[{{elements}}];
{%- endif %}
{%- endif %}
for (unsigned long long i = 0; i < {{elements}}; ++i)
{{field_name}}[i] = 0.23;
if(var_false)
{
dummy({{field_name}});
}
{%- endfor %}
......@@ -92,22 +83,10 @@ int main(int argc, char **argv)
{%- endfor %}
{%- if likwid and openmp %}
#pragma omp parallel
{
likwid_markerRegisterRegion("loop");
#pragma omp barrier
{%- elif likwid %}
likwid_markerRegisterRegion("loop");
{%- endif %}
for(int warmup = 1; warmup >= 0; --warmup) {
int repeat = 2;
if(warmup == 0) {
repeat = atoi(argv[1]);
{%- if likwid %}
likwid_markerStartRegion("loop");
{%- endif %}
}
{%- if timing %}
......@@ -117,7 +96,7 @@ int main(int argc, char **argv)
for (; repeat > 0; --repeat)
{
{{kernelName}}({{call_argument_list}});
{{kernel_name}}({{call_argument_list}});
// Dummy calls
{%- for field_name, dataType, elements, size, offset, alignment in fields %}
......@@ -135,17 +114,6 @@ int main(int argc, char **argv)
}
{%- if likwid %}
likwid_markerStopRegion("loop");
{%- if openmp %}
}
{%- endif %}
{%- endif %}
{%- if likwid %}
likwid_markerClose();
{%- endif %}
{%- for field_name, dataType, elements, size, offset, alignment in fields %}
{%- if alignment > 0 %}
aligned_free({{field_name}});
......
test.ipynb 0 → 100644
+ 144
− 0
View file @ 12a86893
%% Cell type:code id: tags:
``` python
import numpy as np
import pystencils as ps
from pystencils_benchmark import kernel_header, kernel_source, generate_benchmark
from pathlib import Path
```
%% Cell type:code id: tags:
``` python
config = ps.CreateKernelConfig()
```
%% Cell type:code id: tags:
``` python
a, b, c = ps.fields(a=np.ones(4000000), b=np.ones(4000000), c=np.ones(4000000))
```
%% Cell type:code id: tags:
``` python
@ps.kernel
def vadd():
a[0] @= b[0] + c[0]
```
%% Cell type:code id: tags:
``` python
kernel_vadd = ps.create_kernel(vadd, config=config)
ps.show_code(kernel_vadd)
```
%% Output
%% Cell type:code id: tags:
``` python
example_path = Path.cwd() / 'example'
generate_benchmark(kernel_vadd, example_path)
```
%% Output
---------------------------------------------------------------------------
AttributeError Traceback (most recent call last)
/tmp/ipykernel_28091/2748650420.py in <module>
1 example_path = Path.cwd() / 'example'
----> 2 generate_benchmark(kernel_vadd, example_path)
~/git/pystencils-benchmark/pystencils_benchmark/benchmark.py in generate_benchmark(kernel_ast, path, dialect)
31 f.write(header)
32
---> 33 source = kernel_source(kernel_ast, dialect)
34 with open(src_path / f'{kernel_name}.c', 'w+') as f:
35 f.write(source)
~/git/pystencils-benchmark/pystencils_benchmark/benchmark.py in kernel_source(kernel_ast, dialect)
119 function_source = generate_c(kernel_ast, dialect=dialect)
120 headers = get_headers(kernel_ast)
--> 121 headers.add(f'"{kernel_name}.h"')
122
123 jinja_context = {
AttributeError: 'list' object has no attribute 'add'
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment