NV cpu golden conv vcore 计算

最新推荐文章于 2024-04-30 23:23:05 发布
最新推荐文章于 2024-04-30 23:23:05 发布
阅读量71 收藏
点赞数
文章标签： gpu算力 深度学习 卷积神经网络
原文链接：https://github.com/NVIDIA/cutlass/blob/main/tools/util/include/cutlass/util/reference/host/convolution.h
版权

https://github.com/NVIDIA/cutlass/blob/main/tools/util/include/cutlass/util/reference/host/convolution.h

/***************************************************************************************************
 * Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: BSD-3-Clause
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *
 * 1. Redistributions of source code must retain the above copyright notice, this
 * list of conditions and the following disclaimer.
 *
 * 2. Redistributions in binary form must reproduce the above copyright notice,
 * this list of conditions and the following disclaimer in the documentation
 * and/or other materials provided with the distribution.
 *
 * 3. Neither the name of the copyright holder nor the names of its
 * contributors may be used to endorse or promote products derived from
 * this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
 * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
 * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
 * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
 * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
 * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 *
 **************************************************************************************************/

/*! \file
    \brief Reference implementation for convolution in host-side code.
*/

#pragma once

#include "cutlass/coord.h"
#include "cutlass/functional.h"
#include "cutlass/layout/tensor.h"
#include "cutlass/numeric_conversion.h"
#include "cutlass/numeric_types.h"
#include "cutlass/tensor_ref.h"
#include "cutlass/tensor_view.h"
#include "cutlass/conv/convolution.h"
#include "cutlass/conv/conv2d_problem_size.h"
#include "cutlass/conv/conv3d_problem_size.h"
#include <iostream>

