search

mindspore-ai / akg

AKG (Auto Kernel Generator) is an optimizer for operators in Deep Learning Networks, which provides the ability to automatically fuse ops with specific patterns.
https://gitee.com/mindspore/akg
Apache License 2.0
213 stars 38 forks source link

fail to run python test_all.py all #4

Closed wxyhv closed 3 years ago

wxyhv commented 3 years ago

Environment

Hardware Environment(/GPU RTX1060/):

Software Environment:

Describe the current behavior

Operater: fused_relu_grad Time of auto schedule: func_time_required func:random_gaussian, running:16.934264 seconds func_time_required func:random_gaussian, running:17.178366 seconds func_time_required func:random_gaussian, running:17.372735 seconds [ERROR] AKG:2021-03-15-13:49:27.480.610 [unify_loopvars.cc:111] [pass] found undefined variable: threadIdx.x Stack trace: [bt] (0) /akg/build/libakg.so(akg::ir::UnifyLoopVarsMutator::Mutate(air::Variable const, air::Expr const&)+0x3f7) [0x7f03ee75a027] [bt] (1) /akg/build/libakg.so(+0xed5bde) [0x7f03edcc0bde] [bt] (2) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (3) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (4) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Add const, air::Expr const&)+0x88) [0x7f03edcc3d68] [bt] (5) /akg/build/libakg.so(+0xed5d1e) [0x7f03edcc0d1e] [bt] (6) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (7) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (8) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Load const*, air::Expr const&)+0x51) [0x7f03edcc27e1]

Operater: fused_bn_update_grad Time of auto schedule: [ERROR] AKG:2021-03-15-13:49:27.718.533 [unify_loopvars.cc:111] [pass] found undefined variable: threadIdx.x Stack trace: [bt] (0) /akg/build/libakg.so(akg::ir::UnifyLoopVarsMutator::Mutate(air::Variable const, air::Expr const&)+0x3f7) [0x7f03ee75a027] [bt] (1) /akg/build/libakg.so(+0xed5bde) [0x7f03edcc0bde] [bt] (2) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (3) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (4) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Add const, air::Expr const&)+0x88) [0x7f03edcc3d68] [bt] (5) /akg/build/libakg.so(+0xed5d1e) [0x7f03edcc0d1e] [bt] (6) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (7) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (8) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Load const*, air::Expr const&)+0x51) [0x7f03edcc27e1]

Operater: fused_mul_div_rsqrt_mul_isfinite_red Time of auto schedule: func_time_required func:random_gaussian, running:0.140863 seconds func_time_required func:random_gaussian, running:0.150357 seconds [ERROR] AKG:2021-03-15-13:49:28.102.035 [unify_loopvars.cc:111] [pass] found undefined variable: threadIdx.x Stack trace: [bt] (0) /akg/build/libakg.so(akg::ir::UnifyLoopVarsMutator::Mutate(air::Variable const, air::Expr const&)+0x3f7) [0x7f03ee75a027] [bt] (1) /akg/build/libakg.so(+0xed5bde) [0x7f03edcc0bde] [bt] (2) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (3) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (4) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Load const, air::Expr const&)+0x51) [0x7f03edcc27e1] [bt] (5) /akg/build/libakg.so(+0xed5c2e) [0x7f03edcc0c2e] [bt] (6) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (7) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (8) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Mul const*, air::Expr const&)+0x53) [0x7f03edcc40f3]

Run op abs error using auto schedule: Traceback (most recent call last): File "test_all.py", line 529, in op(poly_sch=True, fuzz_shape=fuzz_shape) File "test_all.py", line 274, in abs test_ms_abs((1024, 1024), "float32", poly_sch=poly_sch) File "/akg/tests/operators/gpu/test_ms_abs.py", line 30, in test_ms_abs mod = utils.op_build_test(abs_data, [shape], [dtype], attrs={"target": "cuda"}, kernel_name="abs") File "/akg/python/akg/utils/kernel_exec.py", line 96, in wrapper result = func_name(*args, kwargs) File "/akg/python/akg/utils/kernel_exec.py", line 622, in op_build_test polyhedral, tuning) File "/akg/python/akg/utils/kernel_exec.py", line 1012, in op_build dump_code, tuning) File "/akg/python/akg/utils/kernel_exec.py", line 913, in create_gpu_mod binds=binds) File "/akg/python/akg/utils/validation_check.py", line 135, in in_wrapper return func(*args, *kwargs) File "/akg/python/akg/build_module.py", line 142, in build attrs=attrs, polyhedral=polyhedral, target=target) File "/akg/python/akg/utils/validation_check.py", line 135, in in_wrapper return func(args, kwargs) File "/akg/python/akg/build_module.py", line 135, in build_to_func polyhedral, target, cfg) File "/akg/third_party/incubator-tvm/python/tvm/_ffi/_ctypes/function.py", line 207, in call raise get_last_ffi_error() tvm.ffi.base.TVMError: Traceback (most recent call last): [bt] (8) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::ir::Load const, air::Expr const&)+0x51) [0x7f03edcc27e1] [bt] (7) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (6) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (5) /akg/build/libakg.so(+0xed5d1e) [0x7f03edcc0d1e] [bt] (4) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Add const, air::Expr const&)+0x88) [0x7f03edcc3d68] [bt] (3) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (2) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (1) /akg/build/libakg.so(+0xed5bde) [0x7f03edcc0bde] [bt] (0) /akg/build/libakg.so(akg::ir::UnifyLoopVarsMutator::Mutate_(air::Variable const*, air::Expr const&)+0x3f7) [0x7f03ee75a027] File "/home/xh/projects/akg-binary/src/pass/unify_loop_vars.cc", line 111 TVMError: found undefined variable: threadIdx.x

