# Question about offset\_factor in decl\_buffer

**URL:** <https://discuss.tvm.apache.org/t/question-about-offset-factor-in-decl-buffer/15857>\
**Category:** Questions\
**Created:** [October 24, 2023, 3:21pm UTC](https://discuss.tvm.apache.org/t/question-about-offset-factor-in-decl-buffer/15857 "2023-10-24T15:21:43Z")\
**Posts on this page:** 2\
**Page:** 1

<div class="post-metadata">

**Author:** ![Necrotos](https://discuss.tvm.apache.org/letter_avatar_proxy/v4/letter/n/f9ae1b/32.png) [@Necrotos](https://discuss.tvm.apache.org/u/Necrotos)\
**Post date:** [October 24, 2023, 3:21pm UTC](https://discuss.tvm.apache.org/t/question-about-offset-factor-in-decl-buffer/15857/1 "2023-10-24T15:21:43Z")

</div>

Hello,

when defining access semantics for a tensor I would expect a dependence on stride and offset. So we would get the correct element by using this formula:

idx = offset + stride\_i \* i + stride\_j \* j (assuming an ixj tensor)

If we look at `decl_buffer` then offset corresponds to elem\_offset and strides to strides. The definition of `decl_buffer` also exposes an offset\_factor that can be defined. How does this fit with into the definition? The examples mention that it is relevant to tensorization, which is also the use case that interests me, but not the exact meaning of the value.

---

<div class="post-metadata">

**Author:** ![sywang](https://discuss.tvm.apache.org/user_avatar/discuss.tvm.apache.org/sywang/32/5106_2.png) [@sywang](https://discuss.tvm.apache.org/u/sywang)\
**Post date:** [October 26, 2023, 3:56am UTC](https://discuss.tvm.apache.org/t/question-about-offset-factor-in-decl-buffer/15857/2 "2023-10-26T03:56:18Z")

</div>

The define of **offset\_factor** is in [tvm.tir — tvm 0.15.dev0 documentation](https://tvm.apache.org/docs/reference/api/python/tir.html#tvm.tir.decl_buffer).

It has two meanings:

1. create a Var for **elem\_offset** ;
2. require that **elem\_offset** be a multiple of **offset\_factor**.

Here gives an example of **wmma\_load\_a** tensorization.

```auto

@T.prim_func
def wmma_load_a_desc(a: T.handle, c: T.handle) -> None:
    A = T.match_buffer(a, (16, 16), "float16", align=128, offset_factor=16, scope="shared")
    C = T.match_buffer(c, (16, 16), "float16", align=128, offset_factor=16, scope="wmma.matrix_a")

    with T.block("root"):
        T.reads(A[0:16, 0:16])
        T.writes(C[0:16, 0:16])
        for i, j in T.grid(16, 16):
            with T.block("load"):
                vii, vjj = T.axis.remap("SS", [i, j])
                C[vii, vjj] = A[vii, vjj]

@T.prim_func
def wmma_load_a_impl(a: T.handle, c: T.handle) -> None:
    s1 = T.int32()
    s0 = T.int32()
    A = T.match_buffer(
        a,
        (16, 16),
        "float16",
        align=128,
        offset_factor=16,
        scope="shared",
        strides=[s1, s0],
    )
    C = T.match_buffer(c, (16, 16), "float16", align=128, offset_factor=16, scope="wmma.matrix_a")

    with T.block("root"):
        T.reads(A[0:16, 0:16])
        T.writes(C[0:16, 0:16])
        T.evaluate(
            T.tvm_load_matrix_sync(
                C.data,
                16,
                16,
                16,
                C.elem_offset // 256 + T.floordiv(T.floormod(C.elem_offset, 256), 16),
                A.access_ptr("r"),
                s1,
                "row_major",
                dtype="handle",
            )
        )

```

We can see the Vars **A.access\_ptr(“r”)** and **C.elem\_offset** are used in **wmma\_load\_a\_impl**.

If we don’t set **offset\_factor** , then the **elem\_offset** will be set to 0 by default. It causes

```auto
TVMError: Check failed: (analyzer_.CanProve(lhs == rhs)) is false: The buffer match constraint for A.elem_offset unmet: 0==threadIdx_y // 4 * 2048 + ax0_0 * 1024 + k_0_1 * 32 + ax1_0 * 16.

```

If we set **offset\_factor = 3** and the **elem\_offset** can’t be divisible by 3. It causes

```auto
TVMError: Check failed: (analyzer_.CanProve(truncmod(buffer->elem_offset, buffer->offset_factor) == 0)) is false: The source elem_offset threadIdx_y // 4 * 2048 + ax0_0 * 1024 + k_0_1 * 32 + ax1_0 * 16 does not satisfy the offset_factor 3.

```