[MetaSchedule] [TensorCore]Please help check out this error

MasterJianxing · April 10, 2023, 5:29am

I try to use tensorcore to tune a network. To use tensorcore, I set datatype as “float16”, and I find this error. Besides when I set datatype as “float32”, it runs normally.

Traceback (most recent call last): File “resnet_meta.py”, line 58, in database = ms.tune.tune_tasks( File “/home/pan/tvm/python/tvm/meta_schedule/tune.py”, line 117, in tune_tasks task_scheduler.tune( File “/home/pan/tvm/python/tvm/meta_schedule/task_scheduler/task_scheduler.py”, line 132, in tune _ffi_api.TaskSchedulerTune( # type: ignore # pylint: disable=no-member File “/home/pan/tvm/python/tvm/_ffi/_ctypes/packed_func.py”, line 237, in call raise get_last_ffi_error() tvm.tir.schedule.schedule.ScheduleError: Traceback (most recent call last):

ScheduleError: An error occurred in the schedule primitive ‘compute-at’ … Error message: The scope tir.Block#0 is not a stage pipeline.

I imitate the testing file to write my resnet metaschedule tune file, here is the code:

mod, params = testing.resnet.get_workload( num_layers=50, batch_size=batch_size, image_shape=image_shape, dtype=“float16” )

tune_tasks = ms.relay_integration.extract_tasks(mod, tgt, params)

tasks, task_weights = ms.relay_integration.extracted_tasks_to_tune_contexts( extracted_tasks=tune_tasks, work_dir=work_dir, space=ms.space_generator.PostOrderApply( sch_rules=“cuda-tensorcore”, postprocs=“cuda-tensorcore”, mutator_probs=“cuda-tensorcore”))

database = ms.tune.tune_tasks( tasks=tasks, task_weights=task_weights, work_dir=work_dir, max_trials_per_task=4, max_trials_global=150, )

Please help me check out why this error happens

Many thanks.

junrushao · April 10, 2023, 7:24am

Could you please share the TIR you are tuning?

MasterJianxing · April 10, 2023, 8:09am

Here it is. Please check @junrushao

Besides, here is the original file

github.com

MasterJianxing/metaschedule/blob/main/resnet_meta.py

import numpy as np
import pytest
import tvm
#import tvm.testing
from tvm.contrib import graph_executor
from tvm.relay import testing
from tvm import meta_schedule as ms
from tvm import relay, auto_scheduler
from tvm.meta_schedule.testing import relay_workload
from tvm.meta_schedule.testing.tlcbench import load_quantized_bert_base
from tvm.tir.tensor_intrin import *

batch_size = 1
num_class = 1000
image_shape = (3, 224, 224)
data_shape = (batch_size,) + image_shape
out_shape = (batch_size, num_class)
dtype = "float16"

mod, params = testing.resnet.get_workload(

This file has been truncated. show original

junrushao · April 10, 2023, 9:24pm

I have personally switched to Relax in the unity branch, and my bandwidth is very limited to supporting Relay at this moment, so i cannot guarantee if Relay works or not.

in your specific case, it seems that the stage pipeline assumption is broken. would you like to attach the log file to this workload so that we could investigate without using Relay?

MasterJianxing · April 11, 2023, 5:04am

Here are the tuning logs and print logs, I just upload them.

Thanks for your patience! Please check

MasterJianxing · April 13, 2023, 2:23am

@junrushao Log files are in tune_tmp and log.log, Could you mind having a look?

Many thanks

JackWw · July 26, 2023, 6:32am

Do you solve this problem? i meet the similar problem when tuning fused_matmul_add

Krishna · December 4, 2023, 9:11am

Hello @MasterJianxing ,

I ran the resnet_meta.py code you had shared on my system and I didnt seem to get the same error as you, but I received this error:

 InternalError: Check failed: original_producers.size() == 1u (0 vs. 1) :

The full diagnostic is given below:

Can you please take a look and help as to why this InteralError is occurring? Please help.

Thanks and reagrds,

Krishna

Krishna · March 1, 2024, 6:38am

Hi @MasterJianxing Were you able to fix this? I read about a somehwat similar issue here, please take a look. Thanks.

@zxybazh @junrushao @AndrewZhaoLuo @comaniac Please help.

MasterJianxing · March 1, 2024, 7:31am

It’s too long time… But now i can use metaschedule to tune BERT with latest tvm version. You can update and have a try

Krishna · March 1, 2024, 9:44am

Okay, thank you for the response. Did you tune BERT in fp-16 by targeting cuda-tensorcores?

Krishna · March 4, 2024, 6:46am

Hi, Sorry to bother you, can you please share the metascheduling part of your code? I just wanted to know how you got Metacsheduling to work with BERT. It would greatly help me in my current work. Thanks in advance.

MasterJianxing · March 8, 2024, 11:57am

github.com

MasterJianxing/metaschedule/blob/main/transformer.py

import numpy as np
import pytest
import tvm
import onnx
from tvm.contrib import graph_executor
from tvm import meta_schedule as ms
from tvm import relay, auto_scheduler
from tvm.meta_schedule.testing import relay_workload
from tvm.meta_schedule.testing.tlcbench import load_quantized_bert_base
from tvm.tir.tensor_intrin.cuda import *
from tvm.tir.tensor_intrin.arm_cpu import DP4A_INTRIN
from tvm.tir.tensor_intrin.rocm import AMDGPU_SDOT4_INTRIN
from tvm.tir.tensor_intrin.x86 import VNNI_DOT_16x4_INTRIN as VNNI_INTRIN

@tvm.testing.requires_gpu
@pytest.mark.skip("Slow on CI")
@pytest.mark.parametrize(
    ["model_name", "input_shape"],
    [("bert_base", (8, 128)), ("resnet_18", (16, 3, 224, 224)), ("resnet_50", (16, 3, 224, 224))],
)

This file has been truncated. show original

Krishna · March 18, 2024, 11:34am

Hi @MasterJianxing ,

Thank you so much for sharing the code. I ran the transformer.py file for the BERT model and I am currently facing this RuntimeError in thread bindings. Error :

RuntimeError: parallel_for_dynamic error with [16:48:39] /home/name/tvm/src/tir/transforms/unify_thread_binding.cc:112: Check failed: (ana.CanProveEqual(dom->extent, new_iter_var->dom->extent)) is false: ValueError: All loops that are bound to ``threadIdx.y`` should have the same extent. However, there are two loops with extent T.int64(6) and T.int64(2), which are not equal

What am I doing wrong here? I have not made any modifications to your code, kindly let me know what I am missing here and how I can rectify it.

TIA

Regards, Krishna