def get_semantic_tree(self, arg, output_name):
params = [
SymbolRef(self.array_name, POINTER(c_float)(), _global=True,
_const=True),
SymbolRef(arg.name, POINTER(c_float)(), _global=True, _const=True),
SymbolRef(output_name, POINTER(c_float)(), _global=True)
]
defn = []
defn.extend([
Assign(SymbolRef('element_id%d' % d, c_int()), get_global_id(d))
for d in range(len(arg.data.shape))
])
index = StringTemplate('element_id1 * $len_x + element_id0',
{'len_x': Constant(arg.data.shape[1])})
defn.append(
Assign(
ArrayRef(SymbolRef(params[-1].name), index),
self.original_tree(
ArrayRef(SymbolRef(params[0].name), index),
ArrayRef(SymbolRef(params[1].name), index),
)
)
)
entry_point = unique_kernel_name()
tree = FunctionDecl(None, entry_point, params, defn)
tree.set_kernel()
kernel = OclFile("kernel", [tree])
return kernel
def visit_PointsLoop(self, node):
self.loop_var = node.loop_var
body = list(map(self.visit, node.body))
params = [SymbolRef(x.name, x.type,
_global=self.arg_cfg_dict[x.name].is_global)
for x in self.params]
kernel_name = unique_kernel_name()
kernel = FunctionDecl(None, kernel_name, params, body)
kernel.set_kernel()
self.project.files.append(OclFile(kernel_name, [kernel]))
self.loop_var = None
def transform(self, tree, program_config):
arg_cfg = program_config[0]
self.entry_point = unique_kernel_name()
ctypeObject = c_float()
ctype = c_float
len_x = arg_cfg[0][1][0]
len_y = arg_cfg[0][1][1]
output = unique_name()
params = [
SymbolRef("input", POINTER(ctype)(), _global=True, _const=True),
SymbolRef(output, POINTER(ctype)(), _global=True)
]
defn = []
defn.extend([
Assign(SymbolRef('x', c_int()), get_global_id(0)),
Assign(SymbolRef('y', c_int()), get_global_id(1)),
Assign(SymbolRef('temp', ctypeObject), Constant(0)),
])
body = \
"""
temp = .5 * input[clamp(x/2, 0, (len_x / 2) - 1) * len_y +
clamp(y/2, 0, (len_y / 2) - 1)]
if (x & 0x1):
temp += .25 * input[clamp(x/2 + 1, 0, (len_x / 2) - 1) * len_y +
clamp(y/2, 0, (len_y / 2) - 1)]
else:
temp += .25 * input[clamp(x/2 - 1, 0, (len_x / 2) - 1) * len_y +
clamp(y/2, 0, (len_y / 2) - 1)]
if (y & 0x1):
temp += .25 * input[clamp(x/2, 0, (len_x / 2) - 1) * len_y +
clamp(y/2 + 1, 0, (len_y / 2) - 1)]
else:
temp += .25 * input[clamp(x/2, 0, (len_x / 2) - 1) *len_y +
clamp(y/2 - 1, 0, (len_y / 2) - 1)]
output[x * len_y + y] = temp
"""
body = ast.parse(body).body
name_dict = {
'output': output
}
const_dict = {
'len_x': len_x,
'len_y': len_y,
}
transformation = PyBasicConversions(name_dict, const_dict)
defn.extend(body)
tree = FunctionDecl(None, self.entry_point, params, defn)
tree.set_kernel()
kernel = OclFile("kernel", [tree])
kernel = transformation.visit(kernel)
return kernel
def visit_Lambda(self, node):
if isinstance(node, ast.Lambda):
def_node = ast.FunctionDef(name="default", args=node.args,
body=node.body, decorator_list=None)
params = [self.visit(p) for p in def_node.args.args]
defn = [Return(self.visit(def_node.body))]
decl_node = FunctionDecl(None, def_node.name, params, defn)
Lifter().visit_FunctionDecl(decl_node)
return decl_node
else:
return node
def transform(self, tree, program_config):
# TODO: Have to flip indices, figure out why
arg_cfg, tune_cfg = program_config
output_name = unique_name()
params = [
SymbolRef(self.array_name, POINTER(c_float)(), _global=True,
_const=True),
SymbolRef(arg_cfg[0][0], POINTER(c_float)(), _global=True,
_const=True),
SymbolRef(output_name, POINTER(c_float)(), _global=True)
]
defn = []
defn.extend([
Assign(SymbolRef('element_id%d' % d, c_int()), get_global_id(d))
for d in range(len(arg_cfg[0][2]))
])
index = StringTemplate('element_id1 * $len_x + element_id0',
{'len_x': Constant(arg_cfg[0][2][1])})
defn.append(
Assign(
ArrayRef(SymbolRef(params[-1].name), index),
tree(
ArrayRef(SymbolRef(params[0].name), index),
ArrayRef(SymbolRef(params[1].name), index),
)
)
)
entry_point = unique_kernel_name()
tree = FunctionDecl(None, entry_point, params, defn)
tree.set_kernel()
fn = ArrayOpConcrete(self.array, self.generate_output(output_name))
kernel = OclFile("kernel", [tree])
program = clCreateProgramWithSource(
fn.context, kernel.codegen()
).build()
ptr = program[entry_point]
return fn.finalize(ptr, (arg_cfg[0][2][1], arg_cfg[0][2][0]))
def transform(self, tree, program_config):
# TODO: Have to flip indices, figure out why
arg_cfg = program_config[0]
input_name = arg_cfg[0][0]
self.output_name = unique_name()
params = [
SymbolRef(
input_name, POINTER(c_float)(), _global=True, _const=True),
SymbolRef(self.output_name, POINTER(c_float)(), _global=True)
]
defn = []
defn.extend([
Assign(SymbolRef('element_id%d' % d, c_int()), get_global_id(d))
for d in range(len(arg_cfg[0][2]))
])
out_index = StringTemplate('element_id1 * $len_x + element_id0',
{'len_x': Constant(arg_cfg[0][2][1])})
defn.append(Assign(
ArrayRef(SymbolRef(self.output_name), out_index),
Div(
Add(
ArrayRef(
SymbolRef(input_name),
StringTemplate(
'(element_id1 * 2) * $len_x + (element_id0 * 2)',
{'len_x': Constant(arg_cfg[0][2][1])})
),
Add(
ArrayRef(
SymbolRef(input_name),
StringTemplate(
'(element_id1 * 2) * $len_x + \
(element_id0 * 2 + 1)',
{'len_x': Constant(arg_cfg[0][2][1])})
),
Add(
ArrayRef(
SymbolRef(input_name),
StringTemplate(
'(element_id1 * 2 + 1) * $len_x + \
(element_id0 * 2 + 1)',
{'len_x': Constant(arg_cfg[0][2][1])})
),
Add(
ArrayRef(
SymbolRef(input_name),
StringTemplate(
'(element_id1 * 2 + 1) * $len_x + \
(element_id0 * 2)',
{'len_x': Constant(arg_cfg[0][2][1])})
),
)
)
)
),
Constant(4.0)
)
))
self.entry_point = unique_kernel_name()
tree = FunctionDecl(None, self.entry_point, params, defn)
tree.set_kernel()
kernel = OclFile("kernel", [tree])
return kernel
def test_none(self):
tree = FunctionDecl(None, "foo", ())
self._check_code(tree, "void foo()")
def visit_FunctionDef(self, node):
if ast.get_docstring(node):
node.body.pop(0)
params = [self.visit(p) for p in node.args.args]
defn = [self.visit(s) for s in node.body]
return FunctionDecl(None, node.name, params, defn)
def transform(self, tree, program_config):
call_args = program_config[0]
base_size = call_args.base_shape[0] * call_args.base_shape[1]
border = call_args.border
c_float_type = c_float
c_int_type = c_int
transformer = PyBasicConversions()
output = unique_name()
init_entry_point = unique_kernel_name()
init_params = [
SymbolRef('input', POINTER(c_float_type)(), _global=True, _const=True),
SymbolRef(output, POINTER(c_float_type)(), _global=True),
]
init_defn = []
init_defn.extend([
Assign(SymbolRef('x', c_int()), get_global_id(0)),
Assign(SymbolRef('y', c_int()), get_global_id(1)),
])
body = """{output}[y * {len_x} + x] = input[y * {len_x} + x]""".format(
output=output, len_x=call_args.base_shape[0]
)
print(body)
tree_body = ast.parse(body).body
init_defn.extend(tree_body)
init_tree = FunctionDecl(None, init_entry_point, init_params, init_defn)
init_tree.set_kernel()
init_kernel = OclFile('kernel', [init_tree])
init_kernel = transformer.visit(init_kernel)
print("init kernel codegen")
print(init_kernel.codegen())
compute_entry_point = unique_kernel_name()
compute_params = [
SymbolRef(output, POINTER(c_float_type)(), _global=True),
SymbolRef('power', c_int(), _const=True),
]
compute_defn = []
compute_defn.extend([
Assign(SymbolRef('x', c_int()), get_global_id(0)),
Assign(SymbolRef('y', c_int()), get_global_id(1)),
])
body = """{matrix}[(power+1) * {base_size} + y * {len_x} + x] =
0.1 * {matrix}[
power * {base_size} + clamp(y-1, {border}, {len_y}-{border}-1) * {len_x} + clamp(x, {border}, {len_x}-{border}-1)
] +
0.1 * {matrix}[
power * {base_size} + clamp(y+1, {border}, {len_y}-{border}-1) * {len_x} + clamp(x, {border}, {len_x}-{border}-1)
] +
0.4 * {matrix}[
power * {base_size} + clamp(y, {border}, {len_y}-{border}-1) * {len_x} + clamp(x-1, {border}, {len_x}-{border}-1)
] +
0.4 * {matrix}[
power * {base_size} + clamp(y, {border}, {len_y}-{border}-1) * {len_x} + clamp(x+1, {border}, {len_x}-{border}-1)
] +
1.0 * {matrix}[
power * {base_size} + clamp(y, {border}, {len_y}-{border}-1) * {len_x} + clamp(x, {border}, {len_x}-{border}-1)
]
""".format(
matrix=output,
base_size=base_size,
len_y=call_args.base_shape[0],
len_x=call_args.base_shape[1],
border=border,
)
body = re.sub("""\s\s*""", " ", body)
print(body)
tree_body = ast.parse(body).body
compute_defn.extend(tree_body)
compute_tree = FunctionDecl(None, compute_entry_point, compute_params, compute_defn)
compute_tree.set_kernel()
compute_kernel = OclFile('kernel', [compute_tree])
compute_kernel = transformer.visit(compute_kernel)
print("compute kernel codegen")
print(compute_kernel.codegen())
fn = OclMatrixPowers()
init_program = clCreateProgramWithSource(fn.context, init_kernel.codegen()).build()
init_ptr = init_program[init_entry_point]
compute_program = clCreateProgramWithSource(fn.context, compute_kernel.codegen()).build()
compute_ptr = compute_program[compute_entry_point]
return fn.finalize(init_ptr, compute_ptr, (call_args.base_shape[1], call_args.base_shape[0]))
def transform(self, tree, program_config):
arg_cfg = program_config[0]
len_x = Constant(arg_cfg[0][1][0])
len_y = Constant(arg_cfg[0][1][1])
self.entry_point = unique_kernel_name()
body = FunctionDecl(
None,
self.entry_point,
[
SymbolRef('input', POINTER(c_float)(), _global=True,
_const=True),
SymbolRef('u', POINTER(c_float)(), _global=True, _const=True),
SymbolRef('v', POINTER(c_float)(), _global=True, _const=True),
SymbolRef('output', POINTER(c_float)(), _global=True)
],
[
Assign(SymbolRef('x', c_int()), get_global_id(0)),
Assign(SymbolRef('y', c_int()), get_global_id(1)),
Assign(
SymbolRef('my_x', c_int()),
Cast(c_int(),
ArrayRef(SymbolRef('u'),
Add(SymbolRef('x'),
Mul(SymbolRef('y'), len_x))))),
Assign(
SymbolRef('my_y', c_int()),
Cast(c_int(),
ArrayRef(SymbolRef('v'),
Add(SymbolRef('x'),
Mul(SymbolRef('y'), len_x))))),
Assign(
SymbolRef('xfrac', c_float()),
Sub(ArrayRef(SymbolRef('u'),
Add(SymbolRef('x'),
Mul(len_x, SymbolRef('y')))),
SymbolRef('my_x'))),
Assign(
SymbolRef('yfrac', c_float()),
Sub(ArrayRef(SymbolRef('v'),
Add(SymbolRef('x'),
Mul(len_x, SymbolRef('y')))),
SymbolRef('my_y'))),
If(Lt(
ArrayRef(SymbolRef('u'),
Add(SymbolRef('x'),
Mul(len_x, SymbolRef('y')))),
Constant(0.0)),
[
PostDec(SymbolRef('my_x')),
Assign(SymbolRef('xfrac'),
Add(Constant(1.0),
SymbolRef('xfrac')))
]),
If(Lt(
ArrayRef(SymbolRef('v'),
Add(SymbolRef('x'),
Mul(len_x, SymbolRef('y')))),
Constant(0.0)),
[
PostDec(SymbolRef('my_y')),
Assign(SymbolRef('yfrac'),
Add(Constant(1.0),
SymbolRef('yfrac')))
]),
Assign(SymbolRef('tmp', c_float()), Constant(0.0)),
If(
And(
And(
GtE(Add(SymbolRef('x'),
SymbolRef('my_x')), Constant(0)),
Lt(Add(SymbolRef('x'),
Add(SymbolRef('my_x'), Constant(1))), len_x)
),
And(
GtE(Add(SymbolRef('y'),
SymbolRef('my_y')), Constant(0)),
Lt(Add(SymbolRef('y'),
Add(SymbolRef('my_y'), Constant(1))), len_y)
)
),
[
AddAssign(
SymbolRef('tmp'),
Mul(
Mul(
ArrayRef(
SymbolRef('input'),
Add(
Add(
SymbolRef('x'),
SymbolRef('my_x')),
Mul(
len_x,
Add(SymbolRef('y'),
SymbolRef('my_y'))))),
Sub(Constant(1.0), SymbolRef('xfrac'))),
Sub(Constant(1.0), SymbolRef('yfrac')))),
AddAssign(
SymbolRef('tmp'),
Mul(
Mul(
ArrayRef(
SymbolRef('input'),
Add(
Add(
Add(SymbolRef('x'),
SymbolRef('my_x')),
Constant(1)),
Mul(
len_x,
Add(SymbolRef('y'),
SymbolRef('my_y'))))),
SymbolRef('xfrac')),
Sub(Constant(1.0), SymbolRef('yfrac')))),
AddAssign(
SymbolRef('tmp'),
Mul(
Mul(
ArrayRef(
SymbolRef('input'),
Add(
Add(
SymbolRef('x'),
SymbolRef('my_x')),
Mul(
len_x,
Add(Add(SymbolRef('y'),
SymbolRef('my_y')),
Constant(1))))),
Sub(Constant(1.0), SymbolRef('xfrac'))),
SymbolRef('yfrac'))),
AddAssign(
SymbolRef('tmp'),
Mul(
Mul(ArrayRef(
SymbolRef('input'),
Add(
Add(
Add(SymbolRef('x'),
SymbolRef('my_x')),
Constant(1)),
Mul(
len_x,
Add(Add(SymbolRef('y'),
SymbolRef('my_y')),
Constant(1))))),
SymbolRef('xfrac')),
SymbolRef('yfrac'))),
],
Assign(
SymbolRef('tmp'),
ArrayRef(
SymbolRef('input'),
Add(
FunctionCall(
SymbolRef('clamp'),
[
Add(SymbolRef('x'), SymbolRef('my_x')),
Constant(0),
Sub(len_x, Constant(1))
]
),
Mul(
len_x,
FunctionCall(SymbolRef('clamp'), [
Add(SymbolRef('y'), SymbolRef('my_y')),
Constant(0),
Sub(len_y, Constant(1))
]
),
)
)
)
)
),
Assign(
ArrayRef(SymbolRef('output'),
Add(SymbolRef('x'),
Mul(len_x, SymbolRef('y')))),
SymbolRef('tmp')
)
]
)
body.set_kernel()
kernel = OclFile("kernel", [body])
return kernel