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 18:01:45 UTC
[GitHub] [tvm] tkonolige commented on a change in pull request #9737: [CUTLASS] Add conv2d profiler
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