2062a98c创建于 2024年2月19日历史提交
#!/usr/bin/env python
# -*- coding: UTF-8 -*-
"""
Created on Feb  28 20:56:45 2020
Copyright (c) Huawei Technologies Co., Ltd. 2020-2021. All rights reserved.
"""


def gen_fun_def(title, kernel, argn, arg_type, arg_name):
    entry = []
    entry.append(title)
    entry.append(kernel)
    entry.append('(')
    args = []
    for i in range(0, argn):
        args.append(arg_type + ' ' + arg_name + str(i))
    entry.append(', '.join(args))
    entry.append(')')
    return ' '.join(entry)


def gen_batch_kernel_body(fname, argn, arg_name):
    body = []
    body.append('{')
    fun = []
    fun.append(fname)
    fun.append('(')
    args = []
    for i in range(0, argn):
        args.append(arg_name + str(i))
    fun.append(', '.join(args))
    fun.append(');')
    body.append(' '.join(fun))
    body.append('}')
    return '\n'.join(body)


def gen_mc_kernel_body(kn, argn, arg_name, blknum):
    body = []
    body.append('{')
    body.append('    switch(block_idx) {')
    for blk in range(0, blknum):
        fun = []
        fun.append('{}_blk{:02d}'.format(kn, blk))
        fun.append('(')
        args = []
        for i in range(0, argn):
            args.append(arg_name + str(i))
        fun.append(', '.join(args))
        fun.append(')')
        body.append('        case {}: {}; break;'.format(blk, ' '.join(fun)))
    body.append('        default: break;')
    body.append('    }')
    body.append('}')
    return '\n'.join(body)


def gen_proc_body(argn, arg_name):
    body = []
    body.append('{')
    args = []
    for i in range(0, argn):
        args.append(arg_name + str(i))
    body.append('uint64_t __x = (uint64_t)' + ' + (uint64_t)'.join(args) + ';')
    body.append('__asm__ ("NOP");')
    body.append('__asm__ ("NOP");')
    body.append('__asm__ ("NOP");')
    body.append('}')
    return '\n'.join(body)


def batch_code_gen(kn, argn, argt):
    codes = []
    kernel_name = kn
    proc_name = kernel_name + '_percore'
    arg_num = int(argn)
    data_type = argt
    arg_type = '__gm__ ' + data_type + '* __restrict__'
    arg_name = 'arg'
    kernel_title = 'extern \"C\" __global__ __aicore__ void'
    proc_title = 'extern \"C\" __attribute__((noinline)) __aicore__ void'
    codes.append('#ifndef __aicore__')
    codes.append('#define __aicore__ [aicore]')
    codes.append('#endif')
    codes.append(gen_fun_def(proc_title, proc_name, arg_num, arg_type, arg_name) + ';')
    codes.append(gen_fun_def(kernel_title, kernel_name, arg_num, arg_type, arg_name))
    codes.append(gen_batch_kernel_body(proc_name, arg_num, arg_name))
    codes.append(gen_fun_def(proc_title, proc_name, arg_num, arg_type, arg_name))
    codes.append(gen_proc_body(arg_num, arg_name))
    return '\n'.join(codes) + '\n'


def mc_code_gen(kn, argn, argt, blknum):
    codes = []
    kernel_name = kn
    core_num = int(blknum)
    arg_num = int(argn)
    data_type = argt
    arg_type = '__gm__ ' + data_type + '* __restrict__'
    arg_name = 'arg'
    kernel_title = 'extern \"C\" __global__ __aicore__ void'
    proc_title = 'extern \"C\" __attribute__((noinline)) __aicore__ void'
    codes.append('#ifndef __aicore__')
    codes.append('#define __aicore__ [aicore]')
    codes.append('#endif')
    for i in range(0, core_num):
        proc_name = '{}_blk{:02d}'.format(kernel_name, i)
        codes.append(gen_fun_def(proc_title, proc_name, arg_num, arg_type, arg_name) + ';')
    codes.append(gen_fun_def(kernel_title, kernel_name, arg_num, arg_type, arg_name))
    codes.append(gen_mc_kernel_body(kernel_name, arg_num, arg_name, core_num))
    for i in range(0, core_num):
        proc_name = '{}_blk{:02d}'.format(kernel_name, i)
        codes.append(gen_fun_def(proc_title, proc_name, arg_num, arg_type, arg_name))
        codes.append(gen_proc_body(arg_num, arg_name))
    return '\n'.join(codes) + '\n'