You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2021/12/14 05:33:38 UTC
[GitHub] [tvm] masahi opened a new pull request #9737: [CUTLASS] Add conv2d profiler
masahi opened a new pull request #9737:
URL: https://github.com/apache/tvm/pull/9737
Thanks for contributing to TVM! Please refer to guideline https://tvm.apache.org/docs/contribute/ for useful information and tips. After the pull request is submitted, please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @ them in the pull request thread.
--
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 change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r768968554
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
+ """Emit a C++ source for profiling CUTLASS kernels."""
+
+ def __init__(self):
+ from jinja2 import Template
+
+ self.template = Template(
+ """
+#include <iostream>
+#include "cutlass/cutlass.h"
+#include "cutlass/conv/kernel/default_conv2d_fprop.h"
+#include "cutlass/conv/device/implicit_gemm_convolution.h"
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+
+#define CUTLASS_CHECK(status) \
+ { \
+ cutlass::Status error = status; \
+ if (error != cutlass::Status::kSuccess) { \
+ std::cerr << "Got cutlass error: " << cutlassGetStatusString(error) << " at: " << __LINE__ \
+ << std::endl; \
+ exit(EXIT_FAILURE); \
+ } \
+ }
+
+{{OperatorDef}}
+using ImplicitGemm = cutlass::conv::device::ImplicitGemmConvolution<{{OperatorName}}>;
+
+struct Options {
+ cutlass::Tensor4DCoord input_size;
+ cutlass::Tensor4DCoord filter_size;
+ cutlass::Tensor4DCoord padding;
+ cutlass::MatrixCoord conv_stride;
+ cutlass::MatrixCoord dilation;
+
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+ cmd.get_cmd_line_argument("n", input_size.n());
+ cmd.get_cmd_line_argument("h", input_size.h());
+ cmd.get_cmd_line_argument("w", input_size.w());
+ cmd.get_cmd_line_argument("c", input_size.c());
+ cmd.get_cmd_line_argument("k", filter_size.n());
+ cmd.get_cmd_line_argument("r", filter_size.h());
+ cmd.get_cmd_line_argument("s", filter_size.w());
+ int pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w;
+ cmd.get_cmd_line_argument("pad_h", pad_h);
+ cmd.get_cmd_line_argument("pad_w", pad_w);
+ cmd.get_cmd_line_argument("stride_h", stride_h);
+ cmd.get_cmd_line_argument("stride_w", stride_w);
+ cmd.get_cmd_line_argument("dilation_h", dilation_h);
+ cmd.get_cmd_line_argument("dilation_w", dilation_w);
+ filter_size.c() = input_size.c();
+ padding = {pad_h, pad_h, pad_w, pad_w};
+ conv_stride = {stride_h, stride_w};
+ dilation = {dilation_h, dilation_w};
+ }
+
+ cutlass::Tensor4DCoord output_size() const {
+ auto dilated_h = (filter_size.h() - 1) * dilation.row() + 1;
+ auto dilated_w = (filter_size.w() - 1) * dilation.column() + 1;
+ auto h = (input_size.h() + padding.n() + padding.h() - dilated_h) / conv_stride.row() + 1;
+ auto w = (input_size.w() + padding.w() + padding.c() - dilated_w) / conv_stride.column() + 1;
+ return cutlass::Tensor4DCoord(input_size.n(), h, w, filter_size.n());
+ }
+};
+
+double profile_convolution(Options const &options) {
+ using ElementOutput = typename ImplicitGemm::ElementC;
+ using ElementInputA = typename ImplicitGemm::ElementA;
+ using ElementInputB = typename ImplicitGemm::ElementB;
+ auto oshape = options.output_size();
+ cutlass::HostTensor<ElementInputA, typename ImplicitGemm::LayoutA> tensor_a(options.input_size);
+ cutlass::HostTensor<ElementInputB, typename ImplicitGemm::LayoutB> tensor_b(options.filter_size);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_c(oshape);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_ref_c(oshape);
+
+ cutlass::conv::Conv2dProblemSize problem_size(
+ options.input_size,
+ options.filter_size,
+ options.padding,
+ options.conv_stride,
+ options.dilation,
+ options.output_size(),
+ cutlass::conv::Mode::kCrossCorrelation,
+ 1
+ );
+
+ using ElementComputeEpilogue = typename ImplicitGemm::ElementCompute;
+ typename ImplicitGemm::Arguments arguments{
+ problem_size,
+ tensor_a.device_ref(),
+ tensor_b.device_ref(),
+ tensor_c.device_ref(),
+ tensor_c.device_ref(),
+ {ElementComputeEpilogue(1), ElementComputeEpilogue(0)},
+ };
+
+ ImplicitGemm implicit_gemm_op;
+ size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);
+ cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);
+ auto status = implicit_gemm_op.can_implement(arguments);
+ CUTLASS_CHECK(status);
+
+ status = implicit_gemm_op.initialize(arguments, workspace.get());
+ CUTLASS_CHECK(status);
+ status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+
+ cudaEvent_t events[2];
+ for (auto & event : events) {
+ cudaEventCreate(&event);
+ }
+ cudaEventRecord(events[0]);
+
+ for (int iteration = 0; iteration < 100; ++iteration) {
+ auto status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+ }
+
+ cudaEventRecord(events[1]);
+ cudaEventSynchronize(events[1]);
+ float runtime_ms = 0;
+ cudaEventElapsedTime(&runtime_ms, events[0], events[1]);
+
+ for (auto event : events) {
+ (void)cudaEventDestroy(event);
+ }
+ return double(runtime_ms) / 100.0;
Review comment:
We don't invoke the full TVM compilation pipeline with BYOC yet at this point: The goal of these profiler templates are just to select the best implementation given a workload. So we can't make use of the built-in profiler.
We could invoke the TVM compilation for each candidate kernel, run and record the execution time using the built-in profiler. The advantage of the current approach is that we can compile profiler binaries once and cache them to a work directory, so the compilation cost is amortized over different workloads / networks. We could do the similar thing with the TVM compilation approach, but that requires compiling each module with dynamic shapes, and store `*.so` files instead of executables.
Anyway, this is the way GEMM profiler was already written by @Laurawly, so I inherited the same approach.
--
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 change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r769040090
##########
File path: python/tvm/contrib/cutlass/gen_conv2d.py
##########
@@ -121,27 +131,70 @@ def get_default(self, out_dtype):
data_type = gemm_profile_result["data_type"]
return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ def check_align(self, op_name, C, K):
+ """Filter out kernels that cannot be supported."""
+ aligns = re.findall(r"align[1|2|4|8]", op_name)
+ assert len(aligns) == 1
+ align = int(aligns[0][-1])
+ return all([dim % align == 0 for dim in [C, K]])
+
def profile(
- self, d_shape, w_shape, out_shape, out_dtype, profile_all=True, use_multiprocessing=False
+ self,
+ d_shape,
+ w_shape,
+ padding,
+ stride,
+ dilation,
+ out_dtype,
+ profile_all=True,
+ use_multiprocessing=False,
):
"""Profile and select the best kernel from candidate kernels.
If profile_all is False, return immediately after the first applicable kernel is found.
If use_multiprocessing is True, compile all profiler executables in parallel.
"""
- B, _, _, IC = d_shape
+ N, H, W, IC = d_shape
OC, R, S, _ = w_shape
- _, P, Q, _ = out_shape
+ workload = (
+ N,
+ H,
+ W,
+ IC,
+ OC,
+ R,
+ S,
+ padding[0],
+ padding[1],
+ stride[0],
+ stride[1],
+ dilation[0],
+ dilation[1],
+ )
- M = B * P * Q
- N = OC
- K = R * S * IC
+ if workload in self.cache:
+ return self.cache[workload]
- gemm_profile_result = self.gemm_profiler.profile(
- M, N, K, out_dtype, profile_all=profile_all, use_multiprocessing=use_multiprocessing
- )
+ ops = GENERATOR_FUNC_TABLE[self.sm](out_dtype, op_creator=create_conv2d_operator)
+ ops = list(filter(lambda op: self.check_align(op["name"], IC, OC), ops))
- tile_description = gemm_profile_result["tile_description"]
- alignment = gemm_profile_result["alignment"]
- data_type = gemm_profile_result["data_type"]
+ for op in ops:
+ op["runtime"] = -1
- return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ if profile_all:
+ self.engine.compile_all(ops, use_multiprocessing)
+
+ args = (
+ "--n=%d --h=%d --w=%d --c=%d --k=%d --r=%d --s=%d --pad_h=%d --pad_w=%d "
+ "--stride_h=%d --stride_w=%d --dilation_h=%d --dilation_w=%d"
+ ) % workload
+
+ for op in ops:
+ out = self.engine.evaluate(op, args.split(" "))
+ op["runtime"] = out
+ if out > 0 and profile_all is False:
+ break
+
+ valid_ops = filter(lambda op: op["runtime"] > 0, ops)
+ output = sorted(valid_ops, key=lambda i: i["runtime"])
Review comment:
Thanks for the suggestion, also update `gen_gemm.py` (which which this code was copy-pasted)
--
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 pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#issuecomment-993909962
> @masahi Maybe the reason why the TVM script takes so long is that you are doing 100 iterations per benchmark where as the cutlass script is only doing 20?
I believe `cutlass_profiler` is also doing 100 iterations by default: https://github.com/NVIDIA/cutlass/blob/808c25337a3ed4c97ac21895257b1addc72d6ca8/tools/profiler/src/options.cu#L386
> Also the TVM script is running through the whole tvm compilation pipeline for each workload.
As I commented in https://github.com/apache/tvm/pull/9737#discussion_r768968554, we don't invoke the tvm pipeline when we select cutlass kernels.
One major difference with two scripts are that cutlass compiles all kernels into one giant profiler executable, while we generate separate executables for each kernel. So cutlass can allocate / deallocate memory once and loop through each kernel for a given workload to select the best one. Also, I remember that there is a non-trivial initialization cost (close to 1 sec) for any CUDA apps, when we invoke the first CUDA API call - for `cutlass_profiler` this happens only once while we pay that cost for each profiler binaries (there about 60 of them).
But this still doesn't explain 10x difference, so I believe there is something else going on. We could adopt the same approach as `cutlass_profiler` and compile all candidate kernels into one executable. I didn't do that for `conv2d_profiler` because I just followed how `gemm_profiler` is implemented, but that could be a possible improvement.
--
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 change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r768971377
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
+ """Emit a C++ source for profiling CUTLASS kernels."""
+
+ def __init__(self):
+ from jinja2 import Template
+
+ self.template = Template(
+ """
+#include <iostream>
+#include "cutlass/cutlass.h"
+#include "cutlass/conv/kernel/default_conv2d_fprop.h"
+#include "cutlass/conv/device/implicit_gemm_convolution.h"
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+
+#define CUTLASS_CHECK(status) \
+ { \
+ cutlass::Status error = status; \
+ if (error != cutlass::Status::kSuccess) { \
+ std::cerr << "Got cutlass error: " << cutlassGetStatusString(error) << " at: " << __LINE__ \
+ << std::endl; \
+ exit(EXIT_FAILURE); \
+ } \
+ }
+
+{{OperatorDef}}
+using ImplicitGemm = cutlass::conv::device::ImplicitGemmConvolution<{{OperatorName}}>;
+
+struct Options {
+ cutlass::Tensor4DCoord input_size;
+ cutlass::Tensor4DCoord filter_size;
+ cutlass::Tensor4DCoord padding;
+ cutlass::MatrixCoord conv_stride;
+ cutlass::MatrixCoord dilation;
+
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+ cmd.get_cmd_line_argument("n", input_size.n());
+ cmd.get_cmd_line_argument("h", input_size.h());
+ cmd.get_cmd_line_argument("w", input_size.w());
+ cmd.get_cmd_line_argument("c", input_size.c());
+ cmd.get_cmd_line_argument("k", filter_size.n());
+ cmd.get_cmd_line_argument("r", filter_size.h());
+ cmd.get_cmd_line_argument("s", filter_size.w());
+ int pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w;
+ cmd.get_cmd_line_argument("pad_h", pad_h);
+ cmd.get_cmd_line_argument("pad_w", pad_w);
+ cmd.get_cmd_line_argument("stride_h", stride_h);
+ cmd.get_cmd_line_argument("stride_w", stride_w);
+ cmd.get_cmd_line_argument("dilation_h", dilation_h);
+ cmd.get_cmd_line_argument("dilation_w", dilation_w);
+ filter_size.c() = input_size.c();
+ padding = {pad_h, pad_h, pad_w, pad_w};
+ conv_stride = {stride_h, stride_w};
+ dilation = {dilation_h, dilation_w};
+ }
+
+ cutlass::Tensor4DCoord output_size() const {
+ auto dilated_h = (filter_size.h() - 1) * dilation.row() + 1;
+ auto dilated_w = (filter_size.w() - 1) * dilation.column() + 1;
+ auto h = (input_size.h() + padding.n() + padding.h() - dilated_h) / conv_stride.row() + 1;
+ auto w = (input_size.w() + padding.w() + padding.c() - dilated_w) / conv_stride.column() + 1;
+ return cutlass::Tensor4DCoord(input_size.n(), h, w, filter_size.n());
+ }
+};
+
+double profile_convolution(Options const &options) {
+ using ElementOutput = typename ImplicitGemm::ElementC;
+ using ElementInputA = typename ImplicitGemm::ElementA;
+ using ElementInputB = typename ImplicitGemm::ElementB;
+ auto oshape = options.output_size();
+ cutlass::HostTensor<ElementInputA, typename ImplicitGemm::LayoutA> tensor_a(options.input_size);
+ cutlass::HostTensor<ElementInputB, typename ImplicitGemm::LayoutB> tensor_b(options.filter_size);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_c(oshape);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_ref_c(oshape);
+
+ cutlass::conv::Conv2dProblemSize problem_size(
+ options.input_size,
+ options.filter_size,
+ options.padding,
+ options.conv_stride,
+ options.dilation,
+ options.output_size(),
+ cutlass::conv::Mode::kCrossCorrelation,
+ 1
+ );
+
+ using ElementComputeEpilogue = typename ImplicitGemm::ElementCompute;
+ typename ImplicitGemm::Arguments arguments{
+ problem_size,
+ tensor_a.device_ref(),
+ tensor_b.device_ref(),
+ tensor_c.device_ref(),
+ tensor_c.device_ref(),
+ {ElementComputeEpilogue(1), ElementComputeEpilogue(0)},
+ };
+
+ ImplicitGemm implicit_gemm_op;
+ size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);
+ cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);
+ auto status = implicit_gemm_op.can_implement(arguments);
+ CUTLASS_CHECK(status);
+
+ status = implicit_gemm_op.initialize(arguments, workspace.get());
+ CUTLASS_CHECK(status);
+ status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+
+ cudaEvent_t events[2];
+ for (auto & event : events) {
+ cudaEventCreate(&event);
+ }
+ cudaEventRecord(events[0]);
Review comment:
Anyway, I want to defer addressing this problem in future PR.
--
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 merged pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi merged pull request #9737:
URL: https://github.com/apache/tvm/pull/9737
--
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 change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r768960030
##########
File path: python/tvm/relay/op/strategy/cuda.py
##########
@@ -324,7 +324,10 @@ def conv2d_strategy_cuda(attrs, inputs, out_type, target):
plevel=25,
)
- elif is_depthwise_conv2d(data.shape, layout, kernel.shape, kernel_layout, groups):
+ elif (
+ is_depthwise_conv2d(data.shape, layout, kernel.shape, kernel_layout, groups)
+ and "cudnn" not in target.libs
Review comment:
cuDNN requires a different kernel layout than AutoTVM when the input layout is NHWC, in which case two implementations are not compatible.
But there is no problem with the NCHW layout, so I've refined the condition to
```
elif is_depthwise_conv2d(data.shape, layout, kernel.shape, kernel_layout, groups) and (
layout == "NCHW" or "cudnn" not in target.libs):
```
--
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] tkonolige commented on pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
tkonolige commented on pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#issuecomment-993839104
@masahi Maybe the reason why the TVM script takes so long is that you are doing 100 iterations per benchmark where as the cutlass script is only doing 20? Also the TVM script is running through the whole tvm compilation pipeline for each workload.
--
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 edited a comment on pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi edited a comment on pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#issuecomment-993909962
> @masahi Maybe the reason why the TVM script takes so long is that you are doing 100 iterations per benchmark where as the cutlass script is only doing 20?
I believe `cutlass_profiler` is also doing 100 iterations by default: https://github.com/NVIDIA/cutlass/blob/808c25337a3ed4c97ac21895257b1addc72d6ca8/tools/profiler/src/options.cu#L386
> Also the TVM script is running through the whole tvm compilation pipeline for each workload.
As I commented in https://github.com/apache/tvm/pull/9737#discussion_r768968554, we don't invoke the tvm pipeline when we select cutlass kernels.
One major difference with two scripts are that cutlass compiles all kernels into one giant profiler executable, while we generate separate executables for each kernel. So cutlass can allocate / deallocate memory once and loop through each kernel for a given workload to select the best one. Also, I remember that there is a non-trivial initialization cost (close to 1 sec) for any CUDA apps, when we invoke the first CUDA API call - for `cutlass_profiler` this happens only once while we pay that cost for each profiler binary (there about 60 of them).
But this still doesn't explain 10x difference, so I believe there is something else going on. We could adopt the same approach as `cutlass_profiler` and compile all candidate kernels into one executable. I didn't do that for `conv2d_profiler` because I just followed how `gemm_profiler` is implemented, but that could be a possible improvement.
--
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] jroesch commented on pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
jroesch commented on pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#issuecomment-993326662
cc @tkonolige
--
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 change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r768970900
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
+ """Emit a C++ source for profiling CUTLASS kernels."""
+
+ def __init__(self):
+ from jinja2 import Template
+
+ self.template = Template(
+ """
+#include <iostream>
+#include "cutlass/cutlass.h"
+#include "cutlass/conv/kernel/default_conv2d_fprop.h"
+#include "cutlass/conv/device/implicit_gemm_convolution.h"
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+
+#define CUTLASS_CHECK(status) \
+ { \
+ cutlass::Status error = status; \
+ if (error != cutlass::Status::kSuccess) { \
+ std::cerr << "Got cutlass error: " << cutlassGetStatusString(error) << " at: " << __LINE__ \
+ << std::endl; \
+ exit(EXIT_FAILURE); \
+ } \
+ }
+
+{{OperatorDef}}
+using ImplicitGemm = cutlass::conv::device::ImplicitGemmConvolution<{{OperatorName}}>;
+
+struct Options {
+ cutlass::Tensor4DCoord input_size;
+ cutlass::Tensor4DCoord filter_size;
+ cutlass::Tensor4DCoord padding;
+ cutlass::MatrixCoord conv_stride;
+ cutlass::MatrixCoord dilation;
+
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+ cmd.get_cmd_line_argument("n", input_size.n());
+ cmd.get_cmd_line_argument("h", input_size.h());
+ cmd.get_cmd_line_argument("w", input_size.w());
+ cmd.get_cmd_line_argument("c", input_size.c());
+ cmd.get_cmd_line_argument("k", filter_size.n());
+ cmd.get_cmd_line_argument("r", filter_size.h());
+ cmd.get_cmd_line_argument("s", filter_size.w());
+ int pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w;
+ cmd.get_cmd_line_argument("pad_h", pad_h);
+ cmd.get_cmd_line_argument("pad_w", pad_w);
+ cmd.get_cmd_line_argument("stride_h", stride_h);
+ cmd.get_cmd_line_argument("stride_w", stride_w);
+ cmd.get_cmd_line_argument("dilation_h", dilation_h);
+ cmd.get_cmd_line_argument("dilation_w", dilation_w);
+ filter_size.c() = input_size.c();
+ padding = {pad_h, pad_h, pad_w, pad_w};
+ conv_stride = {stride_h, stride_w};
+ dilation = {dilation_h, dilation_w};
+ }
+
+ cutlass::Tensor4DCoord output_size() const {
+ auto dilated_h = (filter_size.h() - 1) * dilation.row() + 1;
+ auto dilated_w = (filter_size.w() - 1) * dilation.column() + 1;
+ auto h = (input_size.h() + padding.n() + padding.h() - dilated_h) / conv_stride.row() + 1;
+ auto w = (input_size.w() + padding.w() + padding.c() - dilated_w) / conv_stride.column() + 1;
+ return cutlass::Tensor4DCoord(input_size.n(), h, w, filter_size.n());
+ }
+};
+
+double profile_convolution(Options const &options) {
+ using ElementOutput = typename ImplicitGemm::ElementC;
+ using ElementInputA = typename ImplicitGemm::ElementA;
+ using ElementInputB = typename ImplicitGemm::ElementB;
+ auto oshape = options.output_size();
+ cutlass::HostTensor<ElementInputA, typename ImplicitGemm::LayoutA> tensor_a(options.input_size);
+ cutlass::HostTensor<ElementInputB, typename ImplicitGemm::LayoutB> tensor_b(options.filter_size);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_c(oshape);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_ref_c(oshape);
+
+ cutlass::conv::Conv2dProblemSize problem_size(
+ options.input_size,
+ options.filter_size,
+ options.padding,
+ options.conv_stride,
+ options.dilation,
+ options.output_size(),
+ cutlass::conv::Mode::kCrossCorrelation,
+ 1
+ );
+
+ using ElementComputeEpilogue = typename ImplicitGemm::ElementCompute;
+ typename ImplicitGemm::Arguments arguments{
+ problem_size,
+ tensor_a.device_ref(),
+ tensor_b.device_ref(),
+ tensor_c.device_ref(),
+ tensor_c.device_ref(),
+ {ElementComputeEpilogue(1), ElementComputeEpilogue(0)},
+ };
+
+ ImplicitGemm implicit_gemm_op;
+ size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);
+ cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);
+ auto status = implicit_gemm_op.can_implement(arguments);
+ CUTLASS_CHECK(status);
+
+ status = implicit_gemm_op.initialize(arguments, workspace.get());
+ CUTLASS_CHECK(status);
+ status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+
+ cudaEvent_t events[2];
+ for (auto & event : events) {
+ cudaEventCreate(&event);
+ }
+ cudaEventRecord(events[0]);
Review comment:
Yeah, I noticed that too and we should use cuda events there as well. I didn't look deeply into `gemm_profiler` or compare the selected gemm kernels with `cutlass_profiler` like I did with the covn2d profiler in this PR. Any comment @Laurawly?
--
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 change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r769075265
##########
File path: python/tvm/contrib/cutlass/gen_conv2d.py
##########
@@ -121,27 +131,67 @@ def get_default(self, out_dtype):
data_type = gemm_profile_result["data_type"]
return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ def check_align(self, op_name, C, K):
+ """Filter out kernels that cannot be supported."""
+ aligns = re.findall(r"align[1|2|4|8]", op_name)
+ assert len(aligns) == 1
+ align = int(aligns[0][-1])
+ return all([dim % align == 0 for dim in [C, K]])
+
def profile(
- self, d_shape, w_shape, out_shape, out_dtype, profile_all=True, use_multiprocessing=False
+ self,
+ d_shape,
+ w_shape,
+ padding,
+ stride,
+ dilation,
+ out_dtype,
+ profile_all=True,
+ use_multiprocessing=False,
):
"""Profile and select the best kernel from candidate kernels.
If profile_all is False, return immediately after the first applicable kernel is found.
If use_multiprocessing is True, compile all profiler executables in parallel.
"""
- B, _, _, IC = d_shape
+ N, H, W, IC = d_shape
OC, R, S, _ = w_shape
- _, P, Q, _ = out_shape
+ workload = (
+ N,
+ H,
+ W,
+ IC,
+ OC,
+ R,
+ S,
+ padding[0],
+ padding[1],
+ stride[0],
+ stride[1],
+ dilation[0],
+ dilation[1],
+ )
- M = B * P * Q
- N = OC
- K = R * S * IC
+ if workload in self.cache:
+ return self.cache[workload]
- gemm_profile_result = self.gemm_profiler.profile(
- M, N, K, out_dtype, profile_all=profile_all, use_multiprocessing=use_multiprocessing
- )
+ ops = GENERATOR_FUNC_TABLE[self.sm](out_dtype, op_creator=create_conv2d_operator)
+ ops = list(filter(lambda op: self.check_align(op["name"], IC, OC), ops))
- tile_description = gemm_profile_result["tile_description"]
- alignment = gemm_profile_result["alignment"]
- data_type = gemm_profile_result["data_type"]
+ if profile_all:
+ self.engine.compile_all(ops, use_multiprocessing)
- return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ args = (
+ "--n=%d --h=%d --w=%d --c=%d --k=%d --r=%d --s=%d --pad_h=%d --pad_w=%d "
+ "--stride_h=%d --stride_w=%d --dilation_h=%d --dilation_w=%d"
+ ) % workload
+
+ for op in ops:
+ out = self.engine.evaluate(op, args.split(" "))
+ op["runtime"] = out
+ if out > 0 and not profile_all:
Review comment:
oops you are right, changed to `out < float("inf")`
--
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 change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r768961263
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
Review comment:
Yes, but there is not much to share between conv2d and gemm profilers. For now, these twos ops are the only ones we want to support.
--
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] comaniac commented on a change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
comaniac commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r769071320
##########
File path: python/tvm/contrib/cutlass/gen_conv2d.py
##########
@@ -121,27 +131,67 @@ def get_default(self, out_dtype):
data_type = gemm_profile_result["data_type"]
return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ def check_align(self, op_name, C, K):
+ """Filter out kernels that cannot be supported."""
+ aligns = re.findall(r"align[1|2|4|8]", op_name)
+ assert len(aligns) == 1
+ align = int(aligns[0][-1])
+ return all([dim % align == 0 for dim in [C, K]])
+
def profile(
- self, d_shape, w_shape, out_shape, out_dtype, profile_all=True, use_multiprocessing=False
+ self,
+ d_shape,
+ w_shape,
+ padding,
+ stride,
+ dilation,
+ out_dtype,
+ profile_all=True,
+ use_multiprocessing=False,
):
"""Profile and select the best kernel from candidate kernels.
If profile_all is False, return immediately after the first applicable kernel is found.
If use_multiprocessing is True, compile all profiler executables in parallel.
"""
- B, _, _, IC = d_shape
+ N, H, W, IC = d_shape
OC, R, S, _ = w_shape
- _, P, Q, _ = out_shape
+ workload = (
+ N,
+ H,
+ W,
+ IC,
+ OC,
+ R,
+ S,
+ padding[0],
+ padding[1],
+ stride[0],
+ stride[1],
+ dilation[0],
+ dilation[1],
+ )
- M = B * P * Q
- N = OC
- K = R * S * IC
+ if workload in self.cache:
+ return self.cache[workload]
- gemm_profile_result = self.gemm_profiler.profile(
- M, N, K, out_dtype, profile_all=profile_all, use_multiprocessing=use_multiprocessing
- )
+ ops = GENERATOR_FUNC_TABLE[self.sm](out_dtype, op_creator=create_conv2d_operator)
+ ops = list(filter(lambda op: self.check_align(op["name"], IC, OC), ops))
- tile_description = gemm_profile_result["tile_description"]
- alignment = gemm_profile_result["alignment"]
- data_type = gemm_profile_result["data_type"]
+ if profile_all:
+ self.engine.compile_all(ops, use_multiprocessing)
- return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ args = (
+ "--n=%d --h=%d --w=%d --c=%d --k=%d --r=%d --s=%d --pad_h=%d --pad_w=%d "
+ "--stride_h=%d --stride_w=%d --dilation_h=%d --dilation_w=%d"
+ ) % workload
+
+ for op in ops:
+ out = self.engine.evaluate(op, args.split(" "))
+ op["runtime"] = out
+ if out > 0 and not profile_all:
Review comment:
IIUC, now you changed `evaluate` to return `float("inf")` when invalid. Then the fist invalid kernel will be selected since `float("inf") > 0`, right?
--
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] tkonolige commented on a change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r768999703
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
+ """Emit a C++ source for profiling CUTLASS kernels."""
+
+ def __init__(self):
+ from jinja2 import Template
+
+ self.template = Template(
+ """
+#include <iostream>
+#include "cutlass/cutlass.h"
+#include "cutlass/conv/kernel/default_conv2d_fprop.h"
+#include "cutlass/conv/device/implicit_gemm_convolution.h"
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+
+#define CUTLASS_CHECK(status) \
+ { \
+ cutlass::Status error = status; \
+ if (error != cutlass::Status::kSuccess) { \
+ std::cerr << "Got cutlass error: " << cutlassGetStatusString(error) << " at: " << __LINE__ \
+ << std::endl; \
+ exit(EXIT_FAILURE); \
+ } \
+ }
+
+{{OperatorDef}}
+using ImplicitGemm = cutlass::conv::device::ImplicitGemmConvolution<{{OperatorName}}>;
+
+struct Options {
+ cutlass::Tensor4DCoord input_size;
+ cutlass::Tensor4DCoord filter_size;
+ cutlass::Tensor4DCoord padding;
+ cutlass::MatrixCoord conv_stride;
+ cutlass::MatrixCoord dilation;
+
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+ cmd.get_cmd_line_argument("n", input_size.n());
+ cmd.get_cmd_line_argument("h", input_size.h());
+ cmd.get_cmd_line_argument("w", input_size.w());
+ cmd.get_cmd_line_argument("c", input_size.c());
+ cmd.get_cmd_line_argument("k", filter_size.n());
+ cmd.get_cmd_line_argument("r", filter_size.h());
+ cmd.get_cmd_line_argument("s", filter_size.w());
+ int pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w;
+ cmd.get_cmd_line_argument("pad_h", pad_h);
+ cmd.get_cmd_line_argument("pad_w", pad_w);
+ cmd.get_cmd_line_argument("stride_h", stride_h);
+ cmd.get_cmd_line_argument("stride_w", stride_w);
+ cmd.get_cmd_line_argument("dilation_h", dilation_h);
+ cmd.get_cmd_line_argument("dilation_w", dilation_w);
+ filter_size.c() = input_size.c();
+ padding = {pad_h, pad_h, pad_w, pad_w};
+ conv_stride = {stride_h, stride_w};
+ dilation = {dilation_h, dilation_w};
+ }
+
+ cutlass::Tensor4DCoord output_size() const {
+ auto dilated_h = (filter_size.h() - 1) * dilation.row() + 1;
+ auto dilated_w = (filter_size.w() - 1) * dilation.column() + 1;
+ auto h = (input_size.h() + padding.n() + padding.h() - dilated_h) / conv_stride.row() + 1;
+ auto w = (input_size.w() + padding.w() + padding.c() - dilated_w) / conv_stride.column() + 1;
+ return cutlass::Tensor4DCoord(input_size.n(), h, w, filter_size.n());
+ }
+};
+
+double profile_convolution(Options const &options) {
+ using ElementOutput = typename ImplicitGemm::ElementC;
+ using ElementInputA = typename ImplicitGemm::ElementA;
+ using ElementInputB = typename ImplicitGemm::ElementB;
+ auto oshape = options.output_size();
+ cutlass::HostTensor<ElementInputA, typename ImplicitGemm::LayoutA> tensor_a(options.input_size);
+ cutlass::HostTensor<ElementInputB, typename ImplicitGemm::LayoutB> tensor_b(options.filter_size);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_c(oshape);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_ref_c(oshape);
+
+ cutlass::conv::Conv2dProblemSize problem_size(
+ options.input_size,
+ options.filter_size,
+ options.padding,
+ options.conv_stride,
+ options.dilation,
+ options.output_size(),
+ cutlass::conv::Mode::kCrossCorrelation,
+ 1
+ );
+
+ using ElementComputeEpilogue = typename ImplicitGemm::ElementCompute;
+ typename ImplicitGemm::Arguments arguments{
+ problem_size,
+ tensor_a.device_ref(),
+ tensor_b.device_ref(),
+ tensor_c.device_ref(),
+ tensor_c.device_ref(),
+ {ElementComputeEpilogue(1), ElementComputeEpilogue(0)},
+ };
+
+ ImplicitGemm implicit_gemm_op;
+ size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);
+ cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);
+ auto status = implicit_gemm_op.can_implement(arguments);
+ CUTLASS_CHECK(status);
+
+ status = implicit_gemm_op.initialize(arguments, workspace.get());
+ CUTLASS_CHECK(status);
+ status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+
+ cudaEvent_t events[2];
+ for (auto & event : events) {
+ cudaEventCreate(&event);
+ }
+ cudaEventRecord(events[0]);
+
+ for (int iteration = 0; iteration < 100; ++iteration) {
+ auto status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+ }
+
+ cudaEventRecord(events[1]);
+ cudaEventSynchronize(events[1]);
+ float runtime_ms = 0;
+ cudaEventElapsedTime(&runtime_ms, events[0], events[1]);
+
+ for (auto event : events) {
+ (void)cudaEventDestroy(event);
+ }
+ return double(runtime_ms) / 100.0;
Review comment:
You don't need to need to invoke the whole compiler pipeline to use the timing code. It will accept any packedfunc. Here the packedfunc would just run the kernel.
If we only are every going to conv2d and gemm, then it probably doesn't matter.
--
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] tkonolige commented on a change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
tkonolige commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r768913178
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
Review comment:
Do we have to write a new class like this for each op we want to support? If so, we might want to think about making something more reusable.
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
+ """Emit a C++ source for profiling CUTLASS kernels."""
+
+ def __init__(self):
+ from jinja2 import Template
+
+ self.template = Template(
+ """
+#include <iostream>
+#include "cutlass/cutlass.h"
+#include "cutlass/conv/kernel/default_conv2d_fprop.h"
+#include "cutlass/conv/device/implicit_gemm_convolution.h"
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+
+#define CUTLASS_CHECK(status) \
+ { \
+ cutlass::Status error = status; \
+ if (error != cutlass::Status::kSuccess) { \
+ std::cerr << "Got cutlass error: " << cutlassGetStatusString(error) << " at: " << __LINE__ \
+ << std::endl; \
+ exit(EXIT_FAILURE); \
+ } \
+ }
+
+{{OperatorDef}}
+using ImplicitGemm = cutlass::conv::device::ImplicitGemmConvolution<{{OperatorName}}>;
+
+struct Options {
+ cutlass::Tensor4DCoord input_size;
+ cutlass::Tensor4DCoord filter_size;
+ cutlass::Tensor4DCoord padding;
+ cutlass::MatrixCoord conv_stride;
+ cutlass::MatrixCoord dilation;
+
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+ cmd.get_cmd_line_argument("n", input_size.n());
+ cmd.get_cmd_line_argument("h", input_size.h());
+ cmd.get_cmd_line_argument("w", input_size.w());
+ cmd.get_cmd_line_argument("c", input_size.c());
+ cmd.get_cmd_line_argument("k", filter_size.n());
+ cmd.get_cmd_line_argument("r", filter_size.h());
+ cmd.get_cmd_line_argument("s", filter_size.w());
+ int pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w;
+ cmd.get_cmd_line_argument("pad_h", pad_h);
+ cmd.get_cmd_line_argument("pad_w", pad_w);
+ cmd.get_cmd_line_argument("stride_h", stride_h);
+ cmd.get_cmd_line_argument("stride_w", stride_w);
+ cmd.get_cmd_line_argument("dilation_h", dilation_h);
+ cmd.get_cmd_line_argument("dilation_w", dilation_w);
+ filter_size.c() = input_size.c();
+ padding = {pad_h, pad_h, pad_w, pad_w};
+ conv_stride = {stride_h, stride_w};
+ dilation = {dilation_h, dilation_w};
+ }
+
+ cutlass::Tensor4DCoord output_size() const {
+ auto dilated_h = (filter_size.h() - 1) * dilation.row() + 1;
+ auto dilated_w = (filter_size.w() - 1) * dilation.column() + 1;
+ auto h = (input_size.h() + padding.n() + padding.h() - dilated_h) / conv_stride.row() + 1;
+ auto w = (input_size.w() + padding.w() + padding.c() - dilated_w) / conv_stride.column() + 1;
+ return cutlass::Tensor4DCoord(input_size.n(), h, w, filter_size.n());
+ }
+};
+
+double profile_convolution(Options const &options) {
+ using ElementOutput = typename ImplicitGemm::ElementC;
+ using ElementInputA = typename ImplicitGemm::ElementA;
+ using ElementInputB = typename ImplicitGemm::ElementB;
+ auto oshape = options.output_size();
+ cutlass::HostTensor<ElementInputA, typename ImplicitGemm::LayoutA> tensor_a(options.input_size);
+ cutlass::HostTensor<ElementInputB, typename ImplicitGemm::LayoutB> tensor_b(options.filter_size);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_c(oshape);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_ref_c(oshape);
+
+ cutlass::conv::Conv2dProblemSize problem_size(
+ options.input_size,
+ options.filter_size,
+ options.padding,
+ options.conv_stride,
+ options.dilation,
+ options.output_size(),
+ cutlass::conv::Mode::kCrossCorrelation,
+ 1
+ );
+
+ using ElementComputeEpilogue = typename ImplicitGemm::ElementCompute;
+ typename ImplicitGemm::Arguments arguments{
+ problem_size,
+ tensor_a.device_ref(),
+ tensor_b.device_ref(),
+ tensor_c.device_ref(),
+ tensor_c.device_ref(),
+ {ElementComputeEpilogue(1), ElementComputeEpilogue(0)},
+ };
+
+ ImplicitGemm implicit_gemm_op;
+ size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);
+ cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);
+ auto status = implicit_gemm_op.can_implement(arguments);
+ CUTLASS_CHECK(status);
+
+ status = implicit_gemm_op.initialize(arguments, workspace.get());
+ CUTLASS_CHECK(status);
+ status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+
+ cudaEvent_t events[2];
+ for (auto & event : events) {
+ cudaEventCreate(&event);
+ }
+ cudaEventRecord(events[0]);
Review comment:
Using `cudaEventRecord` is good, but I notice the gemm_profiler does not do that. Maybe we should fix it?
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
+ """Emit a C++ source for profiling CUTLASS kernels."""
+
+ def __init__(self):
+ from jinja2 import Template
+
+ self.template = Template(
+ """
+#include <iostream>
+#include "cutlass/cutlass.h"
+#include "cutlass/conv/kernel/default_conv2d_fprop.h"
+#include "cutlass/conv/device/implicit_gemm_convolution.h"
+#include "cutlass/util/command_line.h"
+#include "cutlass/util/host_tensor.h"
+#include "cutlass/util/reference/host/tensor_fill.h"
+
+#define CUTLASS_CHECK(status) \
+ { \
+ cutlass::Status error = status; \
+ if (error != cutlass::Status::kSuccess) { \
+ std::cerr << "Got cutlass error: " << cutlassGetStatusString(error) << " at: " << __LINE__ \
+ << std::endl; \
+ exit(EXIT_FAILURE); \
+ } \
+ }
+
+{{OperatorDef}}
+using ImplicitGemm = cutlass::conv::device::ImplicitGemmConvolution<{{OperatorName}}>;
+
+struct Options {
+ cutlass::Tensor4DCoord input_size;
+ cutlass::Tensor4DCoord filter_size;
+ cutlass::Tensor4DCoord padding;
+ cutlass::MatrixCoord conv_stride;
+ cutlass::MatrixCoord dilation;
+
+ void parse(int argc, char const **args) {
+ cutlass::CommandLine cmd(argc, args);
+ cmd.get_cmd_line_argument("n", input_size.n());
+ cmd.get_cmd_line_argument("h", input_size.h());
+ cmd.get_cmd_line_argument("w", input_size.w());
+ cmd.get_cmd_line_argument("c", input_size.c());
+ cmd.get_cmd_line_argument("k", filter_size.n());
+ cmd.get_cmd_line_argument("r", filter_size.h());
+ cmd.get_cmd_line_argument("s", filter_size.w());
+ int pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w;
+ cmd.get_cmd_line_argument("pad_h", pad_h);
+ cmd.get_cmd_line_argument("pad_w", pad_w);
+ cmd.get_cmd_line_argument("stride_h", stride_h);
+ cmd.get_cmd_line_argument("stride_w", stride_w);
+ cmd.get_cmd_line_argument("dilation_h", dilation_h);
+ cmd.get_cmd_line_argument("dilation_w", dilation_w);
+ filter_size.c() = input_size.c();
+ padding = {pad_h, pad_h, pad_w, pad_w};
+ conv_stride = {stride_h, stride_w};
+ dilation = {dilation_h, dilation_w};
+ }
+
+ cutlass::Tensor4DCoord output_size() const {
+ auto dilated_h = (filter_size.h() - 1) * dilation.row() + 1;
+ auto dilated_w = (filter_size.w() - 1) * dilation.column() + 1;
+ auto h = (input_size.h() + padding.n() + padding.h() - dilated_h) / conv_stride.row() + 1;
+ auto w = (input_size.w() + padding.w() + padding.c() - dilated_w) / conv_stride.column() + 1;
+ return cutlass::Tensor4DCoord(input_size.n(), h, w, filter_size.n());
+ }
+};
+
+double profile_convolution(Options const &options) {
+ using ElementOutput = typename ImplicitGemm::ElementC;
+ using ElementInputA = typename ImplicitGemm::ElementA;
+ using ElementInputB = typename ImplicitGemm::ElementB;
+ auto oshape = options.output_size();
+ cutlass::HostTensor<ElementInputA, typename ImplicitGemm::LayoutA> tensor_a(options.input_size);
+ cutlass::HostTensor<ElementInputB, typename ImplicitGemm::LayoutB> tensor_b(options.filter_size);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_c(oshape);
+ cutlass::HostTensor<ElementOutput, typename ImplicitGemm::LayoutC> tensor_ref_c(oshape);
+
+ cutlass::conv::Conv2dProblemSize problem_size(
+ options.input_size,
+ options.filter_size,
+ options.padding,
+ options.conv_stride,
+ options.dilation,
+ options.output_size(),
+ cutlass::conv::Mode::kCrossCorrelation,
+ 1
+ );
+
+ using ElementComputeEpilogue = typename ImplicitGemm::ElementCompute;
+ typename ImplicitGemm::Arguments arguments{
+ problem_size,
+ tensor_a.device_ref(),
+ tensor_b.device_ref(),
+ tensor_c.device_ref(),
+ tensor_c.device_ref(),
+ {ElementComputeEpilogue(1), ElementComputeEpilogue(0)},
+ };
+
+ ImplicitGemm implicit_gemm_op;
+ size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);
+ cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);
+ auto status = implicit_gemm_op.can_implement(arguments);
+ CUTLASS_CHECK(status);
+
+ status = implicit_gemm_op.initialize(arguments, workspace.get());
+ CUTLASS_CHECK(status);
+ status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+
+ cudaEvent_t events[2];
+ for (auto & event : events) {
+ cudaEventCreate(&event);
+ }
+ cudaEventRecord(events[0]);
+
+ for (int iteration = 0; iteration < 100; ++iteration) {
+ auto status = implicit_gemm_op();
+ CUTLASS_CHECK(status);
+ }
+
+ cudaEventRecord(events[1]);
+ cudaEventSynchronize(events[1]);
+ float runtime_ms = 0;
+ cudaEventElapsedTime(&runtime_ms, events[0], events[1]);
+
+ for (auto event : events) {
+ (void)cudaEventDestroy(event);
+ }
+ return double(runtime_ms) / 100.0;
Review comment:
What is the reasoning for putting the timing code inside this template instead of having the template being just the kernel and using our built in timers instead?
##########
File path: python/tvm/relay/op/strategy/cuda.py
##########
@@ -324,7 +324,10 @@ def conv2d_strategy_cuda(attrs, inputs, out_type, target):
plevel=25,
)
- elif is_depthwise_conv2d(data.shape, layout, kernel.shape, kernel_layout, groups):
+ elif (
+ is_depthwise_conv2d(data.shape, layout, kernel.shape, kernel_layout, groups)
+ and "cudnn" not in target.libs
Review comment:
Don't we still want to consider our built in schedules if cudnn is enabled?
--
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] comaniac commented on a change in pull request #9737: [CUTLASS] Add conv2d profiler
Posted by GitBox <gi...@apache.org>.
comaniac commented on a change in pull request #9737:
URL: https://github.com/apache/tvm/pull/9737#discussion_r769005791
##########
File path: python/tvm/contrib/cutlass/gen_conv2d.py
##########
@@ -121,27 +131,70 @@ def get_default(self, out_dtype):
data_type = gemm_profile_result["data_type"]
return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ def check_align(self, op_name, C, K):
+ """Filter out kernels that cannot be supported."""
+ aligns = re.findall(r"align[1|2|4|8]", op_name)
+ assert len(aligns) == 1
+ align = int(aligns[0][-1])
+ return all([dim % align == 0 for dim in [C, K]])
+
def profile(
- self, d_shape, w_shape, out_shape, out_dtype, profile_all=True, use_multiprocessing=False
+ self,
+ d_shape,
+ w_shape,
+ padding,
+ stride,
+ dilation,
+ out_dtype,
+ profile_all=True,
+ use_multiprocessing=False,
):
"""Profile and select the best kernel from candidate kernels.
If profile_all is False, return immediately after the first applicable kernel is found.
If use_multiprocessing is True, compile all profiler executables in parallel.
"""
- B, _, _, IC = d_shape
+ N, H, W, IC = d_shape
OC, R, S, _ = w_shape
- _, P, Q, _ = out_shape
+ workload = (
+ N,
+ H,
+ W,
+ IC,
+ OC,
+ R,
+ S,
+ padding[0],
+ padding[1],
+ stride[0],
+ stride[1],
+ dilation[0],
+ dilation[1],
+ )
- M = B * P * Q
- N = OC
- K = R * S * IC
+ if workload in self.cache:
+ return self.cache[workload]
- gemm_profile_result = self.gemm_profiler.profile(
- M, N, K, out_dtype, profile_all=profile_all, use_multiprocessing=use_multiprocessing
- )
+ ops = GENERATOR_FUNC_TABLE[self.sm](out_dtype, op_creator=create_conv2d_operator)
+ ops = list(filter(lambda op: self.check_align(op["name"], IC, OC), ops))
- tile_description = gemm_profile_result["tile_description"]
- alignment = gemm_profile_result["alignment"]
- data_type = gemm_profile_result["data_type"]
+ for op in ops:
+ op["runtime"] = -1
- return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ if profile_all:
+ self.engine.compile_all(ops, use_multiprocessing)
+
+ args = (
+ "--n=%d --h=%d --w=%d --c=%d --k=%d --r=%d --s=%d --pad_h=%d --pad_w=%d "
+ "--stride_h=%d --stride_w=%d --dilation_h=%d --dilation_w=%d"
+ ) % workload
+
+ for op in ops:
+ out = self.engine.evaluate(op, args.split(" "))
+ op["runtime"] = out
+ if out > 0 and profile_all is False:
Review comment:
nit
```suggestion
if out > 0 and not profile_all:
```
##########
File path: python/tvm/contrib/cutlass/conv2d_profiler.py
##########
@@ -0,0 +1,163 @@
+# 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=import-outside-toplevel, invalid-name
+"""Instantiate a C++ source for profiling CUTLASS kernels."""
+
+
+class Conv2dProfilerEmitter(object):
Review comment:
I raised this topic before in the GEMM profiler PR, but I agreed with @masahi that it seems not much to share and CUTLASS basically only supports GEMM and Conv2D. Accordingly, it might be a bit overkill to have a common base class at least for now.
##########
File path: python/tvm/contrib/cutlass/gen_conv2d.py
##########
@@ -121,27 +131,70 @@ def get_default(self, out_dtype):
data_type = gemm_profile_result["data_type"]
return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ def check_align(self, op_name, C, K):
+ """Filter out kernels that cannot be supported."""
+ aligns = re.findall(r"align[1|2|4|8]", op_name)
+ assert len(aligns) == 1
+ align = int(aligns[0][-1])
+ return all([dim % align == 0 for dim in [C, K]])
+
def profile(
- self, d_shape, w_shape, out_shape, out_dtype, profile_all=True, use_multiprocessing=False
+ self,
+ d_shape,
+ w_shape,
+ padding,
+ stride,
+ dilation,
+ out_dtype,
+ profile_all=True,
+ use_multiprocessing=False,
):
"""Profile and select the best kernel from candidate kernels.
If profile_all is False, return immediately after the first applicable kernel is found.
If use_multiprocessing is True, compile all profiler executables in parallel.
"""
- B, _, _, IC = d_shape
+ N, H, W, IC = d_shape
OC, R, S, _ = w_shape
- _, P, Q, _ = out_shape
+ workload = (
+ N,
+ H,
+ W,
+ IC,
+ OC,
+ R,
+ S,
+ padding[0],
+ padding[1],
+ stride[0],
+ stride[1],
+ dilation[0],
+ dilation[1],
+ )
- M = B * P * Q
- N = OC
- K = R * S * IC
+ if workload in self.cache:
+ return self.cache[workload]
- gemm_profile_result = self.gemm_profiler.profile(
- M, N, K, out_dtype, profile_all=profile_all, use_multiprocessing=use_multiprocessing
- )
+ ops = GENERATOR_FUNC_TABLE[self.sm](out_dtype, op_creator=create_conv2d_operator)
+ ops = list(filter(lambda op: self.check_align(op["name"], IC, OC), ops))
- tile_description = gemm_profile_result["tile_description"]
- alignment = gemm_profile_result["alignment"]
- data_type = gemm_profile_result["data_type"]
+ for op in ops:
+ op["runtime"] = -1
- return create_conv2d_operator([tile_description], data_type, [alignment])[0]
+ if profile_all:
+ self.engine.compile_all(ops, use_multiprocessing)
+
+ args = (
+ "--n=%d --h=%d --w=%d --c=%d --k=%d --r=%d --s=%d --pad_h=%d --pad_w=%d "
+ "--stride_h=%d --stride_w=%d --dilation_h=%d --dilation_w=%d"
+ ) % workload
+
+ for op in ops:
+ out = self.engine.evaluate(op, args.split(" "))
+ op["runtime"] = out
+ if out > 0 and profile_all is False:
+ break
+
+ valid_ops = filter(lambda op: op["runtime"] > 0, ops)
+ output = sorted(valid_ops, key=lambda i: i["runtime"])
Review comment:
Looks like you could just `output = min(valid_ops, key=lambda i: i["runtime"])`. Moreover, if you directly set the invalid runtime to `float("inf")` after `self.engine.evaluate`, you could also get rid of the filter.
--
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