namespace cutlass {
namespace reference {
namespace host {


/// Forward propagation

/// y = conv2d(x, w)
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ElementD = ElementC,
  typename ConvertOp = NumericConverter<ElementD, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv2dFprop(
  conv::Conv2dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_x,
  TensorRef<ElementB, LayoutB> tensor_w,
  TensorRef<ElementC, LayoutC> tensor_y_in,
  TensorRef<ElementD, LayoutC> tensor_y_out,
  ElementCompute alpha,
  ElementCompute beta) {

  ConvertOp convert_op;
  InnerProductOp inner_product_op;

  // Apply MMA and accumulate ElementAccumulator
  for (int n = 0; n < problem_size.N; ++n) {
    for (int p = 0; p < problem_size.P; ++p) {
      for (int q = 0; q < problem_size.Q; ++q) {
        for (int k = 0; k < problem_size.K; ++k) {

          int group_idx = k / (problem_size.K / problem_size.groups);
          int channels_per_group = problem_size.C / problem_size.groups;

          ElementAccumulator acc = ElementAccumulator();

          for (int r = 0; r < problem_size.R; ++r) {
            for (int s = 0; s < problem_size.S; ++s) {
              for (int c = 0; c < channels_per_group; ++c) {

                int filter_r = r;
                int filter_s = s;

                if (problem_size.mode == cutlass::conv::Mode::kConvolution) {
                  filter_r = problem_size.R - 1 - r;
                  filter_s = problem_size.S - 1 - s;
                }

                int h = p * problem_size.stride_h - problem_size.pad_h + filter_r * problem_size.dilation_h;
                int w = q * problem_size.stride_w - problem_size.pad_w + filter_s * problem_size.dilation_w;

                if (h >= 0 && h < problem_size.H && w >= 0 && w < problem_size.W) {

                  ElementA a = tensor_x.at({n, h, w, c + group_idx * channels_per_group});
                  ElementB b = tensor_w.at({k, r, s, c});

                  acc = inner_product_op(ElementAccumulator(a), ElementAccumulator(b), acc);

                }
              }
            }
          }

          // Apply Epilogue, compute ElementCompute, convert and store ElementC
          ElementC c_ref = ElementC();

          if (beta != ElementCompute()) {
            c_ref = tensor_y_in.at(cutlass::make_Coord(n, p, q, k));
          }

          tensor_y_out.at(cutlass::make_Coord(n, p, q, k)) =
              convert_op(alpha * ElementCompute(acc) + beta * ElementCompute(c_ref));
        }
      }
    }
  }
}

/// Depthwise-separable convolution
template <typename ElementA,
          typename LayoutA,
          typename ElementB,
          typename LayoutB,
          typename ElementC,
          typename LayoutC,
          typename ElementCompute,
          typename ElementAccumulator = ElementCompute,
          typename ElementD = ElementC,
          typename ConvertOp = NumericConverter<ElementD, ElementCompute>,
          typename InnerProductOp = multiply_add<ElementAccumulator>>
void Depsep_Fprop(cutlass::TensorView<ElementA, LayoutA> tensor_A,
                  cutlass::TensorView<ElementB, LayoutB> tensor_B,
                  cutlass::TensorView<ElementC, LayoutC> tensor_C,
                  cutlass::TensorView<ElementD, LayoutC> tensor_D,
                  ElementCompute alpha,
                  ElementCompute beta,
                  cutlass::Tensor4DCoord padding = cutlass::Tensor4DCoord(),
                  cutlass::Coord<2> conv_stride = cutlass::Coord<2>(),
                  cutlass::Coord<2> dilation = cutlass::Coord<2>(),
                  cutlass::conv::Mode mode = cutlass::conv::Mode::kCrossCorrelation) {

  ConvertOp convert_op;
  InnerProductOp inner_product_op;

  // Apply MMA and accumulate ElementAccumulator
  for (int n = 0; n < tensor_C.extent().n(); ++n) {
    for (int p = 0; p < tensor_C.extent().h(); ++p) {
      for (int q = 0; q < tensor_C.extent().w(); ++q) {
        for (int g = 0; g < tensor_C.extent().c(); ++g) {
          ElementAccumulator acc = ElementAccumulator();
          for (int r = 0; r < tensor_B.extent().h(); ++r) {
            for (int s = 0; s < tensor_B.extent().w(); ++s) {
              
              // input activation H and W
              int h = p * conv_stride[0] - padding[0] + r * dilation[0];
              int w = q * conv_stride[1] - padding[2] + s * dilation[1];

              if (h < tensor_A.extent().h() && h >= 0 && w < tensor_A.extent().w() && w >= 0) {
                ElementA a = tensor_A.at(cutlass::make_Coord(n, h, w, g));

                ElementB b = (mode == cutlass::conv::Mode::kCrossCorrelation)
                                   ? tensor_B.at(cutlass::make_Coord(g, r, s, 0))
                                   : tensor_B.at(cutlass::make_Coord(
                                         g, tensor_B.extent().h() - r - 1, tensor_B.extent().w() - s - 1, 0));

                acc = inner_product_op(ElementAccumulator(a), ElementAccumulator(b), acc);
              }
            }
          }

          // Apply Epilogue, compute ElementCompute, convert and store ElementC
          ElementC c_ref = tensor_C.at(cutlass::make_Coord(n, p, q, g));
          tensor_D.at(cutlass::make_Coord(n, p, q, g)) =
              convert_op(alpha * ElementCompute(acc) + beta * ElementCompute(c_ref));
        }
      }
    }
  }
}


/// Dgrad

/// dx = dgrad(dy, w)
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ElementD = ElementC,
  typename ConvertOp = NumericConverter<ElementD, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv2dDgrad(
  cutlass::conv::Conv2dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_dy,
  TensorRef<ElementB, LayoutB> tensor_w,
  TensorRef<ElementC, LayoutC> tensor_dx_in,
  TensorRef<ElementD, LayoutC> tensor_dx_out,
  ElementCompute alpha,
  ElementCompute beta) {

  ConvertOp convert_op;
  InnerProductOp inner_product_op;

  // Apply MMA and accumulate ElementAccumulator
  for (int n = 0; n < problem_size.N; ++n) {
    for (int h = 0; h < problem_size.H; ++h) {
      for (int w = 0; w < problem_size.W; ++w) {
        for (int c = 0; c < problem_size.C; ++c) {

          ElementAccumulator acc = ElementAccumulator();

          for (int r = 0; r < problem_size.R; ++r) {
            for (int s = 0; s < problem_size.S; ++s) {
              for (int k = 0; k < problem_size.K; ++k) {

                int filter_r = r;
                int filter_s = s;

                if (problem_size.mode == cutlass::conv::Mode::kConvolution) {
                  filter_r = problem_size.R - 1 - r;
                  filter_s = problem_size.S - 1 - s;
                }

                int p = h + problem_size.pad_h - filter_r * problem_size.dilation_h;
                int q = w + problem_size.pad_w - filter_s * problem_size.dilation_w;

                if (p >= 0 && (p % problem_size.stride_h) == 0 && 
                    q >= 0 && (q % problem_size.stride_w) == 0) {

                  p = p / problem_size.stride_h;
                  q = q / problem_size.stride_w;
#if 0
                  std::cout << "row:" 
                  << n * problem_size.H * problem_size.W +
                    h * problem_size.W +
                    w << " "
                  << "n, p, q: (" 
                  << n << ", "
                  << p << ", "
                  << q << ") * "
                  << "r, s: (" 
                  << r << ", "
                  << s << ") [" 
                  << ((p < problem_size.P && q < problem_size.Q) ? "true":"false") << "]"        
                  << std::endl;
#endif
                  if (p < problem_size.P && q < problem_size.Q) {

                    ElementA a = tensor_dy.at(cutlass::make_Coord(n, p, q, k));
                    ElementB b = tensor_w.at(cutlass::make_Coord(k, r, s, c));

                    acc = inner_product_op(ElementAccumulator(a), ElementAccumulator(b), acc);
                  }
                }

              } // for (K)
            } // for (S)
          } // for (R)

          // Apply Epilogue, compute ElementCompute, convert and store ElementC
          ElementC c_ref = ElementC();

          if (beta != ElementCompute()) {
            c_ref = tensor_dx_in.at(cutlass::make_Coord(n, h, w, c));
          }

          tensor_dx_out.at(cutlass::make_Coord(n, h, w, c)) =
              convert_op(alpha * ElementCompute(acc) + beta * ElementCompute(c_ref));

        } // for (C)
      } // for (W)
    } // for (H)
  } // for (N)
}


/// Wgrad

/// dw = wgrad(dy, x)
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ElementD = ElementC,
  typename ConvertOp = NumericConverter<ElementD, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv2dWgrad(
  cutlass::conv::Conv2dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_dy,
  TensorRef<ElementB, LayoutB> tensor_x,
  TensorRef<ElementC, LayoutC> tensor_dw_in,
  TensorRef<ElementD, LayoutC> tensor_dw_out,
  ElementCompute alpha,
  ElementCompute beta) {
  
  InnerProductOp inner_product_op;
  ConvertOp convert_op;

  // Apply MMA and accumulate ElementAccumulator
  for (int k = 0; k < problem_size.K; ++k) {
    for (int r = 0; r < problem_size.R; ++r) {
      for (int s = 0; s < problem_size.S; ++s) {
        for (int c = 0; c < problem_size.C; ++c) {

          ElementAccumulator acc = ElementAccumulator();

          for (int n = 0; n < problem_size.N; ++n) {
            for (int p = 0; p < problem_size.P; ++p) {
              for (int q = 0; q < problem_size.Q; ++q) {
                  
                cutlass::Tensor4DCoord b_coord;
                
                int filter_r = r;
                int filter_s = s; 

                if (problem_size.mode == cutlass::conv::Mode::kConvolution) {
                  filter_r = problem_size.R - 1 - r;
                  filter_s = problem_size.S - 1 - s;
                }

                b_coord = make_Coord(
                    n,
                    p * problem_size.stride_h - problem_size.pad_h + filter_r * problem_size.dilation_h,
                    q * problem_size.stride_w - problem_size.pad_w + filter_s * problem_size.dilation_w,
                    c);

                if (b_coord.h() < problem_size.H && b_coord.h() >= 0 &&
                    b_coord.w() < problem_size.W && b_coord.w() >= 0) {

                  ElementAccumulator a = ElementAccumulator(tensor_dy.at(cutlass::make_Coord(n, p, q, k)));
                  ElementAccumulator b = ElementAccumulator(tensor_x.at(b_coord));
                  acc = inner_product_op(a, b, acc);
                }
              }
            }
          }

          // Apply Epilogue, compute ElementCompute, convert and store ElementC
          ElementC c_ref = ElementC();

          if (beta != ElementCompute()) {
            c_ref = tensor_dw_in.at(cutlass::make_Coord(k, r, s, c));
          }

          tensor_dw_out.at(cutlass::make_Coord(k, r, s, c)) =
              convert_op(alpha * ElementCompute(acc) + beta * ElementCompute(c_ref));

        } // for (C)
      } // for (S)
    } // for (R)
  } // for (K)
}

/// Generic 2D convolution targeting Conv2dFprop, Conv2dDgrad, and Conv2dWgrad.
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ElementD = ElementC,
  typename ConvertOp = NumericConverter<ElementD, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv2d(
  conv::Operator convolutional_operator,
  conv::Conv2dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_A,
  TensorRef<ElementB, LayoutB> tensor_B,
  TensorRef<ElementC, LayoutC> tensor_C,
  TensorRef<ElementD, LayoutC> tensor_D,
  ElementCompute alpha,
  ElementCompute beta) {

  switch (convolutional_operator) {
  case conv::Operator::kFprop:
    Conv2dFprop<
      ElementA, LayoutA,
      ElementB, LayoutB,
      ElementC, LayoutC,
      ElementCompute,
      ElementAccumulator,
      ElementD,
      ConvertOp, InnerProductOp
    >(problem_size, tensor_A, tensor_B, tensor_C, tensor_D, alpha, beta);
    break;

  case conv::Operator::kDgrad:
    Conv2dDgrad<
      ElementA, LayoutA,
      ElementB, LayoutB,
      ElementC, LayoutC,
      ElementCompute,
      ElementAccumulator,
      ElementD,
      ConvertOp, InnerProductOp
    >(problem_size, tensor_A, tensor_B, tensor_C, tensor_D, alpha, beta);
    break;

  case conv::Operator::kWgrad:
    Conv2dWgrad<
      ElementA, LayoutA,
      ElementB, LayoutB,
      ElementC, LayoutC,
      ElementCompute,
      ElementAccumulator,
      ElementD,
      ConvertOp, InnerProductOp
    >(problem_size, tensor_A, tensor_B, tensor_C, tensor_D, alpha, beta);
    break;

  default:
    break;  
  }
}


/// 3D convolution 

/// y = conv3d(x, w)
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ConvertOp = NumericConverter<ElementC, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv3dFprop(
  conv::Conv3dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_x,
  TensorRef<ElementB, LayoutB> tensor_w,
  TensorRef<ElementC, LayoutC> tensor_y_in,
  TensorRef<ElementC, LayoutC> tensor_y_out,
  ElementCompute alpha,
  ElementCompute beta) {

  ConvertOp convert_op;
  InnerProductOp inner_product_op;

  // Apply MMA and accumulate ElementAccumulator
  for (int n = 0; n < problem_size.N; ++n) {
    for (int z = 0; z < problem_size.Z; ++z) {
      for (int p = 0; p < problem_size.P; ++p) {
        for (int q = 0; q < problem_size.Q; ++q) {
          for (int k = 0; k < problem_size.K; ++k) {

            ElementAccumulator acc = ElementAccumulator();

            for (int t = 0; t < problem_size.T; ++t) {
              for (int r = 0; r < problem_size.R; ++r) {
                for (int s = 0; s < problem_size.S; ++s) {
                  for (int c = 0; c < problem_size.C; ++c) {

                    int filter_t = t;
                    int filter_r = r;
                    int filter_s = s;

                    if (problem_size.mode == cutlass::conv::Mode::kConvolution) {
                      filter_t = problem_size.T - 1 - t;
                      filter_r = problem_size.R - 1 - r;
                      filter_s = problem_size.S - 1 - s;
                    }

                    int d = z * problem_size.stride_d - problem_size.pad_d + filter_t * problem_size.dilation_d;
                    int h = p * problem_size.stride_h - problem_size.pad_h + filter_r * problem_size.dilation_h;
                    int w = q * problem_size.stride_w - problem_size.pad_w + filter_s * problem_size.dilation_w;

                    if (d >= 0 && d < problem_size.D && 
                      h >=0 && h < problem_size.H && 
                      w >= 0 && w < problem_size.W) {

                      ElementA a = tensor_x.at({n, d, h, w, c});
                      ElementB b = tensor_w.at({k, t, r, s, c});
                      
                      acc = inner_product_op(ElementAccumulator(a), ElementAccumulator(b), acc);
                    }
                  }
                }
              }
            }

            // Apply Epilogue, compute ElementCompute, convert and store ElementC
            ElementC c_ref = ElementC();

            if (beta != ElementCompute()) {
              c_ref = tensor_y_in.at(cutlass::make_Coord(n, z, p, q, k));
            }

            tensor_y_out.at(cutlass::make_Coord(n, z, p, q, k)) =
                convert_op(alpha * ElementCompute(acc) + beta * ElementCompute(c_ref));
          }
        }
      }
    }
  }
}


/// Dgrad

/// dx = dgrad(dy, w)
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ConvertOp = NumericConverter<ElementC, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv3dDgrad(
  cutlass::conv::Conv3dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_dy,
  TensorRef<ElementB, LayoutB> tensor_w,
  TensorRef<ElementC, LayoutC> tensor_dx_in,
  TensorRef<ElementC, LayoutC> tensor_dx_out,
  ElementCompute alpha,
  ElementCompute beta) {

  ConvertOp convert_op;
  InnerProductOp inner_product_op;

  // Apply MMA and accumulate ElementAccumulator
  for (int n = 0; n < problem_size.N; ++n) {
    for (int d = 0; d < problem_size.D; ++d) {
      for (int h = 0; h < problem_size.H; ++h) {
        for (int w = 0; w < problem_size.W; ++w) {
          for (int c = 0; c < problem_size.C; ++c) {

            ElementAccumulator acc = ElementAccumulator();

            for (int t = 0; t < problem_size.T; ++t) {
              for (int r = 0; r < problem_size.R; ++r) {
                for (int s = 0; s < problem_size.S; ++s) {
                  for (int k = 0; k < problem_size.K; ++k) {

                    int filter_t = t;
                    int filter_r = r;
                    int filter_s = s;

                    if (problem_size.mode == cutlass::conv::Mode::kConvolution) {
                      filter_t = problem_size.T - 1 - t;
                      filter_r = problem_size.R - 1 - r;
                      filter_s = problem_size.S - 1 - s;
                    }

                    int z = d + problem_size.pad_d - filter_t * problem_size.dilation_d;
                    int p = h + problem_size.pad_h - filter_r * problem_size.dilation_h;
                    int q = w + problem_size.pad_w - filter_s * problem_size.dilation_w;

                    if (z >= 0 && (z % problem_size.stride_d) == 0 &&
                        p >= 0 && (p % problem_size.stride_h) == 0 && 
                        q >= 0 && (q % problem_size.stride_w) == 0) {

                      z = z / problem_size.stride_d;
                      p = p / problem_size.stride_h;
                      q = q / problem_size.stride_w;
                      
                      if (z < problem_size.Z && p < problem_size.P && q < problem_size.Q) {

                        ElementA a = tensor_dy.at(cutlass::make_Coord(n, z, p, q, k));
                        ElementB b = tensor_w.at(cutlass::make_Coord(k, t, r, s, c));

                        acc = inner_product_op(ElementAccumulator(a), ElementAccumulator(b), acc);
                      }
                    }

                  } // for (K)
                } // for (S)
              } // for (R)
            } // for (T)

            // Apply Epilogue, compute ElementCompute, convert and store ElementC
            ElementC c_ref = ElementC();

            if (beta != ElementCompute()) {
              c_ref = tensor_dx_in.at(cutlass::make_Coord(n, d, h, w, c));
            }

            tensor_dx_out.at(cutlass::make_Coord(n, d, h, w, c)) =
                convert_op(alpha * ElementCompute(acc) + beta * ElementCompute(c_ref));

          } // for (C)
        } // for (W)
      } // for (H)
    } // for (D)
  } // for (N)
}


/// Wgrad

/// dw = wgrad(dy, x)
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ConvertOp = NumericConverter<ElementC, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv3dWgrad(
  cutlass::conv::Conv3dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_dy,
  TensorRef<ElementB, LayoutB> tensor_x,
  TensorRef<ElementC, LayoutC> tensor_dw_in,
  TensorRef<ElementC, LayoutC> tensor_dw_out,
  ElementCompute alpha,
  ElementCompute beta) {
  
  InnerProductOp inner_product_op;
  ConvertOp convert_op;

  // Apply MMA and accumulate ElementAccumulator
  for (int k = 0; k < problem_size.K; ++k) {
    for (int t = 0; t < problem_size.T; ++t) {
      for (int r = 0; r < problem_size.R; ++r) {
        for (int s = 0; s < problem_size.S; ++s) {
          for (int c = 0; c < problem_size.C; ++c) {

            ElementAccumulator acc = ElementAccumulator();

            for (int n = 0; n < problem_size.N; ++n) {
              for (int z = 0; z < problem_size.Z; ++z) {
                for (int p = 0; p < problem_size.P; ++p) {
                  for (int q = 0; q < problem_size.Q; ++q) {
                      
                    int filter_t = t;     
                    int filter_r = r;
                    int filter_s = s; 

                    if (problem_size.mode == cutlass::conv::Mode::kConvolution) {
                      filter_t = problem_size.T - 1 - t;
                      filter_r = problem_size.R - 1 - r;
                      filter_s = problem_size.S - 1 - s;
                    }

                    Tensor5DCoord b_coord = make_Coord(
                        n,
                        z * problem_size.stride_d - problem_size.pad_d + filter_t * problem_size.dilation_d,
                        p * problem_size.stride_h - problem_size.pad_h + filter_r * problem_size.dilation_h,
                        q * problem_size.stride_w - problem_size.pad_w + filter_s * problem_size.dilation_w,
                        c);

                    if (b_coord.d() < problem_size.D && b_coord.d() >= 0 &&
                        b_coord.h() < problem_size.H && b_coord.h() >= 0 &&
                        b_coord.w() < problem_size.W && b_coord.w() >= 0) {

                      ElementAccumulator a = ElementAccumulator(tensor_dy.at(cutlass::make_Coord(n, z, p, q, k)));
                      ElementAccumulator b = ElementAccumulator(tensor_x.at(b_coord));

                      acc = inner_product_op(a, b, acc);
                    }
                  }
                }
              }
            }

            // Apply Epilogue, compute ElementCompute, convert and store ElementC
            ElementC c_ref = ElementC();

            if (beta != ElementCompute()) {
              c_ref = tensor_dw_in.at(cutlass::make_Coord(k, t, r, s, c));
            }

            tensor_dw_out.at(cutlass::make_Coord(k, t, r, s, c)) =
                convert_op(alpha * ElementCompute(acc) + beta * ElementCompute(c_ref));

          } // for (C)
        } // for (S)
      } // for (R)
    } // for (T)
  } // for (K)
}

///

/// Generic 3D convolution targeting Conv2dFprop, Conv2dDgrad, and Conv2dWgrad.
template <
  typename ElementA,
  typename LayoutA,
  typename ElementB,
  typename LayoutB,
  typename ElementC,
  typename LayoutC,
  typename ElementCompute,
  typename ElementAccumulator = ElementCompute,
  typename ConvertOp = NumericConverter<ElementC, ElementCompute>,
  typename InnerProductOp = multiply_add<ElementAccumulator>
>
void Conv3d(
  conv::Operator convolutional_operator,
  conv::Conv3dProblemSize problem_size,
  TensorRef<ElementA, LayoutA> tensor_A,
  TensorRef<ElementB, LayoutB> tensor_B,
  TensorRef<ElementC, LayoutC> tensor_C,
  TensorRef<ElementC, LayoutC> tensor_D,
  ElementCompute alpha,
  ElementCompute beta) {

  switch (convolutional_operator) {
  case conv::Operator::kFprop:
    Conv3dFprop<
      ElementA, LayoutA,
      ElementB, LayoutB,
      ElementC, LayoutC,
      ElementCompute,
      ElementAccumulator,
      ConvertOp, InnerProductOp
    >(problem_size, tensor_A, tensor_B, tensor_C, tensor_D, alpha, beta);
    break;

  case conv::Operator::kDgrad:
    Conv3dDgrad<
      ElementA, LayoutA,
      ElementB, LayoutB,
      ElementC, LayoutC,
      ElementCompute,
      ElementAccumulator, 
      ConvertOp, InnerProductOp
    >(problem_size, tensor_A, tensor_B, tensor_C, tensor_D, alpha, beta);
    break;

  case conv::Operator::kWgrad:
    Conv3dWgrad<
      ElementA, LayoutA,
      ElementB, LayoutB,
      ElementC, LayoutC,
      ElementCompute,
      ElementAccumulator, 
      ConvertOp, InnerProductOp
    >(problem_size, tensor_A, tensor_B, tensor_C, tensor_D, alpha, beta);
    break;

  default:
    break;  
  }
}

/

}  // namespace host
}  // namespace reference
}  // namespace cutlass

