From commits-return-14860-archive-asf-public=cust-asf.ponee.io@tvm.apache.org Tue Jun 9 18:58:14 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 6C27B180663 for ; Tue, 9 Jun 2020 20:58:14 +0200 (CEST) Received: (qmail 80328 invoked by uid 500); 9 Jun 2020 18:58:13 -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 80317 invoked by uid 99); 9 Jun 2020 18:58:13 -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, 09 Jun 2020 18:58:13 +0000 From: =?utf-8?q?GitBox?= To: commits@tvm.apache.org Subject: =?utf-8?q?=5BGitHub=5D_=5Bincubator-tvm=5D_t-vi_commented_on_a_change_in_pul?= =?utf-8?q?l_request_=235752=3A_Make_batch_matrix_multiplication_on_GPU_tuna?= =?utf-8?q?ble?= Message-ID: <159172909372.8471.2210655066950444273.asfpy@gitbox.apache.org> Date: Tue, 09 Jun 2020 18:58:13 -0000 Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: 8bit In-Reply-To: References: t-vi commented on a change in pull request #5752: URL: https://github.com/apache/incubator-tvm/pull/5752#discussion_r437650755 ########## File path: topi/python/topi/cuda/batch_matmul.py ########## @@ -51,55 +62,73 @@ def _schedule(op): C = s.outputs[0].output(0) b, y, x = s[C].op.axis - y_bn = get_max_power2_factor(M, 64) - x_bn = get_max_power2_factor(N, 64) - by, y = s[C].split(y, y_bn) - bx, x = s[C].split(x, x_bn) - y_nthreads = min(y_bn, 8) - x_nthreads = min(x_bn, 8) - ty, yi = s[C].split(y, nparts=y_nthreads) - tx, xi = s[C].split(x, nparts=x_nthreads) - thread_x = te.thread_axis((0, x_nthreads), "threadIdx.x") - thread_y = te.thread_axis((0, y_nthreads), "threadIdx.y") + + cfg.define_split("tile_y", y, num_outputs=3) + cfg.define_split("tile_x", x, num_outputs=3) + cfg.define_knob("auto_unroll_max_step", [8, 16, 32, 64]) + target = tvm.target.Target.current() + if target.target_name in ['nvptx', 'rocm']: + # llvm-based backends cannot do non-explicit unrolling + cfg.define_knob("unroll_explicit", [1]) + else: + cfg.define_knob("unroll_explicit", [0, 1]) + + if cfg.is_fallback: Review comment: I think it should be OK now. ---------------------------------------------------------------- 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