forked from apache/tvm
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[CUTLASS] Add conv2d profiler (apache#9737)
* Add cutlass conv2d profiler commit 1c0bbb2 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 12 18:29:03 2021 +0900 fix lint commit 463574c Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 12 17:28:38 2021 +0900 fixed conv2d check commit 588c5ab Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 12 15:05:27 2021 +0900 update test commit a447b57 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 12 14:54:52 2021 +0900 speed up profiling by removing initialization commit 93cd039 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 12 08:26:29 2021 +0900 fixed nhwc cudnn depthwise conv commit 6db7172 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sat Dec 11 15:39:05 2021 +0900 add cache commit f7d17a1 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sat Dec 11 15:05:38 2021 +0900 removed im2col profiling for conv2d commit b724f44 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 22:57:54 2021 +0900 black commit fe4687b Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 22:49:13 2021 +0900 fixed cmd arguement commit ab114f5 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 22:22:19 2021 +0900 conv2d profiler working commit 49ee61f Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 20:26:15 2021 +0900 add conv2d profiler commit 49e2c89 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 12 08:03:36 2021 +0900 do not offload depthwise conv2d commit cd83677 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 13:20:01 2021 +0900 lint fix commit 870823c Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 12:54:38 2021 +0900 add comment on IC == 3 case commit 6b780db Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 12:48:33 2021 +0900 check align on N dim commit 308c4da Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 12:34:42 2021 +0900 fixed check functions for fused cases, run infer type before mergecomposite commit 8d6a1bf Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 12:10:59 2021 +0900 test IC=3 convolution commit ffce47d Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 12:10:16 2021 +0900 use align1 kernel for unusual channel cases (IC = 3 etc) commit 6cdf205 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 12:06:56 2021 +0900 add dtype and layout check in parttern match commit 7743cc6 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 10:40:53 2021 +0900 add sm75 kernels to sm80 profilings commit efceccb Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 10:40:42 2021 +0900 skip legalize when batch size is dynamic commit 65fbc0a Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 10 10:36:36 2021 +0900 bug fix in im2col encoding * minor fix * lint fix * allow autotvm NCHW depthwise conv2d schedule even if -libs=cudnn * Update python/tvm/contrib/cutlass/gen_conv2d.py Co-authored-by: Cody Yu <comaniac0422@gmail.com> * simplify processing profiler outputs * more simplify * fix runtime check Co-authored-by: Cody Yu <comaniac0422@gmail.com>
- Loading branch information
Showing
7 changed files
with
265 additions
and
50 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -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; | ||
| } | ||
| int main(int argc, char const **args) { | ||
| Options options; | ||
| options.parse(argc, args); | ||
| std::cout << profile_convolution(options) << std::endl; | ||
| return 0; | ||
| } | ||
| """ | ||
| ) | ||
|
|
||
| def emit(self, op_def, op_name): | ||
| src = self.template.render(OperatorDef=op_def, OperatorName=op_name) | ||
| return src |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.