You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "mei-ye (via GitHub)" <gi...@apache.org> on 2023/06/02 01:29:01 UTC
[GitHub] [tvm] mei-ye opened a new pull request, #15011: [Vulkan] Add im2col wmma implementation for CHW direct conv2d
mei-ye opened a new pull request, #15011:
URL: https://github.com/apache/tvm/pull/15011
Currently only 16x16x16 FP16 wmma fragments with FP32 intermediates are supported. Add "min_design_space" as a parameter to give minimum design space for meta scheduler tuning. Add "use_int32_const" as a paramter to use int32 type for constants. Allow target query to be called from the schedules so that samplings are constrained to produce legal schedules. Add a unit test test_meta_schedule_vulkan_wmma.py
--
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15011: [Vulkan] Add im2col wmma implementation for CHW direct conv2d
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15011:
URL: https://github.com/apache/tvm/pull/15011#discussion_r1223488559
##########
tests/python/unittest/test_meta_schedule_vulkan_wmma.py:
##########
@@ -0,0 +1,203 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
+# pylint: disable=missing-docstring
+import tvm
+from tvm import te
+import tvm.testing
+import numpy as np
+from tvm import relay
+from tvm.contrib import graph_executor
+from tvm import auto_scheduler
+from tvm import meta_schedule as ms
+import os
+from os import path as osp
+import argparse
+import sys
+from types import MappingProxyType
+import tempfile
+from tvm.tir.tensor_intrin.cuda import (
+ LDMATRIX_16x16_A_INTRIN,
+ LDMATRIX_16x16_B_INTRIN,
+ LDMATRIX_16x16_B_TRANS_INTRIN,
+ LDMATRIX_16x32_A_INTRIN,
+ LDMATRIX_32x16_B_INTRIN,
+ LDMATRIX_16x32_B_TRANS_INTRIN,
+ MMA_f16f16f32_INTRIN,
+ MMA_f16f16f32_TRANS_INTRIN,
+ MMA_f16f16f16_INTRIN,
+ MMA_f16f16f16_TRANS_INTRIN,
+ MMA_i8i8i32_INTRIN,
+ MMA_i8i8i32_TRANS_INTRIN,
+ MMA_fill_16x16_f32_INTRIN,
+ MMA_fill_16x16_f16_INTRIN,
+ MMA_fill_16x16_i32_INTRIN,
+ MMA_store_16x16_f32_global_INTRIN,
+ MMA_store_16x16_f16_global_INTRIN,
+ MMA_store_16x16_i32_global_INTRIN,
+ shared_16x16_to_ldmatrix_32x8_layout,
+ shared_32x16_to_ldmatrix_32x16_layout,
+ shared_16x32_to_ldmatrix_32x16_layout,
+ WMMA_LOAD_16x16x16_F16_A_INTRIN,
+ WMMA_LOAD_16x16x16_F16_B_INTRIN,
+ WMMA_SYNC_16x16x16_f16f16f32_INTRIN,
+ WMMA_FILL_16x16x16_F32_INTRIN,
+ WMMA_STORE_16x16x16_F32_GLOBAL_INTRIN,
+)
+
+
+def _create_tmp_database(tmpdir: str, mod_eq: str = "structural") -> ms.database.JSONDatabase:
+ path_workload = osp.join(tmpdir, "database_workloads.json")
+ path_tuning_record = osp.join(tmpdir, "database_tuning_records.json")
+ return ms.database.JSONDatabase(path_workload, path_tuning_record, module_equality=mod_eq)
+
+
+@tvm.testing.requires_vulkan
+def test_wmma(
Review Comment:
This test should go into `test_meta_schedule_relay_integration.py`.
--
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [PR] [Vulkan] Add im2col wmma implementation for CHW direct conv2d [tvm]
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi closed pull request #15011: [Vulkan] Add im2col wmma implementation for CHW direct conv2d
URL: https://github.com/apache/tvm/pull/15011
--
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on a diff in pull request #15011: [Vulkan] Add im2col wmma implementation for CHW direct conv2d
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on code in PR #15011:
URL: https://github.com/apache/tvm/pull/15011#discussion_r1257571717
##########
python/tvm/topi/cuda/conv2d.py:
##########
@@ -148,3 +154,352 @@ def conv2d_backward_weight_cudnn(
conv_dtype=conv_dtype,
groups=groups,
)
+
+
+@autotvm.register_topi_compute("conv2d_nchw_mma.cuda")
+def conv2d_nchw_mma(cfg, data, kernel, strides, padding, dilation, out_dtype="float32"):
+ """Compute conv2d nchw using im2col"""
+ assert data.dtype == "float16"
+ out_channels, in_channels, kernel_h, kernel_w = get_const_tuple(kernel.shape)
+
+ if isinstance(dilation, int):
+ dilation_h = dilation_w = dilation
+ else:
+ dilation_h, dilation_w = dilation
+ assert dilation_h == 1 and dilation_w == 1
+
+ if isinstance(strides, int):
+ stride_h = stride_w = strides
+ else:
+ stride_h, stride_w = strides
+
+ batch_size, _, P, Q = get_output_shape(
+ data, kernel, stride_h, stride_w, dilation_h, dilation_w, padding
+ )
+ assert batch_size == 1
+
+ pad_top, pad_left, pad_down, pad_right = get_pad_tuple(padding, (kernel_h, kernel_w))
+ pad_before = [0, 0, pad_top, pad_left]
+ pad_after = [0, 0, pad_down, pad_right]
+
+ if all([v == 0 for v in pad_before]) and all([v == 0 for v in pad_after]):
+ pad_data = data
+ else:
+ pad_data = pad(data, pad_before, pad_after, name="pad_data")
+
+ M = out_channels
+ K = in_channels * kernel_h * kernel_w
+ N = batch_size * P * Q
+
+ if kernel_h * kernel_w == 1:
+ ck = te.reduce_axis((0, K), name="k")
+ A = reshape(kernel, (M, K))
+ B = reshape(pad_data, (K, N))
+ C = te.compute(
+ (batch_size, out_channels, P, Q),
+ lambda b, o, h, w: te.sum(
+ A[o, ck].astype(out_dtype) * B[ck, h * Q + w].astype(out_dtype),
+ axis=[ck],
+ ),
+ name="conv2d_nchw_mma",
+ attrs={
+ "schedule_rule": "conv2d_nchw_mma",
+ },
+ )
+ else:
+ # Convert the kernel of (O,I,H,W) to (N,K) format i.e (OC,IC*KH*KW)
+ A = te.compute(
+ (M, K),
+ lambda x, y: kernel[
+ x, (y // (kernel_h * kernel_w)), (y // kernel_w) % kernel_h, y % kernel_w
+ ],
+ name="T_reshape",
+ )
+
+ # Convert the data of (N,C,H,W) to (K,M) format i.e (IC*KH*KW,OH*OW)
+ B = te.compute(
+ (K, N),
+ lambda y, x: pad_data[
+ 0,
+ y // (kernel_h * kernel_w),
+ stride_h * (x // Q) + ((y // kernel_w) % kernel_h),
+ stride_w * (x % Q) + y % kernel_w,
+ ],
+ name="T_reshape_1",
+ )
+
+ # Apply GEMM operation. The result will be of (N,O,H,W) format
+ ck = te.reduce_axis((0, K), name="k")
+ C = te.compute(
+ (batch_size, out_channels, P, Q),
+ lambda b, o, h, w: te.sum(
+ A[o, ck].astype(out_dtype) * B[ck, h * Q + w].astype(out_dtype),
+ axis=[ck],
+ ),
+ name="conv2d_nchw_mma",
+ attrs={
+ "schedule_rule": "conv2d_nchw_mma",
+ },
+ )
+ return C
+
+
+def schedule_rule_conv2d_nchw_mma(sch: Schedule, block: BlockRV):
Review Comment:
Rather than implementing a manual schedule like this, I recommend trying to get auto tensorization for tensor core working for vulkan. I expect the resulting schedule to be very similar.
--
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] mei-ye commented on pull request #15011: [Vulkan] Add im2col wmma implementation for CHW direct conv2d
Posted by "mei-ye (via GitHub)" <gi...@apache.org>.
mei-ye commented on PR #15011:
URL: https://github.com/apache/tvm/pull/15011#issuecomment-1574737609
@tvm-bot rerun
--
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tvm-bot commented on pull request #15011: [Vulkan] Add im2col wmma implementation for CHW direct conv2d
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #15011:
URL: https://github.com/apache/tvm/pull/15011#issuecomment-1573000298
<!---bot-comment-->
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<!--bot-comment-ccs-start-->
* No users to tag found in teams: `vulkan` <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy commented on pull request #15011: [Vulkan] Add im2col wmma implementation for CHW direct conv2d
Posted by "Hzfengsy (via GitHub)" <gi...@apache.org>.
Hzfengsy commented on PR #15011:
URL: https://github.com/apache/tvm/pull/15011#issuecomment-1573017933
cc @junrushao @vinx13 @masahi
--
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.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org