[GitHub] [incubator-tvm] roastduck opened a new issue #5374: [TE] Warp memory in InferBound

[GitBox <gi...@apache.org>]

roastduck opened a new issue #5374: [TE] Warp memory in InferBound
URL: https://github.com/apache/incubator-tvm/issues/5374
 
 
   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?

----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
 
For queries about this service, please contact Infrastructure at:
users@infra.apache.org


With regards,
Apache Git Services

[GitHub] [incubator-tvm] tqchen commented on issue #5374: [TE] Warp memory in InferBound

[GitBox <gi...@apache.org>]

tqchen commented on issue #5374: [TE] Warp memory in InferBound
URL: https://github.com/apache/incubator-tvm/issues/5374#issuecomment-616012671
 
 
   given that this topic is more like a development discussion, please open a new thread on https://discuss.tvm.ai/

----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
 
For queries about this service, please contact Infrastructure at:
users@infra.apache.org


With regards,
Apache Git Services

[GitHub] [incubator-tvm] tqchen edited a comment on issue #5374: [TE] Warp memory in InferBound

[GitBox <gi...@apache.org>]

tqchen edited a comment on issue #5374: [TE] Warp memory in InferBound
URL: https://github.com/apache/incubator-tvm/issues/5374#issuecomment-616012671
 
 
   thanks @roastduck given that this topic is more like a development discussion, please open a new thread on https://discuss.tvm.ai/

----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
 
For queries about this service, please contact Infrastructure at:
users@infra.apache.org


With regards,
Apache Git Services

[GitHub] [incubator-tvm] tqchen closed issue #5374: [TE] Warp memory in InferBound

[GitBox <gi...@apache.org>]

tqchen closed issue #5374: [TE] Warp memory in InferBound
URL: https://github.com/apache/incubator-tvm/issues/5374
 
 
   

----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
 
For queries about this service, please contact Infrastructure at:
users@infra.apache.org


With regards,
Apache Git Services