def clSetKernelArg(kernel, arg_index, arg_size, arg_value):
if isinstance(kernel, str): kernel = SymbolRef(kernel)
if isinstance(arg_index, int): arg_index = Constant(arg_index)
if isinstance(arg_size, int): arg_size = Constant(arg_size)
if isinstance(arg_value, str): arg_value = Ref(SymbolRef(arg_value))
return FunctionCall(SymbolRef("clSetKernelArg"),
[kernel, arg_index, arg_size, arg_value])
def test_array_ref(self):
tree = MultiNode([
SymbolRef("foo", ctypes.POINTER(ctypes.c_double)()),
Assign(SymbolRef("____temp__x"), ArrayRef(SymbolRef("foo"), Constant(0)))
])
DeclarationFiller().visit(tree)
self._check_code(tree, "\ndouble* foo;\n"
"double ____temp__x = foo[0];\n")
def test_complex(self):
node = ArrayDef(
SymbolRef('myArray', ct.c_int()), Constant(2),
Array(body=[
Add(SymbolRef('b'), SymbolRef('c')),
Mul(Sub(Constant(99), SymbolRef('d')), Constant(200))
]))
self._check_code(node, "int myArray[2] = {b + c, (99 - d) * 200}")
def test_dot(self):
op = SymbolRef("op")
setattr(op, "get_type", lambda: ctypes.c_char())
foo = SymbolRef("foo")
setattr(foo, "get_type", lambda: ctypes.c_double())
tree = Assign(SymbolRef("x"), Dot(foo, op))
DeclarationFiller().visit(tree)
self._check_code(tree, "char x = foo . op")
def clEnqueueReadBuffer(queue, buf, blocking, offset, cb, ptr, num_events=0, evt_list_ptr=None, evt=None):
if isinstance(buf, str): buf = SymbolRef(buf)
if isinstance(blocking, bool): blocking = Constant(int(blocking))
if isinstance(ptr, str): ptr = SymbolRef(ptr)
if not isinstance(offset, ast.AST): offset = Constant(offset)
if not isinstance(cb, ast.AST): cb = Constant(cb)
if not isinstance(num_events, ast.AST): num_events = Constant(num_events)
if not isinstance(evt_list_ptr, ast.AST): event_list_ptr = NULL()
if not isinstance(evt, ast.AST): evt = NULL()
return FunctionCall(SymbolRef('clEnqueueReadBuffer'), [
queue, buf, blocking, offset, cb, ptr, num_events, event_list_ptr, evt])
def clEnqueueCopyBuffer(queue, src_buf, dst_buf, src_offset=0, dst_offset=0, cb=0):
if isinstance(src_buf, str): src_buf = SymbolRef(src_buf)
if isinstance(dst_buf, str): dst_buf = SymbolRef(dst_buf)
if isinstance(src_offset, int): src_offset = Constant(src_offset)
if isinstance(dst_offset, int): dst_offset = Constant(dst_offset)
if isinstance(cb, int): cb = Constant(cb)
num_events = Constant(0)
event_list_ptr = NULL()
evt = NULL()
return FunctionCall(SymbolRef('clEnqueueCopyBuffer'), [
queue, src_buf, dst_buf, src_offset, dst_offset, cb, num_events, event_list_ptr, evt])
def printf(fmt, *args):
"""
Makes a printf call. Args must be CtreeNodes.
"""
for arg in args:
assert isinstance(arg, CtreeNode)
return FunctionCall(SymbolRef("printf"), [String(fmt)] + list(args))
def test_long(self):
tree = SymbolRef("i", ctypes.c_long())
if sys.maxsize > 2 ** 32:
self._check_code(tree, "long i")
else:
# int == long
self._check_code(tree, "int i")
def test_simple_array_def(self):
self._check_code(
ArrayDef(
SymbolRef('hi', ct.c_int()),
Constant(2),
Array(body=[Constant(0), Constant(1)]),
), "int hi[2] = {0, 1}")
def visit_For(self, node):
"""restricted, for now, to range as iterator with long-type args"""
if isinstance(node, ast.For) and \
isinstance(node.iter, ast.Call) and \
isinstance(node.iter.func, ast.Name) and \
node.iter.func.id == 'range':
Range = node.iter
nArgs = len(Range.args)
if nArgs == 1:
stop = self.visit(Range.args[0])
start, step = Constant(0), Constant(1)
elif nArgs == 2:
start, stop = map(self.visit, Range.args)
step = Constant(1)
elif nArgs == 3:
start, stop, step = map(self.visit, Range.args)
else:
raise Exception("Cannot convert a for...range with %d args." % nArgs)
# TODO allow any expressions castable to Long type
assert isinstance(stop.get_type(), c_long), "Can only convert range's with stop values of Long type."
assert isinstance(start.get_type(), c_long), "Can only convert range's with start values of Long type."
assert isinstance(step.get_type(), c_long), "Can only convert range's with step values of Long type."
target = SymbolRef(node.target.id, c_long())
for_loop = For(
Assign(target, start),
Lt(target.copy(), stop),
AddAssign(target.copy(), step),
[self.visit(stmt) for stmt in node.body],
)
return for_loop
node.body = list(map(self.visit, node.body))
return node
def clEnqueueNDRangeKernel(queue, kernel, work_dim=1, work_offset=0, global_size=0, local_size=0):
assert isinstance(queue, SymbolRef)
assert isinstance(kernel, SymbolRef)
global_size_sym = SymbolRef('global_size', c_size_t())
local_size_sym = SymbolRef('local_size', c_size_t())
call = FunctionCall(SymbolRef("clEnqueueNDRangeKernel"), [
queue, kernel,
work_dim, work_offset,
Ref(global_size_sym.copy()), Ref(local_size_sym.copy()),
0, NULL(), NULL()
])
return Block([
Assign(global_size_sym, Constant(global_size)),
Assign(local_size_sym, Constant(local_size)),
call
])
def test_template_parent_pointers(self):
from ctree.c.nodes import SymbolRef
symbol = SymbolRef("hello")
template = "char *str = $val"
template_args = {
'val': symbol,
}
node = StringTemplate(template, template_args)
self.assertIs(symbol.parent, node)
def test_template_with_transformer(self):
from ctree.visitors import NodeTransformer
from ctree.c.nodes import String, SymbolRef
template = "char *str = $val"
template_args = {
'val': SymbolRef("hello"),
}
tree = StringTemplate(template, template_args)
self._check(tree, 'char *str = hello')
class SymbolsToStrings(NodeTransformer):
def visit_SymbolRef(self, node):
return String(node.name)
tree = SymbolsToStrings().visit(tree)
self._check(tree, 'char *str = "hello"')
def test_template_parent_pointers_with_transformer(self):
from ctree.visitors import NodeTransformer
from ctree.c.nodes import String, SymbolRef
template = "char *str = $val"
template_args = {
'val': SymbolRef("hello"),
}
class SymbolsToStrings(NodeTransformer):
def visit_SymbolRef(self, node):
return String(node.name)
tree = StringTemplate(template, template_args)
tree = SymbolsToStrings().visit(tree)
template_node, string = tree, tree.val
self.assertIs(string.parent, template_node)
def CL_SUCCESS():
return SymbolRef("CL_SUCCESS")
def CL_DEVICE_TYPE_ALL():
return SymbolRef("CL_DEVICE_TYPE_ALL")
def CL_DEVICE_TYPE_DEFAULT():
return SymbolRef("CL_DEVICE_TYPE_DEFAULT")
def CL_DEVICE_TYPE_ACCELERATOR():
return SymbolRef("CL_DEVICE_TYPE_ACCELERATOR")
def CL_DEVICE_TYPE_CPU():
return SymbolRef("CL_DEVICE_TYPE_CPU")
def test_bad_type(self):
class Bad(object): pass
with self.assertRaises(ValueError):
SymbolRef("i", Bad()).codegen()
def test_pointer(self):
tree = SymbolRef("i", ctypes.POINTER(ctypes.c_double)())
self._check_code(tree, "double* i")
def get_num_groups(id):
return FunctionCall(SymbolRef('get_num_groups'), [Constant(id)])
def get_global_id(id):
return FunctionCall(SymbolRef('get_global_id'), [Constant(id)])
def CLK_LOCAL_MEM_FENCE():
return SymbolRef("CLK_LOCAL_MEM_FENCE")
def barrier(arg):
return FunctionCall(SymbolRef('barrier'), [arg])
def test_none(self):
tree = SymbolRef("i", ctypes.c_void_p())
self._check_code(tree, "void* i")
def get_local_size(id):
return FunctionCall(SymbolRef('get_local_size'), [Constant(id)])
def test_bool(self):
tree = SymbolRef("i", ctypes.c_bool())
self._check_code(tree, "bool i")
def clReleaseMemObject(arg):
return FunctionCall(SymbolRef('clReleaseMemObject'), [arg])
def test_string(self):
tree = SymbolRef("i", ctypes.c_char_p())
self._check_code(tree, "char* i")
