{"action":"created","issue":{"url":"https://api.github.com/repos/dmlc/tvm/issues/2088","repository_url":"https://api.github.com/repos/dmlc/tvm","labels_url":"https://api.github.com/repos/dmlc/tvm/issues/2088/labels{/name}","comments_url":"https://api.github.com/repos/dmlc/tvm/issues/2088/comments","events_url":"https://api.github.com/repos/dmlc/tvm/issues/2088/events","html_url":"https://github.com/dmlc/tvm/issues/2088","id":379576777,"node_id":"MDU6SXNzdWUzNzk1NzY3Nzc=","number":2088,"title":"TOPI LSTM recipe fails to compile due to inserting sync is disallowed inside unnecessary `IfThenElse`","user":{"login":"junrushao1994","id":22515877,"node_id":"MDQ6VXNlcjIyNTE1ODc3","avatar_url":"https://avatars2.githubusercontent.com/u/22515877?v=4","gravatar_id":"","url":"https://api.github.com/users/junrushao1994","html_url":"https://github.com/junrushao1994","followers_url":"https://api.github.com/users/junrushao1994/followers","following_url":"https://api.github.com/users/junrushao1994/following{/other_user}","gists_url":"https://api.github.com/users/junrushao1994/gists{/gist_id}","starred_url":"https://api.github.com/users/junrushao1994/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/junrushao1994/subscriptions","organizations_url":"https://api.github.com/users/junrushao1994/orgs","repos_url":"https://api.github.com/users/junrushao1994/repos","events_url":"https://api.github.com/users/junrushao1994/events{/privacy}","received_events_url":"https://api.github.com/users/junrushao1994/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"assignees":[],"milestone":null,"comments":2,"created_at":"2018-11-11T23:37:15Z","updated_at":"2018-11-18T00:54:33Z","closed_at":null,"author_association":"CONTRIBUTOR","body":"**TL;DR.** TVM doesn't allow syncing threads inside an `IfThenElse`, because it is very prone to deadlock or infinite wait once the conditional diverges. However, some conditionals should be eliminated before doing this check, e.g. `if (0 <= xxxx.idx) { // certainly true thing }`.\r\n\r\n## Details and causes of the error\r\n\r\nThe LSTM recipe in the current master [link](https://github.com/dmlc/tvm/blob/c468558ef9fc0481e70670c0c027c8595d7a0258/topi/recipe/rnn/lstm.py) fails to compile, complaining stuff like:\r\n\r\n```\r\n>>> python $TVM_HOME/topi/recipe/rnn/lstm.py\r\nTraceback (most recent call last):\r\n  File \"topi/recipe/rnn/lstm.py\", line 182, in <module>\r\n    lstm()\r\n  File \"topi/recipe/rnn/lstm.py\", line 179, in lstm\r\n    check_device(\"cuda\")\r\n  File \"topi/recipe/rnn/lstm.py\", line 150, in check_device\r\n    target)\r\n  File \"/home/junrus/Projects/tvm/python/tvm/build_module.py\", line 585, in build\r\n    fhost, mdev = _build_for_device(flist, tar, target_host)\r\n  File \"/home/junrus/Projects/tvm/python/tvm/build_module.py\", line 417, in _build_for_device\r\n    func = ir_pass.ThreadSync(func, \"global\")\r\n  File \"/home/junrus/Projects/tvm/python/tvm/_ffi/_ctypes/function.py\", line 185, in __call__\r\n    ctypes.byref(ret_val), ctypes.byref(ret_tcode)))\r\n  File \"/home/junrus/Projects/tvm/python/tvm/_ffi/base.py\", line 68, in check_call\r\n    raise TVMError(py_str(_LIB.TVMGetLastError()))\r\ntvm._ffi.base.TVMError: [18:29:14] /home/junrus/Projects/tvm/src/pass/storage_sync.cc:76: Check failed: condition_counter() == 0 (1 vs. 0) Cannot insert syncs inside condition\r\n\r\nStack trace returned 10 entries:\r\n[bt] (0) /home/junrus/Projects/tvm/build/libtvm.so(dmlc::StackTrace[abi:cxx11]()+0x1a9) [0x7f40217e5989]\r\n[bt] (1) /home/junrus/Projects/tvm/build/libtvm.so(dmlc::LogMessageFatal::~LogMessageFatal()+0x18) [0x7f40217e6878]\r\n[bt] (2) /home/junrus/Projects/tvm/build/libtvm.so(tvm::ir::ThreadSyncPlanner::Summarize(std::vector<tvm::ir::StorageAccessVisitor::StmtEntry, std::allocator<tvm::ir::StorageAccessVisitor::StmtEntry> >, HalideIR::Internal::For const*)+0x8f3) [0x7f4021a4dce3]\r\n[bt] (3) /home/junrus/Projects/tvm/build/libtvm.so(tvm::ir::StorageAccessVisitor::Visit_(HalideIR::Internal::IfThenElse const*)+0x25d) [0x7f4021a29f1d]\r\n[bt] (4) /home/junrus/Projects/tvm/build/libtvm.so(std::_Function_handler<void (tvm::NodeRef const&, tvm::ir::IRVisitor*), tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>& tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>::set_dispatch<HalideIR::Internal::IfThenElse>(std::function<void (HalideIR::Internal::IfThenElse const*, tvm::ir::IRVisitor*)>)::{lambda(tvm::NodeRef const&, tvm::ir::IRVisitor*)#1}>::_M_invoke(std::_Any_data const&, tvm::NodeRef const&, tvm::ir::IRVisitor*&&)+0x3b) [0x7f40219b932b]\r\n[bt] (5) /home/junrus/Projects/tvm/build/libtvm.so(tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>::operator()(tvm::NodeRef const&, tvm::ir::IRVisitor*) const+0x3f6) [0x7f402183d076]\r\n[bt] (6) /home/junrus/Projects/tvm/build/libtvm.so(std::_Function_handler<void (tvm::NodeRef const&, tvm::ir::IRVisitor*), tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>& tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>::set_dispatch<HalideIR::Internal::ProducerConsumer>(std::function<void (HalideIR::Internal::ProducerConsumer const*, tvm::ir::IRVisitor*)>)::{lambda(tvm::NodeRef const&, tvm::ir::IRVisitor*)#1}>::_M_invoke(std::_Any_data const&, tvm::NodeRef const&, tvm::ir::IRVisitor*&&)+0x3b) [0x7f40219b9c8b]\r\n[bt] (7) /home/junrus/Projects/tvm/build/libtvm.so(tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>::operator()(tvm::NodeRef const&, tvm::ir::IRVisitor*) const+0x3f6) [0x7f402183d076]\r\n[bt] (8) /home/junrus/Projects/tvm/build/libtvm.so(std::_Function_handler<void (tvm::NodeRef const&, tvm::ir::IRVisitor*), tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>& tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>::set_dispatch<HalideIR::Internal::Block>(std::function<void (HalideIR::Internal::Block const*, tvm::ir::IRVisitor*)>)::{lambda(tvm::NodeRef const&, tvm::ir::IRVisitor*)#1}>::_M_invoke(std::_Any_data const&, tvm::NodeRef const&, tvm::ir::IRVisitor*&&)+0x3b) [0x7f40219b9d7b]\r\n[bt] (9) /home/junrus/Projects/tvm/build/libtvm.so(tvm::IRFunctor<void (tvm::NodeRef const&, tvm::ir::IRVisitor*)>::operator()(tvm::NodeRef const&, tvm::ir::IRVisitor*) const+0x3f6) [0x7f402183d076]\r\n```\r\n\r\nI dig into the problem and find the then branch of this line raised the error: ```if ((0 <= lstm_scan.idx))``` in the lowered code below:\r\n\r\n```\r\nproduce lstm_scan {\r\n  // attr [iter_var(blockIdx.x, Range(min=0, extent=24), blockIdx.x)] thread_extent = 24\r\n  // attr [Wh2h.local] storage_scope = \"local\"\r\n  allocate Wh2h.local[float32 * 4 * 1 * 72]\r\n  // attr [placeholder.shared] storage_scope = \"shared\"\r\n  allocate placeholder.shared[float32 * 1 * 1 * 24]\r\n  // attr [s_h2h] storage_scope = \"local\"\r\n  allocate s_h2h[float32 * 1 * 1 * 4 * 1]\r\n  // attr [placeholder.shared] storage_scope = \"shared\"\r\n  allocate placeholder.shared[float32 * 1 * 1 * 576]\r\n  // attr [s_h2h.rf] storage_scope = \"local\"\r\n  allocate s_h2h.rf[float32 * 1 * 1 * 1 * 1 * 1]\r\n  // attr [reduce_temp0] storage_scope = \"local\"\r\n  allocate reduce_temp0[float32 * 1]\r\n  // attr [next_c] storage_scope = \"local\"\r\n  allocate next_c[float32 * 1 * 1 * 1]\r\n  // attr [next_h] storage_scope = \"local\"\r\n  allocate next_h[float32 * 1 * 1 * 1]\r\n  // attr [iter_var(threadIdx.y, Range(min=0, extent=8), threadIdx.y)] thread_extent = 8\r\n  // attr [iter_var(threadIdx.x, Range(min=0, extent=24), threadIdx.x)] thread_extent = 24\r\n  produce Wh2h.local {\r\n    unrolled (ax0, 0, 4) {\r\n      unrolled (ax2, 0, 72) {\r\n        Wh2h.local[((ax0*72) + ax2)] = Wh2h[((((((blockIdx.x*192) + threadIdx.y) + (threadIdx.x*8)) + (ax0*4608))*72) + ax2)]\r\n      }\r\n    }\r\n  }\r\n  lstm_scan.v0[((blockIdx.x*24) + threadIdx.x)] = 0.000000f\r\n  lstm_scan.v1[((blockIdx.x*24) + threadIdx.x)] = 0.000000f\r\n  for (lstm_scan.idx, 0, (num_step + -1)) {\r\n    produce placeholder.shared {\r\n      unrolled (ax2, 0, 24) {\r\n        placeholder.shared[ax2] = lstm_scan.v1[(((blockIdx.x + (lstm_scan.idx*24))*24) + ax2)]\r\n      }\r\n    }\r\n    if ((0 <= lstm_scan.idx)) {\r\n      produce s_h2h {\r\n        for (ax2, 0, 4) {\r\n          produce s_h2h.rf {\r\n            produce placeholder.shared {\r\n              unrolled (ax2.outer, 0, 3) {\r\n                placeholder.shared[(((threadIdx.y*24) + threadIdx.x) + (ax2.outer*192))] = lstm_scan.v0[((((threadIdx.y*24) + threadIdx.x) + (lstm_scan.idx*576)) + (ax2.outer*192))]\r\n              }\r\n            }\r\n            s_h2h.rf[0] = 0.000000f\r\n            unrolled (ki2h.inner, 0, 72) {\r\n              s_h2h.rf[0] = (s_h2h.rf[0] + (placeholder.shared[((threadIdx.y*72) + ki2h.inner)]*Wh2h.local[((ax2*72) + ki2h.inner)]))\r\n            }\r\n          }\r\n          // attr [comm_reducer(result=[(x + y)], lhs=[x], rhs=[y], identity_element=[0.000000f])] reduce_scope = reinterpret((uint64)0)\r\n          tvm_thread_allreduce((uint32)1, s_h2h.rf[0], (uint1)1, reduce_temp0, threadIdx.y)\r\n          s_h2h[ax2] = reduce_temp0[0]\r\n        }\r\n      }\r\n      produce next_c {\r\n        if ((threadIdx.y == 0)) {\r\n          next_c[0] = ((sigmoid((Xi2h[((((blockIdx.x*24) + threadIdx.x) + (lstm_scan.idx*2304)) + 3456)] + s_h2h[2]))*placeholder.shared[threadIdx.x]) + (sigmoid((Xi2h[((((blockIdx.x*24) + threadIdx.x) +\r\n(lstm_scan.idx*2304)) + 2304)] + s_h2h[0]))*tanh((Xi2h[((((blockIdx.x*24) + threadIdx.x) + (lstm_scan.idx*2304)) + 2880)] + s_h2h[1]))))\r\n        }\r\n      }\r\n      produce next_h {\r\n        if ((threadIdx.y == 0)) {\r\n          next_h[0] = (sigmoid((Xi2h[((((blockIdx.x*24) + threadIdx.x) + (lstm_scan.idx*2304)) + 4032)] + s_h2h[3]))*tanh(next_c[0]))\r\n        }\r\n      }\r\n      if ((threadIdx.y == 0)) {\r\n        lstm_scan.v0[((((blockIdx.x*24) + threadIdx.x) + (lstm_scan.idx*576)) + 576)] = next_h[0]\r\n        lstm_scan.v1[((((blockIdx.x*24) + threadIdx.x) + (lstm_scan.idx*576)) + 576)] = next_c[0]\r\n      }\r\n    }\r\n  }\r\n}\r\n```\r\n\r\nI think this conditional should be eliminated before the ThreadSyncPlanner does the check in [this line](https://github.com/dmlc/tvm/blob/master/src/pass/storage_sync.cc#L76).\r\n\r\n## Possible Solution\r\nI am not quite familiar with passes in TVM, but in which pass we could eliminate the conditionals? I think it would be better to be arranged before the `ThreadSync` pass.\r\n"},"comment":{"url":"https://api.github.com/repos/dmlc/tvm/issues/comments/439659006","html_url":"https://github.com/dmlc/tvm/issues/2088#issuecomment-439659006","issue_url":"https://api.github.com/repos/dmlc/tvm/issues/2088","id":439659006,"node_id":"MDEyOklzc3VlQ29tbWVudDQzOTY1OTAwNg==","user":{"login":"junrushao1994","id":22515877,"node_id":"MDQ6VXNlcjIyNTE1ODc3","avatar_url":"https://avatars2.githubusercontent.com/u/22515877?v=4","gravatar_id":"","url":"https://api.github.com/users/junrushao1994","html_url":"https://github.com/junrushao1994","followers_url":"https://api.github.com/users/junrushao1994/followers","following_url":"https://api.github.com/users/junrushao1994/following{/other_user}","gists_url":"https://api.github.com/users/junrushao1994/gists{/gist_id}","starred_url":"https://api.github.com/users/junrushao1994/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/junrushao1994/subscriptions","organizations_url":"https://api.github.com/users/junrushao1994/orgs","repos_url":"https://api.github.com/users/junrushao1994/repos","events_url":"https://api.github.com/users/junrushao1994/events{/privacy}","received_events_url":"https://api.github.com/users/junrushao1994/received_events","type":"User","site_admin":false},"created_at":"2018-11-18T00:54:33Z","updated_at":"2018-11-18T00:54:33Z","author_association":"CONTRIBUTOR","body":"It seems that the boundary checking is introduced by the following code snippet, which I don't quite understand.\r\n\r\nhttps://github.com/dmlc/tvm/blob/c5e1da936a4b583a957c3f325aea0da36aeaf8d3/src/schedule/message_passing.cc#L490-L499\r\n\r\nThe logic here is: (iv be a IterVar, e.g. `lstm_scan.idx`)\r\n1) Let `value = iv - min(iv)`.\r\n2) Let `vmin = min(value)`, and `vmax = max(value)`.\r\n3) Add boundary checks for `value >= 0`  ===> `iv >= min(iv)` if line 494 is true.\r\n4) Add boundary checks for `value + min(iv) < extend(iv)` ===> `iv >= extend(iv)` if line 498 is true.\r\n\r\nThere are 2 things I think weird\r\n1) The first boundary check, i.e. `iv >= min(iv)`, should be always true. It should be the main cause of the issue in this thread. But I found that it was introduced by PR #1091 which fixes Issue #1014. I will spend some time looking into it.\r\n2) The second boundary check, i.e. `iv >= extend(iv)` seems to compare a var to its extent, which does not look very usual.\r\n\r\nPing @tqchen @were "}}
{"actor":{"display_login":"junrushao1994"}}