def convert_type(loc, type_enum, val):
'''
type_enum: Type.INT, Type.FLOAT, ...
val: string that will be converted to type 'type_enum'
Return 'val' as a value of the type corresponding to 'type_enum'. For
example:
convert_type(loc, Type.INT, "67") -> 67
convert_type(loc, Type.FLOAT, "0.43") -> 0.43
convert_type(loc, Type.STRING, "hello") -> "hello"
'''
try:
if type_enum == Type.INT:
return int(val)
elif type_enum == Type.FLOAT:
return float(val)
elif type_enum == Type.STRING:
if val[0] != '\'' or val[-1] != '\'':
raise error.Syntax(loc, f'invalid string: {val}')
return str(val[1:-1])
else:
error_str = f'convert_type: unknown type {type_enum}'
raise error.InternalError(loc, error_str)
except ValueError:
type_str = Type.type_to_str(loc, type_enum)
raise error.IncompatibleType(loc, type_str, val)
def __validate_list_at_expr_type(self, expr):
''' Typecheck and add environment data for a single ListAt Expr. '''
if expr.type != Type.UNDETERMINED:
error_str = f'expected ListAt to have UNDETERMINED type'
raise error.InternalError(expr.loc, error_str)
# Get the type of the list.
list_type = self._env.lookup_variable(expr.loc, expr.name)
# Determine the type of the list element.
elem_type = None
if list_type == Type.LIST_INT:
elem_type = Type.INT
elif list_type == Type.LIST_FLOAT:
elem_type = Type.FLOAT
elif list_type == Type.LIST_STRING:
elem_type = Type.STRING
else:
error_str = f'expected list type but got {list_type}'
raise error.Type(expr.loc, error_str)
# Make sure the index has type int.
self.__validate_single_expr(expr.index)
self.__validate_type_of_expr(expr.index.loc, Type.INT, expr.index)
# ListAt returns a list element, so it has the list element type.
expr.type = elem_type
def __validate_list_set_expr_type(self, expr):
''' Typecheck and add environment data for a single ListSet Expr. '''
if expr.type != Type.NONE:
error_str = f'expected ListSet to have NONE type'
raise error.InternalError(expr.loc, error_str)
# Get the type of the list.
list_type = self._env.lookup_variable(expr.loc, expr.name)
# Determine the type of the list element.
elem_type = None
if list_type == Type.LIST_INT:
elem_type = Type.INT
elif list_type == Type.LIST_FLOAT:
elem_type = Type.FLOAT
elif list_type == Type.LIST_STRING:
elem_type = Type.STRING
else:
error_str = f'expected list type but got {list_type}'
raise error.Type(expr.loc, error_str)
# Make sure the index has type int.
self.__validate_single_expr(expr.index)
self.__validate_type_of_expr(expr.index.loc, Type.INT, expr.index)
# Make sure the type of the new value matches the element type.
self.__validate_single_expr(expr.val)
self.__validate_type_of_expr(expr.val.loc, elem_type, expr.val)
def __validate_single_expr(self, expr):
''' Validate one parsed expression and add environment info. '''
if expr.exprClass == ExprEnum.LITERAL:
self.__validate_literal_expr_type(expr)
elif expr.exprClass == ExprEnum.CREATE_VAR:
self.__validate_create_var_expr_type(expr)
elif expr.exprClass == ExprEnum.SET_VAR:
self.__validate_set_var_expr_type(expr)
elif expr.exprClass == ExprEnum.GET_VAR:
self.__validate_get_var_expr_type(expr)
elif expr.exprClass == ExprEnum.DEFINE:
self.__validate_define_expr_type(expr)
elif expr.exprClass == ExprEnum.CALL:
self.__validate_call_expr_type(expr)
elif expr.exprClass == ExprEnum.IF:
self.__validate_if_expr_type(expr)
elif expr.exprClass == ExprEnum.LOOP:
self.__validate_loop_expr_type(expr)
elif expr.exprClass == ExprEnum.LIST:
self.__validate_list_expr_type(expr)
elif expr.exprClass == ExprEnum.LIST_AT:
self.__validate_list_at_expr_type(expr)
elif expr.exprClass == ExprEnum.LIST_SET:
self.__validate_list_set_expr_type(expr)
elif expr.exprClass == ExprEnum.PRIM_FUNC:
self.__validate_prim_func_type(expr)
else:
error_str = f'unknown expression type: {expr.exprClass}'
raise error.InternalError(expr.loc, error_str)
# Set the environment for this expressions if it was not already set.
if expr.env is None:
expr.env = self._env.copy()
def build_productions():
"""
From the productions string, creates the functions needed by yacc() to
generate the parser. Also completes the hierarchy of Node classes in the
ir module by creating Node subclasses for all non-terminal symbols that
don't already have one and setting the _symbols field (the list of symbols
in the relevant production) in each Node subclass.
"""
funcdefs = []
classdefs = []
for prod in productions.strip().split('\n\n'):
symbols = [s for s in prod.split() if s not in (':', '|', '%prec')]
prodname = symbols[0]
funcname = 'p_' + prodname
classname = ''.join([s.capitalize() for s in prodname.split('_')])
if not hasattr(ir, classname):
exec('class %s(Node): pass\n' % classname) in ir.__dict__
cls = getattr(ir, classname)
if (not isinstance(cls, type)) or (not issubclass(cls, ir.Node)):
raise error.InternalError('object %s is not a Node' % classname)
cls._symbols = set(symbols)
funcdoc = prod.replace('\n\t', ' ', 1)
funcdefs.append("def %s(p):\n '''%s'''\n p[0] = ir.%s(p)\n" %
(funcname, funcdoc, classname))
exec ''.join(funcdefs) in globals()
def type_to_str(loc, t):
if t == Type.INT:
return 'int'
elif t == Type.FLOAT:
return 'float'
elif t == Type.STRING:
return 'string'
else:
raise error.InternalError(loc, f'type_to_str: unknown type: {t}')
def __validate_get_var_expr_type(self, expr):
''' Typecheck and add environment data for a single GetVar Expr. '''
if expr.type != Type.UNDETERMINED:
error_str = f'expected GetVar to have UNDETERMINED type'
raise error.InternalError(expr.loc, error_str)
# Get the type of the variable and set expr.type.
expr.type = self._env.lookup_variable(expr.loc, expr.name)
def main():
if config.env <> 'VMware':
cmd = 'bash /etc/rc.key'
os.system(cmd)
cmd = 'rm -rf /dev/zstor'
os.system(cmd)
cmd = 'rm -rf /home/zonion/.llog'
os.system(cmd)
cmd = 'gpg --ignore-time-conflict --homedir /home/gpg --verify /home/zonion/license/key.sig > /tmp/key 2>&1'
_, o = execute(cmd, False, logging=False)
cmd = 'cat /tmp/key'
_, o = execute(cmd, False)
m = re.search('Good', o)
if not m:
cmd = 'touch /home/zonion/.llog'
_, o = execute(cmd, False)
print 'nokey'
cmd = 'cat /home/zonion/license/key'
_, o = execute(cmd, False, logging=False)
sn = util.get_sn()
if sn != o.strip():
cmd = 'touch /home/zonion/.llog'
_, o = execute(cmd, False)
print 'nokey'
cmd = 'rm -rf /tmp/key'
os.system(cmd)
cmd = 'rm /home/zonion/bitmap.disk/*.bitmap -rf'
execute(cmd, False)
try:
db.create_tables()
except Exception as e:
log.error(caused(e).detail)
raise error.InternalError()
if os.path.exists(config.boot.startup_file):
log.journal_warning('System may not be poweroffed normally!')
if 'monfs' not in config.feature:
cmd = 'bash /home/zonion/command/ipsan.sh'
os.system(cmd)
is_die = os.path.exists('/home/zonion/.llog')
if is_die:
log.info('NO license.......')
return
cxt = api.APIContext()
cxt.scan_all()
log.info('System startup.')
def __validate_if_expr_type(self, expr):
''' Typecheck and add environment data for a single If Expr. '''
if expr.type != Type.NONE:
error_str = f'expected If to have NONE type'
raise error.InternalError(expr.loc, error_str)
# An If expression has no type, so just typecheck its subexpressions.
self.__validate_single_expr(expr.cond)
for e in expr.then:
self.__validate_single_expr(e)
for e in expr.otherwise:
self.__validate_single_expr(e)
def enum_to_c_type(loc, type_enum):
''' Return a string of the C++ type for the type 'type_enum'. '''
if type_enum == Type.INT:
return 'int'
elif type_enum == Type.FLOAT:
return 'float'
elif type_enum == Type.STRING:
return 'char *'
elif type_enum == Type.LIST_INT:
return 'int_list'
elif type_enum == Type.LIST_FLOAT:
return 'float_list'
elif type_enum == Type.LIST_STRING:
return 'char **'
else:
error.InternalError(loc, f'bad enum_to_c_type param: {type_enum}')
def __validate_set_var_expr_type(self, expr):
''' Typecheck and add environment data for a single SetVar Expr. '''
if expr.type != Type.UNDETERMINED:
error_str = f'expected SetVar to have UNDETERMINED type'
raise error.InternalError(expr.loc, error_str)
# Get the type of the variable.
expr_type = self._env.lookup_variable(expr.loc, expr.name)
# Validate the expression that the value is set to.
self.__validate_single_expr(expr.val)
# Validate the types and update the expression type. There is no need
# to update the environment because no function/variable was created.
self.__validate_type_of_expr(expr.loc, expr_type, expr.val)
expr.type = expr_type
def pop_scope(self, loc):
''' Pop off 'env' until a None element is popped. '''
# Find the index of the last None element.
ind = None
for i in range(len(self.env) - 1, -1, -1):
if self.env[i] == None:
ind = i
break
# Make sure a None element was found.
if ind == None:
error_str = 'tried popping scope without previous scope'
raise error.InternalError(loc, error_str)
# Pop the scope.
self.env = self.env[:ind]
def __validate_prim_func_type(self, expr):
''' Add environment data for a single PrimFunc Expr. '''
# PrimFunc expressions don't have a body so there is not much to
# validate, but the environment must be updated so later code can call
# the primitive functions.
if expr.type == Type.UNDETERMINED or expr.type == Type.NONE:
error_str = f'primitive function {expr.name} has bad '
error_str += f'return type: {type_lst[0]}'
raise error.InternalError(expr.loc, error_str)
type_lst = [expr.type]
for arg_type in expr.arg_types:
type_lst.append(arg_type)
# Update the environment.
self._env.add(expr.name, type_lst, False)
def __validate_loop_expr_type(self, expr):
''' Typecheck and add environment data for a single Loop Expr. '''
if expr.type != Type.NONE:
error_str = f'expected Loop to have NONE type'
raise error.InternalError(expr.loc, error_str)
# The whole loop is in a new scope.
self._env.push_scope()
# A Loop expression has no type, so just typecheck its subexpressions.
self.__validate_single_expr(expr.init)
self.__validate_single_expr(expr.test)
self.__validate_single_expr(expr.update)
for e in expr.body:
self.__validate_single_expr(e)
# The body scope ended.
self._env.pop_scope(expr.loc)
def __validate_call_expr_type(self, expr):
''' Typecheck and add environment data for a single Call Expr. '''
if expr.type != Type.UNDETERMINED:
error_str = f'expected Call to have UNDETERMINED type'
raise error.InternalError(expr.loc, error_str)
# The print function can take many arguments after the initial string,
# so relax the type checking for it becuase optional arguments are not
# yet supported.
if expr.name == 'print':
# The first argument should be a string.
if len(expr.params) < 1:
error_str = f'print() expected one or more arguments'
raise error.Call(expr.loc, error_str)
arg1 = expr.params[0]
self.__validate_single_expr(arg1)
self.__validate_type_of_expr(arg1.loc, Type.STRING, arg1)
expr.type = Type.NONE
return
# The function is not print, so it must be in the environment.
# Get the type of the function.
(ret_type, args_types) = self._env.lookup_function(expr.loc, expr.name)
# Check that the number of arguments is correct.
if len(expr.params) != len(args_types):
error_str = f'{expr.name}() expected {len(args_types)} arguments '
error_str += f'but found {len(expr.params)}'
raise error.Call(expr.loc, error_str)
# Check that the arguments have the correct type.
for (param_e, arg_type) in zip(expr.params, args_types):
self.__validate_single_expr(param_e)
self.__validate_type_of_expr(param_e.loc, arg_type, param_e)
# Type checking succeeded.
expr.type = ret_type
def __init__(self, prod):
"Creates a new Node from prod (a yacc.YaccProduction instance)"
# Validate input (just a sanity check)
if not isinstance(prod, yacc.YaccProduction):
raise error.InternalError('expecting YaccProduction, got ' +
prod.__class__.__name__)
# Store line number and create child_list
self.lineno = prod.lineno(0)
self.child_list = list(prod)[1:]
#
# Create child_dict
#
self.child_dict = {}
for item in prod.slice[1:]:
if item.type not in self.child_dict:
# No entry for this symbol in child_dict, so make one
self.child_dict[item.type] = item.value
else:
# There's already an entry for this symbol in child_dict, so the
# entry's value must be a list
if isinstance(self.child_dict[item.type], list):
# Already a list, so append to it
self.child_dict[item.type].append(item.value)
else:
# Make list from previous and current value
self.child_dict[item.type] = ([
self.child_dict[item.type], item.value
])
# For each valid symbol that isn't involved in this instance, make an
# entry in child_dict with value None
for symbol in self._symbols:
if symbol not in self.child_dict:
self.child_dict[symbol] = None
def __validate_list_expr_type(self, expr):
''' Typecheck and add environment data for a single List Expr. '''
if expr.type != Type.UNDETERMINED:
error_str = f'expected List to have UNDETERMINED type'
raise error.InternalError(expr.loc, error_str)
# The size of the list should be an integer.
self.__validate_single_expr(expr.size)
self.__validate_type_of_expr(expr.size.loc, Type.INT, expr.size)
# Determine the type of the list from the element type.
if expr.elem_type == Type.INT:
expr.type = Type.LIST_INT
elif expr.elem_type == Type.FLOAT:
expr.type = Type.LIST_FLOAT
elif expr.elem_type == Type.STRING:
expr.type = Type.LIST_STRING
else:
raise error.InvalidType(expr.loc, expr.elem_type)
# Add the new list to the environment.
self._env.add(expr.name, [expr.type], True)
def lookup_function(self, loc, expr_name):
'''
Find a function in the environment.
Return a tuple with two elements. The first is the return type of the
function and the second is a list of types for the arguments.
'''
# Make sure the function already exists.
(name, type_lst, is_var) = self.get_entry_for_name(expr_name)
if name is None:
error_str = f'function {expr_name} does not exist'
raise error.Name(loc, error_str)
if is_var:
error_str = f'{expr_name} is a variable; expected a function'
raise error.Syntax(loc, error_str)
if len(type_lst) == 0:
error_str = f'{expr_name} has no return type'
raise error.InternalError(loc, error_str)
# Make sure the function has a return type.
ret_type = type_lst[0]
if ret_type == Type.UNDETERMINED or ret_type == Type.NONE:
error_str = f'function {expr_name} has bad return type: {ret_type}'
raise error.Type(loc, error_str)
# Make sure the parameter types are valid.
for param_type in type_lst[1:]:
if param_type == Type.UNDETERMINED or param_type == Type.NONE:
error_str = f'function {expr_name} has bad '
error_str += f'paramter type: {param_type}'
raise error.Type(expr.loc, error_str)
return (ret_type, type_lst[1:])
def __translate_expr(self, expr, end=True):
''' Get a single parsed expression and return the equivalent C++ and
CUDA code.
If 'end' is false, then the final characters of the expression,
like semi-colons and newlines, are not added.
'''
if expr.exprClass == ExprEnum.LITERAL:
return self.__translate_literal_expr(expr, end)
elif expr.exprClass == ExprEnum.CREATE_VAR:
return self.__translate_create_var_expr(expr, end)
elif expr.exprClass == ExprEnum.SET_VAR:
return self.__translate_set_var_expr(expr, end)
elif expr.exprClass == ExprEnum.GET_VAR:
return self.__translate_get_var_expr(expr, end)
elif expr.exprClass == ExprEnum.DEFINE:
return self.__translate_define_expr(expr, end)
elif expr.exprClass == ExprEnum.CALL:
return self.__translate_call_expr(expr, end)
elif expr.exprClass == ExprEnum.IF:
return self.__translate_if_expr(expr, end)
elif expr.exprClass == ExprEnum.LOOP:
return self.__translate_loop_expr(expr, end)
elif expr.exprClass == ExprEnum.LIST:
return self.__translate_list_expr(expr, end)
elif expr.exprClass == ExprEnum.LIST_AT:
return self.__translate_list_at_expr(expr, end)
elif expr.exprClass == ExprEnum.LIST_SET:
return self.__translate_list_set_expr(expr, end)
elif expr.exprClass == ExprEnum.PRIM_FUNC:
# There is no need to translate primitive functions into C++/CUDA.
return ('', '')
elif expr.exprClass == ExprEnum.PARA_LOOP:
return self.__translate_parallel_loop_expr(expr, end)
else:
error_str = f'unknown expression type: {expr.exprClass}'
raise error.InternalError(expr.loc, error_str)
misc = static.StaticBrowseableDir("/misc",
site_path("misc"),
misc_template,
root_path="/",
site_menu=site_menu,
site_menu_active_name="/misc")
# Static bits just taken from the static directory
static_resources = static.StaticDir("/", site_path("static"))
notfound_handler = error.NotFound(notfound_template,
root_path="/",
site_menu=site_menu,
site_menu_active_name=None)
internal_error_handler = error.InternalError(internal_error_template,
root_path="/",
site_menu=site_menu,
site_menu_active_name=None)
# Set up the application
app = web.application(urls, globals())
app.notfound = notfound_handler
if __name__ == "__main__":
# If running standalone, boot up the testing server
app.run()
else:
# If running via FCGI then this is probably production, show sanitised errors
app.internalerror = internal_error_handler
def __parse_define(self, start_point_loc):
'''
Private function to parse a single DEFINE expression. The _file
variable should be in a state starting with (without quotes):
'<name> (<arg1> <arg2> ...) <body1> <body2> ...)'
That is, the '(define ' section has alread been read.
Returns an instance of Define()
'''
# Get the return type and function name.
(l, def_return_type) = self.__parse_word_with_loc()
if def_return_type == 'list':
list_type = self.__parse_type()
if list_type == Type.INT:
def_return_type = Type.LIST_INT
elif list_type == Type.FLOAT:
def_return_type = Type.LIST_FLOAT
elif list_type == Type.STRING:
def_return_type = Type.LIST_STRING
else:
error_str = '__parse_define cannot get here'
raise error.InternalError(l, error_str)
else:
def_return_type = Type.str_to_type(l, def_return_type)
def_name = self.__parse_word()
# Parse the ':' between function name and start of arguments.
c = self.__eat_whitespace()
if c == '':
# The end of file was reached.
loc = self.__get_point_loc(True).span(self.__get_point_loc())
raise error.Syntax(loc, f'unexpected end of file')
if c != ':':
loc = self.__get_point_loc(True).span(self.__get_point_loc())
raise error.Syntax(loc, f'expected ":" but found "{c}"')
# Parse the arguments.
def_args = []
while True:
(l, maybe_type) = self.__parse_word_with_loc()
# Check if the end of the arguments has been reached.
if maybe_type == ':':
break
# There is another argument to parse.
arg_type = None
if maybe_type == 'list':
list_type = self.__parse_type()
if list_type == Type.INT:
arg_type = Type.LIST_INT
elif list_type == Type.FLOAT:
arg_type = Type.LIST_FLOAT
elif list_type == Type.STRING:
arg_type = Type.LIST_STRING
else:
error_str = '__parse_define cannot get here'
raise error.InternalError(l, error_str)
else:
arg_type = Type.str_to_type(l, maybe_type)
arg_name = self.__parse_word()
def_args.append((arg_type, arg_name))
# Parse the body expressions.
def_body = []
while True:
e = self.__parse_expr(parsing_exprs_list=True)
if e is None:
break
def_body.append(e)
# Now the entire function has been parsed.
loc = start_point_loc.span(self.__get_point_loc())
return Define(loc, def_return_type, def_name, def_args, def_body)
def __translate_parallel_loop_expr(self, expr, end=True):
''' Get a single parsed PARA_LOOP expression and return the equivalent
C++ and CUDA code.
If 'end' is false, then the final characters of the expression,
like semi-colons and newlines, are not added.
'''
def sub_expr_str(expr):
return f'{self.__translate_expr(expr, end=False)[0]}'
cpp = ''
cuda = ''
# Get a unique name for the cuda kernel.
self._para_loop_ind += 1
cuda_kernel_name = f'cuda_loop{self._para_loop_ind}_kernel'
# Determine the number of blocks and threads to use.
iters_str = f'({sub_expr_str(expr.end_index)} - ' + \
f'{sub_expr_str(expr.start_index)})'
threads_per_block = f'min(512, {iters_str})'
blocks = f'min(32, 1 + {iters_str} / {threads_per_block})'
# Setup the function to call the kernel.
args = []
for var_name in expr.used_vars:
var_type = expr.env.lookup_variable(expr.loc, var_name)
c_type = Type.enum_to_c_type(expr.loc, var_type)
args.append(f'{c_type} {var_name}')
cuda += f'void call_{cuda_kernel_name}'
cuda += f'({", ".join(args)})'
self.cuda_prototypes.append(cuda + ';\n')
# Call this kernel-calling fucntion in the cpp code.
cpp += f'call_{cuda_kernel_name}'
cpp += f'({", ".join(expr.used_vars)});\n'
cuda += ' {\n'
cuda_body = ''
# Make device variables if necessary.
dev_vars = []
for var_name in expr.used_vars:
dev_name = f'dev_{var_name}'
data_name = f'{dev_name}_data' # Only used for lists.
var_type = expr.env.lookup_variable(expr.loc, var_name)
if var_type == Type.INT or var_type == Type.FLOAT:
dev_vars.append(var_name)
elif var_type == Type.LIST_INT:
dev_vars.append(dev_name)
# Allocate memory.
size = f'{var_name}.size * sizeof(int)'
cuda_body += f'int *{data_name};\n'
cuda_body += f'cudaMalloc((void **) &{data_name}, {size});\n'
# Copy the data from host to device.
cuda_body += f'cudaMemcpy({data_name}, {var_name}.data, ' + \
f'{size}, cudaMemcpyHostToDevice);\n'
cuda_body += f'int_list {dev_name} = {"{"}{var_name}.size, ' + \
f'{data_name}{"}"};\n\n'
elif var_type == Type.LIST_FLOAT:
dev_vars.append(dev_name)
# Allocate memory.
size = f'{var_name}.size * sizeof(float)'
cuda_body += f'float *{data_name};\n'
cuda_body += f'cudaMalloc((void **) &{data_name}, {size});\n'
# Copy the data from host to device.
cuda_body += f'cudaMemcpy({data_name}, {var_name}.data, ' + \
f'{size}, cudaMemcpyHostToDevice);\n'
cuda_body += f'float_list {dev_name} = {"{"}' + \
f'{var_name}.size, {data_name}{"}"};\n\n'
elif var_type == Type.STRING or var_type == Type.LIST_STRING:
error_str = 'strings not yet allowed in parallelization'
raise error.InternalError(expr.loc, error_str)
# Call the kernel.
cuda_body += f'{cuda_kernel_name}<<<{blocks}, {threads_per_block}>>>'
cuda_body += f'({", ".join(dev_vars)});\n\n'
# Copy the data back from device to host.
for var_name in expr.used_vars:
dev_name = f'dev_{var_name}'
data_name = f'{dev_name}_data'
var_type = expr.env.lookup_variable(expr.loc, var_name)
if var_type == Type.INT or var_type == Type.FLOAT:
# There is no need to copy the variable back, because C++
# passes it by value and so the value did not change.
pass
elif var_type == Type.LIST_INT:
size = f'{var_name}.size * sizeof(int)'
# Copy the data from host to device.
cuda_body += f'cudaMemcpy({var_name}.data, {data_name}, {size}, '
cuda_body += f'cudaMemcpyDeviceToHost);\n'
elif var_type == Type.LIST_FLOAT:
size = f'{var_name}.size * sizeof(float)'
# Copy the data from host to device.
cuda_body += f'cudaMemcpy({var_name}.data, {data_name}, {size}, '
cuda_body += f'cudaMemcpyDeviceToHost);\n'
elif var_type == Type.STRING or var_type == Type.LIST_STRING:
error_str = 'strings not yet allowed in parallelization'
raise error.InternalError(expr.loc, error_str)
self._increase_indent()
cuda_body = self._make_indented(cuda_body)
self._decrease_indent()
cuda += cuda_body + '}\n\n'
# Setup the cuda code.
# No return value is needed because results are returned via the input
# lists.
cuda_kernel = f'__global__ void {cuda_kernel_name}'
# Add the arguments.
cuda_kernel += f'({", ".join(args)})'
# Add this function prototype for use in a header file.
self.cuda_prototypes.append(cuda_kernel + ';\n')
cuda_kernel += ' {\n'
# Determine the index in the loop.
index = expr.index_name
cuda_kernel += f' int {index} = blockIdx.x * blockDim.x + '
cuda_kernel += f'threadIdx.x + {sub_expr_str(expr.start_index)};\n\n'
# Loop over all indices that this thread is responsible for.
max_index = sub_expr_str(expr.end_index)
cuda_kernel += f' while ({index} < {max_index}) {"{"}\n'
for e in expr.body:
(c, _) = self.__translate_expr(e);
cuda_kernel += c
cuda_kernel += f' {index} += gridDim.x * blockDim.x;\n'
cuda_kernel += f' {"}"}\n'
cuda_kernel += f'{"}"}\n\n'
cuda += cuda_kernel
return (cpp, cuda)