Run op add error using auto schedule: Traceback (most recent call last): File "test_all.py", line 529, in op(poly_sch=True, fuzz_shape=fuzz_shape) File "test_all.py", line 78, in add test_ms_add((1, 1024), (1, 1024), 'float32', poly_sch=poly_sch) File "/akg/tests/operators/gpu/test_ms_add.py", line 31, in test_ms_add mod = utils.op_build_test(add, (shape1, shape2), (dtype, dtype), kernel_name="add", attrs={"target": "cuda"}) File "/akg/python/akg/utils/kernel_exec.py", line 96, in wrapper result = func_name(*args, kwargs) File "/akg/python/akg/utils/kernel_exec.py", line 622, in op_build_test polyhedral, tuning) File "/akg/python/akg/utils/kernel_exec.py", line 1012, in op_build dump_code, tuning) File "/akg/python/akg/utils/kernel_exec.py", line 913, in create_gpu_mod binds=binds) File "/akg/python/akg/utils/validation_check.py", line 135, in in_wrapper return func(*args, *kwargs) File "/akg/python/akg/build_module.py", line 142, in build attrs=attrs, polyhedral=polyhedral, target=target) File "/akg/python/akg/utils/validation_check.py", line 135, in in_wrapper return func(args, kwargs) File "/akg/python/akg/build_module.py", line 135, in build_to_func polyhedral, target, cfg) File "/akg/third_party/incubator-tvm/python/tvm/_ffi/_ctypes/function.py", line 207, in call raise get_last_ffi_error() tvm.ffi.base.TVMError: Traceback (most recent call last): [bt] (8) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::ir::Add const, air::Expr const&)+0x53) [0x7f03edcc3d33] [bt] (7) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (6) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (5) /akg/build/libakg.so(+0xed5c2e) [0x7f03edcc0c2e] [bt] (4) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Load const, air::Expr const&)+0x51) [0x7f03edcc27e1] [bt] (3) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (2) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (1) /akg/build/libakg.so(+0xed5bde) [0x7f03edcc0bde] [bt] (0) /akg/build/libakg.so(akg::ir::UnifyLoopVarsMutator::Mutate_(air::Variable const*, air::Expr const&)+0x3f7) [0x7f03ee75a027] File "/home/xh/projects/akg-binary/src/pass/unify_loop_vars.cc", line 111 TVMError: found undefined variable: threadIdx.x

Steps to reproduce the issue

  1. get docker 1.1.2
  2. git clone akg repo
  3. build akg
  4. cd /akg/tests/operators/gpu
  5. python test_all.py all

Related log / screenshot

tvm.ffi.base.TVMError: Traceback (most recent call last): [bt] (8) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::ir::Add const, air::Expr const&)+0x53) [0x7f03edcc3d33] [bt] (7) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (6) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (5) /akg/build/libakg.so(+0xed5c2e) [0x7f03edcc0c2e] [bt] (4) /akg/build/libakg.so(air::ir::IRMutator::Mutate_(air::ir::Load const, air::Expr const&)+0x51) [0x7f03edcc27e1] [bt] (3) /akg/build/libakg.so(air::ir::IRMutator::Mutate(air::Expr)+0x5d) [0x7f03ed4d625d] [bt] (2) /akg/build/libakg.so(air::NodeFunctor<air::Expr (air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator)>::operator()(air::runtime::ObjectRef const&, air::Expr const&, air::ir::IRMutator) const+0x62) [0x7f03ed4d60e2] [bt] (1) /akg/build/libakg.so(+0xed5bde) [0x7f03edcc0bde] [bt] (0) /akg/build/libakg.so(akg::ir::UnifyLoopVarsMutator::Mutate_(air::Variable const*, air::Expr const&)+0x3f7) [0x7f03ee75a027] File "/home/xh/projects/akg-binary/src/pass/unify_loop_vars.cc", line 111 TVMError: found undefined variable: threadIdx.x

anyrenwei commented 3 years ago

HI wxyhv! Seems that you have built an ascend-backend version because the libakg.so linked akg-binary which is a library contained the ascend-backend passes. You can rebuild a gpu-backend version using "bash build.sh -e gpu".

anyrenwei commented 3 years ago

Here are some docs which could help you. https://gitee.com/mindspore/akg/blob/master/README.md https://gitee.com/mindspore/akg/wikis/Quick%20Start%20Tutorial%20for%20Generating%20Kernels%20Automatically?sort_id=3597685

anyrenwei commented 3 years ago

@wxyhv This issue will be closed if it can work normally.