[TVM Discuss] [Development] Warp memory in InferBound

[Tang Shizhi via TVM Discuss <no...@discuss.tvm.ai>]

I'm working with a buffer that bound to warp scope. In `src/te/schedule/message_passing.cc:208`:

```c++
PrimExpr outer = state.at(s->outer);
PrimExpr inner = state.at(s->inner);
PrimExpr factor = dom_map.at(s->inner)->extent;
PrimExpr parent_min = dom_map.at(s->parent)->min;
state[s->parent] = inner + outer * factor;
// add min if they exist
if (!is_zero(parent_min)) {
    state[s->parent] = state[s->parent] + parent_min;
}
```

I found `threadIdx.y` is presented both in `state[s->parent]` and `parent_min` in my application, so the result becomes `threadIdx.y + threadIdx.y + ...`, which leads to a wrong boundary checking in the end.

I tracked down `state[s->parent]`. In `src/te/operation/op_util.cc:167`, there is a code piece that handles different thread indices for different storage scope:

```c++
runtime::ThreadScope ts = runtime::ThreadScope::make(bind_iv->thread_tag);
if (stage->scope == "" || stage->scope == "warp" ||
    static_cast<int>(runtime::StorageScope::make(stage->scope).rank) <= ts.rank) {
    value_map[iv] = var;
} else {
    value_map[iv] = dom->min;
}
```

I think the purpose of the code above is like:

- Both `threadIdx` and `blockIdx` should be indices of a global memory buffer.
- `threadIdx` should be a indies of a shared memory buffer, but `blockIdx` should not.

I think here is a defect on warp memory. `threadIdx.x` (suppose the extent of `threadIdx.x` equals to the warp size) should indeed be a index of a warp buffer, but `threadIdx.y` should not. Currently it seems that both `threadIdx.x` and `threadIdx.y` are counted as indices.

I have not figured out the whole picture yet, and I have not constructed a simple enough counter-example. I think the code piece above is not the only code that handles warp memory in bound inference. Where does `parent_min` decided? And should be consider the situation that extent of `threadIdx.x` < warp size?





---
[Visit Topic](https://discuss.tvm.ai/t/warp-memory-in-inferbound/6421/1) to respond.

You are receiving this because you enabled mailing list mode.

To unsubscribe from these emails, [click here](https://discuss.tvm.ai/email/unsubscribe/5e82e33add8acabe5d836677f3cce0b7d43cadd9351a730235dc8af978ff935c).

[TVM Discuss] [Development] Warp memory in InferBound

[Tang Shizhi via TVM Discuss <no...@discuss.tvm.ai>]

The problem is basically fixed in [PR #5382](https://github.com/apache/incubator-tvm/pull/5382), but I'm still not very clear of the whole picture.





---
[Visit Topic](https://discuss.tvm.ai/t/warp-memory-in-inferbound/6421/2) to respond.

You are receiving this because you enabled mailing list mode.

To unsubscribe from these emails, [click here](https://discuss.tvm.ai/email/unsubscribe/5eab4bf328e3088b3b3e5f8f0cdb2e1dcdb7ff09bb7055dde5806622dbddffc1).