(type_args, compute_capability)
return kernel
generate compound kernel for the optree from type_args
generate compound kernel for the optree from type_args
generate compound kernel for the optree from type_args
def _get_compound_kernel(type_args, compute_capability):
    """
    generate compound kernel for the optree from type_args
    """

    # from the stack, rebuild a mutable tree
    tree = _build_tree(type_args)
    # _print_tree(tree)
    # exit()

    # split all reductions and post reduction scalar operations out of the tree
    # sub-trees are converted to stacks and pushed onto stages list
    stages = _split_stages(tree)
    # _print_tree(tree)
    # exit()

    # set the final stage type to type of output (scalar or elementwise)
    last_stage = "red_out" if tree[1] == 1 else "ew_out"
    # convert the remainder of tree to stack
    stages.append((last_stage, _post_order(tree)))

    # for stage, stage_data in enumerate(stages):
    #     print stage_data[0], stage
    #     for s in stage_data[1]: print s
    #     print
    # exit()

    stack = list()
    placeholders = list()
    stage_out_reg = dict()
    arg_dict = dict()
    array_ids = set()
    fp16In = False
    rand_init = False
    rand_func = False
    threads = type_args[-1][3]
    template = _ew_template
    template_vals = {
        "threads": threads,
        "name": _get_kernel_name(),
        "common": list(),
        "inits": list(),
        "finish": list(),
    }

    for stage, stage_data in enumerate(stages):

        stage_type, stage_stack = stage_data
        new_placeholders = list()

        # build out the template as we process stages
        if stage_type == "reduction":

            new_placeholders.append("loads%d" % stage)
            new_placeholders.append("ops%d" % stage)
            new_placeholders.append("shfl_red%d" % stage)
            template += _stage_template["loop"].format(stage)
            if threads > 32:
                new_placeholders.append("var_red%d" % stage)
                new_placeholders.append("share1_red%d" % stage)
                new_placeholders.append("share2_red%d" % stage)
                template += _stage_template["red"].format(stage)
            else:
                template += _stage_template["red32"].format(stage)

        elif stage_type == "scalar":

            new_placeholders.append("ops%d" % stage)
            template += _stage_template["red_ops"].format(stage)

        elif stage_type == "red_out":

            new_placeholders.append("ops%d" % stage)
            template += _stage_template["red_out"].format(stage)

        else:  # ew_out

            new_placeholders.append("loads%d" % stage)
            new_placeholders.append("ops%d" % stage)
            template += _stage_template["loop"].format(stage)

        for key in new_placeholders:
            template_vals[key] = []
        placeholders.extend(new_placeholders)

        for arg_i, arg in enumerate(stage_stack):

            arg_type, arg_id = arg[0:2]

            # Array operands
            if arg_type is ng.GPUTensor:

                dtype, take_axis = arg[2:4]

                is_out_tensor = True if stage == len(
                    stages) - 1 and arg_i == 0 else False

                # first arg is output array, don't put on stack
                if is_out_tensor:
                    out_dtype = dtype
                    out_take = take_axis
                else:
                    stack.append("a%d" % arg_id)

                # 0: arg_id, 1: stage, 2: type, 3: cvt
                ew_dtype = _ew_types[dtype]
                fmt = (arg_id, stage, ew_dtype["type"], ew_dtype["cvt"])

                # First time we see a tensor initialize everything
                if arg_id not in array_ids:

                    array_ids.add(arg_id)
                    array_ids.add((arg_id, stage))

                    sig = "Pii"
                    if take_axis > 0:
                        sig += "P"

                    # output tensor
                    if is_out_tensor:
                        ew_out = _ew_strings["out%d" % take_axis]
                        arguments = ew_out["arguments"].format(*fmt)
                        template_vals["inits"].append(
                            ew_out["inits"].format(*fmt))
                    # input tensors
                    else:
                        ew_in = _ew_strings["in%d" % take_axis]
                        loads = "loads%d" % stage
                        arguments = ew_in["arguments"].format(*fmt)
                        template_vals["inits"].append(
                            ew_in["inits"].format(*fmt))
                        template_vals[loads].append(
                            ew_in["loads"].format(*fmt))

                    if dtype == 'f2' and not fp16In:
                        template_vals["common"].append(_common_fp16_to_fp32)
                        fp16In = True

                    arg_dict[arg] = (sig, arguments)

                # Subsequent times we see a tensor just initialize inits and
                # loads
                elif (arg_id, stage) not in array_ids:
                    array_ids.add((arg_id, stage))
                    ew_in = _ew_strings["in%d" % take_axis]
                    loads = "loads%d" % stage
                    template_vals["inits"].append(ew_in["inits"].format(*fmt))
                    template_vals[loads].append(ew_in["loads"].format(*fmt))

            # Constant operands
            elif arg_type is float:

                stack.append("c%d" % arg_id)
                if arg not in arg_dict:
                    arg_dict[arg] = (
                        "f", _ew_strings["const"]["arguments"].format(arg_id))

            # Operations (arg_type = op_name)
            else:

                if arg_type == "assign":

                    ops = "ops%d" % stage

                    # loop end condition for last stage
                    sig = "i"
                    arguments = ["const int n%d" % stage]

                    # rounding mode
                    if arg[2]:
                        mode = "random"
                        sig += "i"
                        arguments.append("const int mantissa_bits")
                        if not rand_init:
                            rand_init = _init_rand(template_vals)
                        template_vals["inits"].append(_init_rand_round_func)
                    else:
                        mode = "nearest"

                    arg_dict[arg] = (sig, ", ".join(arguments))

                    out_val = stack.pop()
                    # if the last stack value came from an argmax/min just do
                    # implicit type conversion
                    if out_val[0] == "i" and out_dtype[0] in "iu":
                        ew_round = None
                    else:
                        ew_round = _ew_strings["round"][
                            mode].get(out_dtype, None)
                        ew_common = _common_round[mode].get(out_dtype, None)
                        if ew_common:
                            template_vals["common"].append(ew_common)

                    if ew_round:
                        round_val = "r%d" % arg_id
                        template_vals[ops].append(
                            ew_round.format(round_val, out_val))
                    else:
                        round_val = out_val

                    template_vals[ops].append(
                        _ew_strings["out%d" % out_take]["output"].format(round_val))

                elif arg in stage_out_reg:

                    stack.append(stage_out_reg[arg])

                elif arg_type in _float_ops:

                    if len(template_vals["name"]) < 16:
                        template_vals["name"].append(arg_type)

                    ops = "ops%d" % stage

                    (num_ops, op_code) = _float_ops[arg_type]

                    if arg_type == "rand":
                        if not rand_init:
                            rand_init = _init_rand(template_vals)
                        if not rand_func:
                            template_vals["common"].append(_common_frand)
                            rand_func = True

                    op_list = ["r%d" % arg_id]

                    # build the operands from the stack
                    for i in range(num_ops):
                        op_list.append(stack.pop())

                    if arg_type == "onehot":

                        hot_axis = arg[2]
                        test_val = "i" if hot_axis else "bid"

                        ew_in = _ew_strings[arg_type + native_str(hot_axis)]
                        loads = "loads%d" % stage
                        template_vals["inits"].append(
                            ew_in["inits"].format(arg_id))
                        template_vals[loads].append(
                            ew_in["loads"].format(arg_id))
                        op_list.append("onehot%d" % arg_id)
                        op_list.append(test_val)

                        arg_dict[arg] = (
                            "P", ew_in["arguments"].format(arg_id))

                    template_vals[ops].append(op_code.format(*op_list))

                    # if this is the last op on the current stack, store its register stage
                    # in the stage output dict
                    if arg_i == len(stage_stack) - 1:
                        stage_out_reg[arg] = op_list[0]
                    # otherwise push the reg onto the stack as normal
                    else:
                        stack.append(op_list[0])

                elif arg_type in _reduction_ops:

                    if len(template_vals["name"]) < 16:
                        template_vals["name"].append(arg_type)

                    # loop end condition for current stage
                    # add regardless of duplicate reduction stage
                    arg_dict[arg] = ("i", "const int n%d" % stage)

                    # avoid float conversion for argmax/min
                    reg = "i" if "arg" == arg_type[0:3] else "r"

                    ops = "ops%d" % stage
                    shfl_red = "shfl_red%d" % stage
                    red_arg = "%s%d" % (reg, arg_id)
                    red_strings = _reduction_ops[arg_type]
                    stack_arg = stack.pop()

                    template_vals["inits"].append(
                        red_strings["inits"].format(red_arg))
                    template_vals[ops].append(
                        red_strings["ops"].format(red_arg, stack_arg))
                    template_vals[shfl_red].append(
                        red_strings["shfl_red"].format(red_arg))
                    if threads > 32:
                        var_red = "var_red%d" % stage
                        shr1_red = "share1_red%d" % stage
                        shr2_red = "share2_red%d" % stage
                        template_vals[var_red].append(red_arg)
                        template_vals[shr1_red].append(
                            red_strings["share1_red"].format(red_arg))
                        template_vals[shr2_red].append(
                            red_strings["share2_red"].format(red_arg))

                    # reduction ops are always the last on the stack
                    # just store the register state in the stage output dict
                    stage_out_reg[arg] = red_arg

                else:
                    raise ValueError("Bad op type.")

    if (compute_capability[0] == 3 and compute_capability[1] < 5) or compute_capability[0] < 3:
        template_vals["common"].append(_common_kepler)

    template += _fin_template

    # since we reorderd the operations we need to generate the argument list
    # in the original order
    sig = "P"
    arguments = list()
    unused = 1
    for arg in type_args:
        params = arg_dict.get(arg, False)
        if params:
            sig += params[0]
            arguments.append(params[1])
            del arg_dict[arg]
        # fill in the loop counter for the duplicate reductions that were
        # removed
        elif arg[0] in _reduction_ops:
            sig += "i"
            arguments.append("const int unused%d" % unused)
            unused += 1

    # convert lists to strings
    template_vals["name"] = "_".join(template_vals["name"])
    template_vals["common"] = "\n".join(template_vals["common"])
    template_vals["arguments"] = ",\n    ".join(arguments)
    template_vals["inits"] = "\n    ".join(template_vals["inits"])
    template_vals["finish"] = "\n".join(template_vals["finish"])

    # add the dynamic placeholders: loads#, ops#, reduction#
    for key in placeholders:
        template_vals[key] = "\n        ".join(template_vals[key])

    # populate the template
    code = template % template_vals

    # debugging:
    # print "Compiling %s" % template_vals["name"]
    # f = open("kernel.cu", "w")
    # f = open("%s.cu" % template_vals["name"], "w")
    # print >>f, code
    # f.close()

    # ,"-G" , keep=False
    # module = SourceModule(code, options=["--use_fast_math"])
    module = SourceModule(code, options=[])
    kernel = module.get_function(template_vals["name"])
    kernel.name = template_vals["name"]
    kernel.prepare(sig)

    return kernel
