From commits-return-12854-archive-asf-public=cust-asf.ponee.io@tvm.apache.org Tue May 5 15:13:30 2020 Return-Path: X-Original-To: archive-asf-public@cust-asf.ponee.io Delivered-To: archive-asf-public@cust-asf.ponee.io Received: from mail.apache.org (hermes.apache.org [207.244.88.153]) by mx-eu-01.ponee.io (Postfix) with SMTP id 2B5AA180626 for ; Tue, 5 May 2020 17:13:30 +0200 (CEST) Received: (qmail 46693 invoked by uid 500); 5 May 2020 15:13:29 -0000 Mailing-List: contact commits-help@tvm.apache.org; run by ezmlm Precedence: bulk List-Help: List-Unsubscribe: List-Post: List-Id: Reply-To: dev@tvm.apache.org Delivered-To: mailing list commits@tvm.apache.org Received: (qmail 46684 invoked by uid 99); 5 May 2020 15:13:29 -0000 Received: from ec2-52-202-80-70.compute-1.amazonaws.com (HELO gitbox.apache.org) (52.202.80.70) by apache.org (qpsmtpd/0.29) with ESMTP; Tue, 05 May 2020 15:13:29 +0000 From: =?utf-8?q?GitBox?= To: commits@tvm.apache.org Subject: =?utf-8?q?=5BGitHub=5D_=5Bincubator-tvm=5D_roastduck_commented_on_a_change_i?= =?utf-8?q?n_pull_request_=235382=3A_=5BTE=5D_Fix_MakeLoopNest_for_warp_memo?= =?utf-8?q?ry?= Message-ID: <158869160948.26397.16948780124000636865.asfpy@gitbox.apache.org> Date: Tue, 05 May 2020 15:13:29 -0000 Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: 8bit References: In-Reply-To: roastduck commented on a change in pull request #5382: URL: https://github.com/apache/incubator-tvm/pull/5382#discussion_r420187744 ########## File path: src/te/operation/op_util.cc ########## @@ -164,9 +164,21 @@ MakeLoopNest(const Stage& stage, value_map[iv] = dom->min; } else { runtime::ThreadScope ts = runtime::ThreadScope::make(bind_iv->thread_tag); - if (stage->scope == "" || stage->scope == "warp" || + if (stage->scope == "" || static_cast(runtime::StorageScope::make(stage->scope).rank) <= ts.rank) { value_map[iv] = var; + } else if (stage->scope == "warp" && ts.rank == 1) { + // To determine whether a thread index is inside or outside a warp, we need + // to know the thread extent. We leave a warning for now. + if (ts.dim_index == 0) { + value_map[iv] = var; + } else { Review comment: Let me explain this piece of code: Line 167: If we are using a local storage, then each iteration variable `iv` accounts to the accessing offset. For example, when accessing a local buffer `a[i, j]`, `i` and `j` are both offsets of `a`. Line 168: There are two cases: 1. If we are using a global storage, then all indices account to the offset. For example, `b[blockIdx.x, threadIdx.x, i]`. 2. If we are using a shared storage, then **variables inside a block**, i.e. iteration variables and thread indices (`<= ts.rank`) account to the offset, but block indices (`> ts.rank`) do not. For example, when accessing a shared buffer `c`, it should be `c[threadIdx.x, i]`, instead of `c[blockIdx.x, threadIdx.x, i]`. Now we come to warp storage. Ideally, all the variables **inside a warp** should account to the offset, but those **outside a warp** do not. It is complex to determine whether a variable is inside a warp, but usually only `threadIdx.x` is inside, so we made the assumption, and give a warning otherwise. ---------------------------------------------------------------- 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