/

zou丢的生化环
关注
  • 0
    点赞
  • 踩
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
conv_core_HLS_
10-04
HLS卷积运算源代码，包括.cpp文件，.h文件以及测试文件
cutlass 第一节：下载 编译 运行 cutlass 的一些备忘信息 ubuntu docker
eloudy的专栏
02-19 2668
—————v100———cutlass—start——————————————————————————————————————————— __________web site__________________________________________________ https://developer.nvidia.com/blog/cutlass-linear-algebra-cuda/ https://github.com/NVIDIA/cutlass 参考READ.md ______
使用cutlass实现多种精度的GEMM，支持cuda core与tensor core
weixin_42764932的博客
03-27 705
使用cutlass实现多种精度的GEMM，附有完整代码与cmakelist
项目复现---nvdiffrec
j1829634的博客
07-04 2538
复现项目GitHub - NVlabs/nvdiffrec: Official code for the CVPR 2022 (oral) paper "Extracting Triangular 3D Models, Materials, and Lighting From Images".1、前面下载安装顺利，运行如下命令出现报错 上面的命令是一个结合的命令，首先会用git去下载项目，然后进入项目中setup.py文件所在的位置，接着运行这个文件，安装项目，故可以手动下载项目，然后进行项目的安装2、
Linux Ubuntu20.04系统下安装TVM
m0_48403174的博客
12-13 4096
本文主要讲述了在Linux系统下安装TVM以及前期CMake、LLVM等依赖的配置流程和可能遇到的问题及其解决方法。
conv2d原理、计算过程、pytorch使用
热门推荐
纸上得来终觉浅，绝知此事要躬行！
12-30 1万+
文章目录conv2d卷积过程单通道三（多）通道多个输出通道（**即改变输出通道**）对应代码 【学习于刘二大大视频：https://www.bilibili.com/video/BV1Y7411d7Ys?p=10】 conv2d 卷积过程 向两个维度进行卷积 单通道 三（多）通道 一个通道需要一个kernel 画图： 多个输出通道（即改变输出通道） 卷积核的数量==输出通道的通道数 每个卷积核的数量==输入通道的通道数 对应代码 self.conv1 = nn.Conv2d(in_channel
在Pytorch中计算卷积方法的区别详解(conv2d的区别)
01-20
class torch.nn.Conv2d(in_channels, out_channels, kernel_size, stride=1, padding=0, dilation=1, groups=1, bias=True) 对由多个特征平面组成的输入信号进行2D的卷积操作。详解 torch.nn.functional.conv2d...
tf.nn.conv1d
01-20
tf.nn.conv1d(value, filters, stride, padding) value: [batch, in_width, in_channels] batch是输入文本的个数 in_width是样本的单词个数 in_channels是每个单词的维度 (float32 or float64，必须是np.array，因为...
TensorFlow tf.nn.conv2d实现卷积的方式
01-20
tf.nn.conv2d(input, filter, strides, padding, use_cudnn_on_gpu=None, name=None) 除去name参数用以指定该操作的name，与方法有关的一共五个参数： input： 指需要做卷积的输入图像，它要求是一个Tensor，具有...
Pytorch.nn.conv2d 过程验证方式(单,多通道卷积过程)
12-23
今天在看文档的时候，发现pytorch 的conv操作不是很明白，于是有了一下记录 首先提出两个问题： 1.输入图片是单通道情况下的filters是如何操作的？ 即一通道卷积核卷积过程 2.输入图片是多通道情况下的filters是如何...
GpuMall的GPU算力资源池化技术有何优势？
GpuMall的博客
04-26 363
GpuMall的GPU算力资源池化技术具有显著的优势，这些优势使得其在智算云领域脱颖而出，为用户提供了高效、灵活且可靠的GPU算力服务。
【深度学习】YOLOv5，金属表面的缺陷检测，GC10-DET数据集
q742971636的博客
04-26 404
数据集的lable文件夹下是每个图片的框和类别标记，是xml格式。
深度学习突破：LLaMA-MoE模型的高效训练策略
人工智能前沿分享
04-26 1322
在人工智能领域，大模型（LLM）的崛起带来了前所未有的进步，但随之而来的是巨大的计算资源需求。为了解决这一问题，Mixture-of-Expert（MoE）模型架构应运而生，而LLaMA-MoE正是这一架构下的重要代表。LLaMA-MoE是一种基于LLaMA系列和SlimPajama的MoE模型，它通过将LLaMA的前馈网络（FFNs）划分为稀疏专家，并为每层专家插入top-K个门，从而显著减小模型大小，降低训练成本。这种方法不仅保持了模型的语言能力，同时实现了输入的高效处理。
基于自注意力机制的长短期记忆神经网络(LSTM-SelfAttention)的回归预测
沅的博客
04-27 995
它结合了LSTM和自注意力机制的优势，可以有效地处理时序数据中的长期依赖关系和关键部分。这个模型结合了两个不同的结构，即长短期记忆网络（LSTM）和自注意力机制，从而在时序数据预测中发挥各自的优势。LSTM-SelfAttention模型的优势在于它既能捕捉到时序数据中的长期依赖关系（通过LSTM），又能关注输入数据中的关键部分（通过自注意力机制）。LSTM在时序数据预测中表现出色，因为它可以捕捉到数据中的长期依赖关系。自注意力层：在LSTM层之后，引入自注意力机制，以关注时序数据中的不同部分。
Sarcasm detection论文解析 |用于微博讽刺检测的上下文增强卷积神经网络
最新发布
人工智能学习生
04-30 810
在卷积神经网络的基础上，我们为这项任务提出了两种不同的语境增强神经网络模型。与之前的研究相比，我们的神经网络模型将推文内容本身的特征和上下文特征以单词向量的形式整合到了一个模型中。实验结果表明，与最先进的离散模型和基于上下文的模型相比，我们提出的上下文增强神经网络模型具有更好的性能，证明了上下文增强神经网络模型在这项任务中的有效性。
目标检测（任务概述、机器学习时代的方法、深度学习时代的目标检测模型）
hanmo22357的博客
04-27 1092
目标检测任务的概念：尝试从一张输入图像中找出感兴趣的物体所在的区域。感兴趣物体所在的区域：一般通过RoI（Region of Interest）进行表示，往往是一个包含有感兴趣的物体的水平的矩形，有时候也会是倾斜的矩形或一个表示物体轮廓线的任意几何形状。目标检测和OCR的关系：目标检测是一种常见的图像处理任务，是OCR任务的基础。目标检测技术的发展：在21世纪初，研究者们提出了很多基于机器学习的目标检测算法；十多年后，深度学习将目标检测引领向了一个新时代。目标检测方法的分类。
深度学习之基于Matlab NN的伦敦房价预测
Q1744828575的博客
04-30 468
一、项目背景房价预测是房地产领域的一个重要问题，对于投资者、开发商以及政策制定者等都具有重要的指导意义。随着深度学习技术的不断发展，其在房价预测领域的应用也越来越广泛。本项目旨在利用Matlab神经网络（NN）工具箱，构建一个基于深度学习的伦敦房价预测模型，以实现对伦敦地区房价的准确预测。二、项目目标数据收集与预处理：收集伦敦地区的房价数据，包括房屋类型、面积、地理位置、周边环境等信息，并进行数据清洗、去噪、标准化等预处理操作，以提高数据质量。
NLP（9）--rnn实现中文分词
njh1147394013的博客
04-27 286
rnn分词
pytorch中conv1d的计算过程
03-27
Conv1d的计算过程可以分为以下几个步骤： 1. 输入数据格式转换：将输入的一维数据（通常是一个向量）转换为一个三维的Tensor，包含batch_size、通道数（channel）和序列长度（sequence length）。这是为了方便卷积操作的实现。 2. 卷积核初始化：初始化卷积核的权重和偏置项，一般使用随机初始化的方式。 3. 卷积操作：使用卷积核对输入的数据进行卷积操作，得到卷积结果。卷积的过程可以看做是以步长为1的滑动窗口在输入数据上进行的，每个窗口的大小为卷积核的大小。 4. 激活函数处理：对卷积结果进行激活函数处理，通常使用ReLU等激活函数，可以使得网络具有非线性映射的能力。 5. 池化操作：对激活函数处理后的结果进行池化操作，通常使用最大池化或平均池化等操作，可以减小特征图的大小，提取更加重要的特征。 6. 输出数据处理：将池化后的结果进行展开，得到一个向量，作为下一层的输入。 以上就是Conv1d的计算过程。在实际应用中，Conv1d被广泛应用于语音识别、自然语言处理等领域。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值