Buffer bind scope mismatch

Questions
#1

I am wandering what does this error mean?

Check failed: arg->scope == value->scope: Argument local.in_buffer Buffer bind scope mismatch

Here I code a matrix multiplication and I want to intrisic to tensorize, as the code follows:

data = tvm.te.placeholder((64, 64), dtype="int8", name="data")
kernel = tvm.te.placeholder((64, 64), dtype="int8", name="kernel")

out_type = "int16"
oshape = (64, 64)

# define the conv2d operator over the padded data
k = te.reduce_axis((0, 64), name="k_w")
hstride, wstride = 1, 1
res = te.compute(
    oshape,
    lambda bo, co: te.sum(
        data[bo, k].astype(out_type) * kernel[k, co].astype(out_type),
        axis=k,
    ),
    name="res",
)

s = tvm.te.create_schedule(res.op)

env = get_env()

x, y= s[res].op.axis

s[res].tensorize(x, gemm(env, 0, 4, 4, 4))
s[res].pragma(x, "import_llvm", intrin_impl())

code = tvm.lower(s, [data, kernel, res], simple_mode=True)
print(code)

While the in intrisic, I declared the buffer, as the code follows:

inp_layout = tvm.tir.decl_buffer(
    inp.shape,
    inp.dtype,
    env.in_scope,
    scope=env.in_scope,
    offset_factor=in_lanes,
    data_alignment=in_lanes,
)

I am wandering what does the error mean? what is arg? what is value? If anyone have any idea, please share with me. I will be very happy to accept any advice. And if you have any question about my code, please let me know.

#2

Have you figured out this question? I have just met the same problem.

#3

sorry, i may forget how i solve the problem

#4

It may be too late for you, but I believe my post would help others. I have encountered the same error and I found out it occurs when the memory scope of the in/out data applied intrinsic to is not matched with the definition of the intrinsic.