def _get_compound_kernel ( type_args , compute_capability ) : # from the stack, rebuild a mutable tree tree = _build_tree ( type_args ) # _print_tree(tree) # exit() # split all reductions and post reduction scalar operations out of the tree # sub-trees are converted to stacks and pushed onto stages list stages = _split_stages ( tree ) # _print_tree(tree) # exit() # set the final stage type to type of output (scalar or elementwise) last_stage = "red_out" if tree [ 1 ] == 1 else "ew_out" # convert the remainder of tree to stack stages . append ( ( last_stage , _post_order ( tree ) ) ) # for stage, stage_data in enumerate(stages): #     print stage_data[0], stage #     for s in stage_data[1]: print s #     print # exit() stack = list ( ) placeholders = list ( ) stage_out_reg = dict ( ) arg_dict = dict ( ) array_ids = set ( ) fp16In = False rand_init = False rand_func = False threads = type_args [ - 1 ] [ 3 ] template = _ew_template template_vals = { "threads" : threads , "name" : _get_kernel_name ( ) , "common" : list ( ) , "inits" : list ( ) , "finish" : list ( ) , } for stage , stage_data in enumerate ( stages ) : stage_type , stage_stack = stage_data new_placeholders = list ( ) # build out the template as we process stages if stage_type == "reduction" : new_placeholders . append ( "loads%d" % stage ) new_placeholders . append ( "ops%d" % stage ) new_placeholders . append ( "shfl_red%d" % stage ) template += _stage_template [ "loop" ] . format ( stage ) if threads > 32 : new_placeholders . append ( "var_red%d" % stage ) new_placeholders . append ( "share1_red%d" % stage ) new_placeholders . append ( "share2_red%d" % stage ) template += _stage_template [ "red" ] . format ( stage ) else : template += _stage_template [ "red32" ] . format ( stage ) elif stage_type == "scalar" : new_placeholders . append ( "ops%d" % stage ) template += _stage_template [ "red_ops" ] . format ( stage ) elif stage_type == "red_out" : new_placeholders . append ( "ops%d" % stage ) template += _stage_template [ "red_out" ] . format ( stage ) else : # ew_out new_placeholders . append ( "loads%d" % stage ) new_placeholders . append ( "ops%d" % stage ) template += _stage_template [ "loop" ] . format ( stage ) for key in new_placeholders : template_vals [ key ] = [ ] placeholders . extend ( new_placeholders ) for arg_i , arg in enumerate ( stage_stack ) : arg_type , arg_id = arg [ 0 : 2 ] # Array operands if arg_type is ng . GPUTensor : dtype , take_axis = arg [ 2 : 4 ] is_out_tensor = True if stage == len ( stages ) - 1 and arg_i == 0 else False # first arg is output array, don't put on stack if is_out_tensor : out_dtype = dtype out_take = take_axis else : stack . append ( "a%d" % arg_id ) # 0: arg_id, 1: stage, 2: type, 3: cvt ew_dtype = _ew_types [ dtype ] fmt = ( arg_id , stage , ew_dtype [ "type" ] , ew_dtype [ "cvt" ] ) # First time we see a tensor initialize everything if arg_id not in array_ids : array_ids . add ( arg_id ) array_ids . add ( ( arg_id , stage ) ) sig = "Pii" if take_axis > 0 : sig += "P" # output tensor if is_out_tensor : ew_out = _ew_strings [ "out%d" % take_axis ] arguments = ew_out [ "arguments" ] . format ( * fmt ) template_vals [ "inits" ] . append ( ew_out [ "inits" ] . format ( * fmt ) ) # input tensors else : ew_in = _ew_strings [ "in%d" % take_axis ] loads = "loads%d" % stage arguments = ew_in [ "arguments" ] . format ( * fmt ) template_vals [ "inits" ] . append ( ew_in [ "inits" ] . format ( * fmt ) ) template_vals [ loads ] . append ( ew_in [ "loads" ] . format ( * fmt ) ) if dtype == 'f2' and not fp16In : template_vals [ "common" ] . append ( _common_fp16_to_fp32 ) fp16In = True arg_dict [ arg ] = ( sig , arguments ) # Subsequent times we see a tensor just initialize inits and # loads elif ( arg_id , stage ) not in array_ids : array_ids . add ( ( arg_id , stage ) ) ew_in = _ew_strings [ "in%d" % take_axis ] loads = "loads%d" % stage template_vals [ "inits" ] . append ( ew_in [ "inits" ] . format ( * fmt ) ) template_vals [ loads ] . append ( ew_in [ "loads" ] . format ( * fmt ) ) # Constant operands elif arg_type is float : stack . append ( "c%d" % arg_id ) if arg not in arg_dict : arg_dict [ arg ] = ( "f" , _ew_strings [ "const" ] [ "arguments" ] . format ( arg_id ) ) # Operations (arg_type = op_name) else : if arg_type == "assign" : ops = "ops%d" % stage # loop end condition for last stage sig = "i" arguments = [ "const int n%d" % stage ] # rounding mode if arg [ 2 ] : mode = "random" sig += "i" arguments . append ( "const int mantissa_bits" ) if not rand_init : rand_init = _init_rand ( template_vals ) template_vals [ "inits" ] . append ( _init_rand_round_func ) else : mode = "nearest" arg_dict [ arg ] = ( sig , ", " . join ( arguments ) ) out_val = stack . pop ( ) # if the last stack value came from an argmax/min just do # implicit type conversion if out_val [ 0 ] == "i" and out_dtype [ 0 ] in "iu" : ew_round = None else : ew_round = _ew_strings [ "round" ] [ mode ] . get ( out_dtype , None ) ew_common = _common_round [ mode ] . get ( out_dtype , None ) if ew_common : template_vals [ "common" ] . append ( ew_common ) if ew_round : round_val = "r%d" % arg_id template_vals [ ops ] . append ( ew_round . format ( round_val , out_val ) ) else : round_val = out_val template_vals [ ops ] . append ( _ew_strings [ "out%d" % out_take ] [ "output" ] . format ( round_val ) ) elif arg in stage_out_reg : stack . append ( stage_out_reg [ arg ] ) elif arg_type in _float_ops : if len ( template_vals [ "name" ] ) < 16 : template_vals [ "name" ] . append ( arg_type ) ops = "ops%d" % stage ( num_ops , op_code ) = _float_ops [ arg_type ] if arg_type == "rand" : if not rand_init : rand_init = _init_rand ( template_vals ) if not rand_func : template_vals [ "common" ] . append ( _common_frand ) rand_func = True op_list = [ "r%d" % arg_id ] # build the operands from the stack for i in range ( num_ops ) : op_list . append ( stack . pop ( ) ) if arg_type == "onehot" : hot_axis = arg [ 2 ] test_val = "i" if hot_axis else "bid" ew_in = _ew_strings [ arg_type + native_str ( hot_axis ) ] loads = "loads%d" % stage template_vals [ "inits" ] . append ( ew_in [ "inits" ] . format ( arg_id ) ) template_vals [ loads ] . append ( ew_in [ "loads" ] . format ( arg_id ) ) op_list . append ( "onehot%d" % arg_id ) op_list . append ( test_val ) arg_dict [ arg ] = ( "P" , ew_in [ "arguments" ] . format ( arg_id ) ) template_vals [ ops ] . append ( op_code . format ( * op_list ) ) # if this is the last op on the current stack, store its register stage # in the stage output dict if arg_i == len ( stage_stack ) - 1 : stage_out_reg [ arg ] = op_list [ 0 ] # otherwise push the reg onto the stack as normal else : stack . append ( op_list [ 0 ] ) elif arg_type in _reduction_ops : if len ( template_vals [ "name" ] ) < 16 : template_vals [ "name" ] . append ( arg_type ) # loop end condition for current stage # add regardless of duplicate reduction stage arg_dict [ arg ] = ( "i" , "const int n%d" % stage ) # avoid float conversion for argmax/min reg = "i" if "arg" == arg_type [ 0 : 3 ] else "r" ops = "ops%d" % stage shfl_red = "shfl_red%d" % stage red_arg = "%s%d" % ( reg , arg_id ) red_strings = _reduction_ops [ arg_type ] stack_arg = stack . pop ( ) template_vals [ "inits" ] . append ( red_strings [ "inits" ] . format ( red_arg ) ) template_vals [ ops ] . append ( red_strings [ "ops" ] . format ( red_arg , stack_arg ) ) template_vals [ shfl_red ] . append ( red_strings [ "shfl_red" ] . format ( red_arg ) ) if threads > 32 : var_red = "var_red%d" % stage shr1_red = "share1_red%d" % stage shr2_red = "share2_red%d" % stage template_vals [ var_red ] . append ( red_arg ) template_vals [ shr1_red ] . append ( red_strings [ "share1_red" ] . format ( red_arg ) ) template_vals [ shr2_red ] . append ( red_strings [ "share2_red" ] . format ( red_arg ) ) # reduction ops are always the last on the stack # just store the register state in the stage output dict stage_out_reg [ arg ] = red_arg else : raise ValueError ( "Bad op type." ) if ( compute_capability [ 0 ] == 3 and compute_capability [ 1 ] < 5 ) or compute_capability [ 0 ] < 3 : template_vals [ "common" ] . append ( _common_kepler ) template += _fin_template # since we reorderd the operations we need to generate the argument list # in the original order sig = "P" arguments = list ( ) unused = 1 for arg in type_args : params = arg_dict . get ( arg , False ) if params : sig += params [ 0 ] arguments . append ( params [ 1 ] ) del arg_dict [ arg ] # fill in the loop counter for the duplicate reductions that were # removed elif arg [ 0 ] in _reduction_ops : sig += "i" arguments . append ( "const int unused%d" % unused ) unused += 1 # convert lists to strings template_vals [ "name" ] = "_" . join ( template_vals [ "name" ] ) template_vals [ "common" ] = "\n" . join ( template_vals [ "common" ] ) template_vals [ "arguments" ] = ",\n    " . join ( arguments ) template_vals [ "inits" ] = "\n    " . join ( template_vals [ "inits" ] ) template_vals [ "finish" ] = "\n" . join ( template_vals [ "finish" ] ) # add the dynamic placeholders: loads#, ops#, reduction# for key in placeholders : template_vals [ key ] = "\n        " . join ( template_vals [ key ] ) # populate the template code = template % template_vals # debugging: # print "Compiling %s" % template_vals["name"] # f = open("kernel.cu", "w") # f = open("%s.cu" % template_vals["name"], "w") # print >>f, code # f.close() # ,"-G" , keep=False # module = SourceModule(code, options=["--use_fast_math"]) module = SourceModule ( code , options = [ ] ) kernel = module . get_function ( template_vals [ "name" ] ) kernel . name = template_vals [ "name" ] kernel . prepare ( sig ) return kernel