【模型推理】談談 caffe 的 conv 運算元

語言: CN / TW / HK

  一起養成寫作習慣!這是我參與「掘金日新計劃 · 4 月更文挑戰」的第17天,點選檢視活動詳情

歡迎關注我的公眾號 [極智視界],獲取我的更多筆記分享

  大家好,我是極智視界。本文主要聊一下 caffe 框架中的 conv 運算元。

   caffe 中的卷積主要有兩種實現:一種是 img2col + gemm + bias,另一種是呼叫 cudnn_conv,關於 img2col 加速卷積計算的原理可以參考我之前寫的這篇 《【模型推理】一文看懂Img2Col卷積加速演算法》,這篇文章裡寫的比較清楚了。下面進行 caffe 中 conv 實現的介紹。

1、caffe conv img2col + gemm + bias

   跟 caffe conv img2col + gemm + bias 相關的幾個 cpp 有:conv_layer.cpp、base_conv_layer.cpp、im2col.cpp、math_functions.cpp,以 forward_cpu 為例,這個時候的 conv 呼叫流程如下:

   可以看到如上呼叫過程其實是單卷積運算元,不是運算元融合的情況,把 gemm 和 bias 的計算分開了,簡單來看一下原理(下面的這組圖在我的這篇文章裡有《【模型推理】談談為什麼卷積加速更喜歡 NHWC Layout》):

   原理過程很清晰,下面來看實現,邏輯按照上面的呼叫圖。

1.1 conv forward

  這是 caffe conv 運算元最頂層實現,主要看 include/caffe/layers/conv_layer.hpp 頭 和 src/caffe/layers/conv_layer.cpp 實現。讓我們先來看一下頭:

```c++ /// conv_layer.hpp

template class ConvolutionLayer : public BaseConvolutionLayer { public: explicit ConvolutionLayer(const LayerParameter& param) : BaseConvolutionLayer(param) {}

virtual inline const char* type() const { return "Convolution"; }

protected: virtual void Forward_cpu(const vector*>& bottom, const vector*>& top); virtual void Forward_gpu(const vector*>& bottom, const vector*>& top); virtual void Backward_cpu(const vector*>& top, const vector& propagate_down, const vector*>& bottom); virtual void Backward_gpu(const vector*>& top, const vector& propagate_down, const vector*>& bottom); virtual inline bool reverse_dimensions() { return false; } virtual void compute_output_shape(); }; ```

  可以看到前向推理幾個關鍵的函式:Forward_cpu、Forward_gpu。先看 Forward_cpu 的實現:

c++ template <typename Dtype> void ConvolutionLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) { const Dtype* weight = this->blobs_[0]->cpu_data(); for (int i = 0; i < bottom.size(); ++i) { const Dtype* bottom_data = bottom[i]->cpu_data(); Dtype* top_data = top[i]->mutable_cpu_data(); for (int n = 0; n < this->num_; ++n) { /// compute gemm this->forward_cpu_gemm(bottom_data + n * this->bottom_dim_, weight, top_data + n * this->top_dim_); /// compute bias if (this->bias_term_) { /// 有 bias 的話 const Dtype* bias = this->blobs_[1]->cpu_data(); this->forward_cpu_bias(top_data + n * this->top_dim_, bias); } } } }

  這裡兩個關鍵的操作是 this->forward_cpu_gemm、this->forward_cpu_bias。關於 this 指標有一個比較經典的回答,當你進入一個房子後,你可以看見桌子、椅子、電視等,但是房子這個整體你卻看不完整,對於一個類的例項來說,你可以看到它的成員函式、成員變數,但是例項本身呢?this 是一個指標,它時刻指向你這個示例本身。所以,拿 this->forward_cpu_gemm 來說,咋一看找不到 forward_cpu_gemm 的實現在哪裡,結合對 this 指標的理解,然後再看一下ConvolutionLayer 類的宣告:

c++ class ConvolutionLayer : public BaseConvolutionLayer<Dtype> {...}

  可以看到 ConvolutionLayer 類為 BaseConvolutionLayer 類的派生類,繼承方式為 public,公有繼承時基類的公有成員和保護成員被繼承到派生類中仍作為派生類的公有成員和保護成員,看到這裡可能就豁然開朗了,若 BaseConvolutionLayer 類中有 forward_cpu_gemm,這樣就捋順了。來看下 BaseConvolutionLayer 的宣告:

```c++ template class BaseConvolutionLayer : public Layer { public: ...

protected: void forward_cpu_gemm(const Dtype input, const Dtype weights, Dtype output, bool skip_im2col = false); void forward_cpu_bias(Dtype output, const Dtype* bias); ... } ```

   可以看到如我們所願,確實有 forward_cpu_gemm、forward_cpu_bias,這裡已經跳轉到了 include/caffe/layers/base_conv_layer.hpp,這個我們下面再看。

   以上我們看了 conv_forward_cpu 的實現,再來看看 conv_forward_gpu 的實現:

c++ template <typename Dtype> void ConvolutionLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) { const Dtype* weight = this->blobs_[0]->gpu_data(); for (int i = 0; i < bottom.size(); ++i) { const Dtype* bottom_data = bottom[i]->gpu_data(); Dtype* top_data = top[i]->mutable_gpu_data(); for (int n = 0; n < this->num_; ++n) { this->forward_gpu_gemm(bottom_data + n * this->bottom_dim_, weight, top_data + n * this->top_dim_); if (this->bias_term_) { const Dtype* bias = this->blobs_[1]->gpu_data(); this->forward_gpu_bias(top_data + n * this->top_dim_, bias); } } } }

   已經封裝的和 conv_forward_cpu 看起來差不多了,主要是 this->forward_gpu_gemm、this->forward_gpu_bias,這個同樣指向了 include/caffe/layers/base_conv_layer.hpp,下一節見分曉。

  在 ConvolutionLayer 類的宣告中注意到還有個介面 compute_output_shape,這個函式主要用於計算輸出特徵的 shape,其實現如下:

c++ template <typename Dtype> void ConvolutionLayer<Dtype>::compute_output_shape() { const int* kernel_shape_data = this->kernel_shape_.cpu_data(); const int* stride_data = this->stride_.cpu_data(); const int* pad_data = this->pad_.cpu_data(); const int* dilation_data = this->dilation_.cpu_data(); this->output_shape_.clear(); for (int i = 0; i < this->num_spatial_axes_; ++i) { // i + 1 to skip channel axis const int input_dim = this->input_shape(i + 1); const int kernel_extent = dilation_data[i] * (kernel_shape_data[i] - 1) + 1; const int output_dim = (input_dim + 2 * pad_data[i] - kernel_extent) / stride_data[i] + 1; this->output_shape_.push_back(output_dim); } }

   以上主要通過一些卷積運算元的引數來計算卷積後的尺寸,包括 kernel_size、stride、pad、dilation(空洞卷積的配置)。


1.2 gemm forward

  這裡我們主要來看一下上面提到的 this->forward_cpu_gemm:

c++ template <typename Dtype> void BaseConvolutionLayer<Dtype>::forward_cpu_gemm(const Dtype* input, const Dtype* weights, Dtype* output, bool skip_im2col) { const Dtype* col_buff = input; if (!is_1x1_) { if (!skip_im2col) { conv_im2col_cpu(input, col_buffer_.mutable_cpu_data()); } col_buff = col_buffer_.cpu_data(); } for (int g = 0; g < group_; ++g) { caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, conv_out_channels_ / group_, conv_out_spatial_dim_, kernel_dim_, (Dtype)1., weights + weight_offset_ * g, col_buff + col_offset_ * g, (Dtype)0., output + output_offset_ * g); } }

   可以看到 forward_cpu_gemm 有兩個過程,先做 im2col 的展開,然後做 gemm 矩陣乘,這比較理所當然。看程式碼還可以知道,除 1 x 1 卷積 和 手動指定不要做 im2col 外,其餘都需要做 im2col,其中 1 x 1 卷積是沒有必要做展開,也比較好理解。下面順序來看這兩個過程,首先是 im2col。看下 conv_im2col_cpu 的宣告:

c++ inline void conv_im2col_cpu(const Dtype* data, Dtype* col_buff) { if (!force_nd_im2col_ && num_spatial_axes_ == 2) { im2col_cpu(data, conv_in_channels_, conv_input_shape_.cpu_data()[1], conv_input_shape_.cpu_data()[2], kernel_shape_.cpu_data()[0], kernel_shape_.cpu_data()[1], pad_.cpu_data()[0], pad_.cpu_data()[1], stride_.cpu_data()[0], stride_.cpu_data()[1], dilation_.cpu_data()[0], dilation_.cpu_data()[1], col_buff); } else { im2col_nd_cpu(data, num_spatial_axes_, conv_input_shape_.cpu_data(), col_buffer_shape_.data(), kernel_shape_.cpu_data(), pad_.cpu_data(), stride_.cpu_data(), dilation_.cpu_data(), col_buff); } }

  主要看 im2col_cpu() 的實現,實現在 src/caffe/util/im2col.cpp,下面結合圖例進行說明。

  如下圖,黑實線為 feature map,虛線為填充 0,紅色虛線框為 3 x 3 卷積核作用區域,步長和 dilation 都為 1。

  把如上卷積過程展開後得到如下示意圖:

  其中 output_rows 的 for 迴圈對應上圖中的藍色箭頭範圍,output_col 的 for 迴圈對應上圖中的橙色半框範圍,用程式碼實現如下:

c++ template <typename Dtype> void im2col_cpu(const Dtype* data_im, const int channels, const int height, const int width, const int kernel_h, const int kernel_w, const int pad_h, const int pad_w, const int stride_h, const int stride_w, const int dilation_h, const int dilation_w, Dtype* data_col) { const int output_h = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1; const int output_w = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1; const int channel_size = height * width; for (int channel = channels; channel--; data_im += channel_size) { for (int kernel_row = 0; kernel_row < kernel_h; kernel_row++) { for (int kernel_col = 0; kernel_col < kernel_w; kernel_col++) { int input_row = -pad_h + kernel_row * dilation_h; for (int output_rows = output_h; output_rows; output_rows--) { if (!is_a_ge_zero_and_a_lt_b(input_row, height)) { for (int output_cols = output_w; output_cols; output_cols--) { *(data_col++) = 0; } } else { int input_col = -pad_w + kernel_col * dilation_w; for (int output_col = output_w; output_col; output_col--) { if (is_a_ge_zero_and_a_lt_b(input_col, width)) { *(data_col++) = data_im[input_row * width + input_col]; } else { *(data_col++) = 0; } input_col += stride_w; } } input_row += stride_h; } } } } }

   如上就完成了 im2col 的過程,下面進入 gemm,主要由 caffe_cpu_gemm() 介面實現,這個實現在 src/caffe/util/math_functions.cpp:

c++ template<> void caffe_cpu_gemm<float>(const CBLAS_TRANSPOSE TransA, const CBLAS_TRANSPOSE TransB, const int M, const int N, const int K, const float alpha, const float* A, const float* B, const float beta, float* C) { int lda = (TransA == CblasNoTrans) ? K : M; int ldb = (TransB == CblasNoTrans) ? N : K; cblas_sgemm(CblasRowMajor, TransA, TransB, M, N, K, alpha, A, lda, B, ldb, beta, C, N); }

   這裡主要呼叫了 Intel 數學核心庫 MKL 中的 cblas_sgemm 來實現高效能的矩陣乘,函式介面為:

c++ void cblas_sgemm(const CBLAS_LAYOUT Layout, const CBLAS_TRANSPOSE transa, const CBLAS_TRANSPOSE transb, const MKL_INT m, const MKL_INT n, const MKL_INT k, const float alpha, const float *a, const MKL_INT lda, const float *b, const MKL_INT ldb, const float beta, float *c, const MKL_INT ldc);

   計算方式為 c := alpha * op(a) * op(b) + beta * c,引數:

  • Layeout:表示二維矩陣儲存是按行優先(CblasRowMajor)還是列優先(CblasColMajor), C++ 裡是行優先儲存的,fortran 是列優先儲存資料;
  • transa、transb:可為 CblasNoTrans、CblasTrans、CblasConjTrans;
  • m:矩陣 a 和 c 的行數;
  • n:矩陣 b 和 c 的列數;
  • k:矩陣 a 的列數,矩陣 c 的行數;
  • lda:行優先 & 不轉置時, lda >= max(1, k);行優先 & 轉置時,lda >= max(1, m);
  • ldb:行優先 & 不轉置時, ldb x k 的矩陣, b 矩陣左上角包含 n x k 的 B 矩陣;行優先 & 轉置時,ldb x n 的矩陣, b 矩陣左上角包含 k x n 的 B 矩陣;
  • ldc:行優先時, ldc >= max(1, n);

   再來看一下 this->forward_gpu_gemm:

c++ template <typename Dtype> void BaseConvolutionLayer<Dtype>::forward_gpu_gemm(const Dtype* input, const Dtype* weights, Dtype* output, bool skip_im2col) { const Dtype* col_buff = input; if (!is_1x1_) { if (!skip_im2col) { conv_im2col_gpu(input, col_buffer_.mutable_gpu_data()); } col_buff = col_buffer_.gpu_data(); } for (int g = 0; g < group_; ++g) { caffe_gpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, conv_out_channels_ / group_, conv_out_spatial_dim_, kernel_dim_, (Dtype)1., weights + weight_offset_ * g, col_buff + col_offset_ * g, (Dtype)0., output + output_offset_ * g); } }

  同樣看一下 conv_im2col_gpu 宣告:

c++ inline void conv_im2col_gpu(const Dtype* data, Dtype* col_buff) { if (!force_nd_im2col_ && num_spatial_axes_ == 2) { im2col_gpu(data, conv_in_channels_, conv_input_shape_.cpu_data()[1], conv_input_shape_.cpu_data()[2], kernel_shape_.cpu_data()[0], kernel_shape_.cpu_data()[1], pad_.cpu_data()[0], pad_.cpu_data()[1], stride_.cpu_data()[0], stride_.cpu_data()[1], dilation_.cpu_data()[0], dilation_.cpu_data()[1], col_buff); } else { im2col_nd_gpu(data, num_spatial_axes_, num_kernels_im2col_, conv_input_shape_.gpu_data(), col_buffer_.gpu_shape(), kernel_shape_.gpu_data(), pad_.gpu_data(), stride_.gpu_data(), dilation_.gpu_data(), col_buff); } }

   來看 im2col_gpu(),這個實現在 src/caffe/util/im2col.cu:

c++ template <typename Dtype> void im2col_gpu(const Dtype* data_im, const int channels, const int height, const int width, const int kernel_h, const int kernel_w, const int pad_h, const int pad_w, const int stride_h, const int stride_w, const int dilation_h, const int dilation_w, Dtype* data_col) { // We are going to launch channels * height_col * width_col kernels, each // kernel responsible for copying a single-channel grid. int height_col = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1; int width_col = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1; int num_kernels = channels * height_col * width_col; // NOLINT_NEXT_LINE(whitespace/operators) im2col_gpu_kernel<Dtype><<<CAFFE_GET_BLOCKS(num_kernels), CAFFE_CUDA_NUM_THREADS>>>( num_kernels, data_im, height, width, kernel_h, kernel_w, pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w, height_col, width_col, data_col); CUDA_POST_KERNEL_CHECK; }

  來看 im2col_gpu_kernel 這個核函式:

c++ template <typename Dtype> __global__ void im2col_gpu_kernel(const int n, const Dtype* data_im, const int height, const int width, const int kernel_h, const int kernel_w, const int pad_h, const int pad_w, const int stride_h, const int stride_w, const int dilation_h, const int dilation_w, const int height_col, const int width_col, Dtype* data_col) { CUDA_KERNEL_LOOP(index, n) { const int h_index = index / width_col; const int h_col = h_index % height_col; const int w_col = index % width_col; const int c_im = h_index / height_col; const int c_col = c_im * kernel_h * kernel_w; const int h_offset = h_col * stride_h - pad_h; const int w_offset = w_col * stride_w - pad_w; Dtype* data_col_ptr = data_col; data_col_ptr += (c_col * height_col + h_col) * width_col + w_col; const Dtype* data_im_ptr = data_im; data_im_ptr += (c_im * height + h_offset) * width + w_offset; for (int i = 0; i < kernel_h; ++i) { for (int j = 0; j < kernel_w; ++j) { int h_im = h_offset + i * dilation_h; int w_im = w_offset + j * dilation_w; *data_col_ptr = (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) ? data_im_ptr[i * dilation_h * width + j * dilation_w] : 0; data_col_ptr += height_col * width_col; } } } }

   以上完成了 gpu 版的 im2col 的過程,接下來是 gpu 版的 gemm。

c++ template <typename Dtype> void caffe_gpu_gemm(const CBLAS_TRANSPOSE TransA, const CBLAS_TRANSPOSE TransB, const int M, const int N, const int K, const Dtype alpha, const Dtype* A, const Dtype* B, const Dtype beta, Dtype* C);

  其實現在 src/caffe/util/math_functions.cu 中:

c++ template <> void caffe_gpu_gemm<float>(const CBLAS_TRANSPOSE TransA, const CBLAS_TRANSPOSE TransB, const int M, const int N, const int K, const float alpha, const float* A, const float* B, const float beta, float* C) { // Note that cublas follows fortran order. int lda = (TransA == CblasNoTrans) ? K : M; int ldb = (TransB == CblasNoTrans) ? N : K; cublasOperation_t cuTransA = (TransA == CblasNoTrans) ? CUBLAS_OP_N : CUBLAS_OP_T; cublasOperation_t cuTransB = (TransB == CblasNoTrans) ? CUBLAS_OP_N : CUBLAS_OP_T; CUBLAS_CHECK(cublasSgemm(Caffe::cublas_handle(), cuTransB, cuTransA, N, M, K, &alpha, B, ldb, A, lda, &beta, C, N)); }

  可以看到和 gemm_cpu 不同的是 gemm_gpu 呼叫了 cublas 中的高效矩陣乘方法 cublasSgemm。

1.3 bias forward

   一般的卷積帶有偏置 bias,上面已經談了談 im2col + gemm 的實現了,下面再說一下 bias 的實現。

   forward_cpu_bias 的實現在 src/caffe/layers/base_conv_layer.cpp:

c++ template <typename Dtype> void BaseConvolutionLayer<Dtype>::forward_cpu_bias(Dtype* output, const Dtype* bias) { caffe_cpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, num_output_, out_spatial_dim_, 1, (Dtype)1., bias, bias_multiplier_.cpu_data(), (Dtype)1., output); }

  同樣是呼叫了 caffe_cpu_gemm,前面介紹過這個接口裡又呼叫了 cblas_sgemm,區別是通過傳參控制 bias 計算為 矩陣加,而 gemm 為 矩陣乘,所以不多說了。

  forward_gpu_bias 的實現在 src/caffe/layers/base_conv_layer.cpp:

c++ template <typename Dtype> void BaseConvolutionLayer<Dtype>::forward_gpu_bias(Dtype* output, const Dtype* bias) { caffe_gpu_gemm<Dtype>(CblasNoTrans, CblasNoTrans, num_output_, out_spatial_dim_, 1, (Dtype)1., bias, bias_multiplier_.gpu_data(), (Dtype)1., output); }

  裡面呼叫了 caffe_gpu_gemm,最終呼叫了 cublas 中的矩陣運算操作。

   以上介紹了 caffe conv img2col + gemm + bias 的整個實現過程。下面介紹一下 caffe cudnn_conv 的實現。

2、caffe cudnn conv

   caffe 框架原始碼中除了有 img2col + gemm +bias 的卷積高效實現方法,還有 cudnn conv 的實現可供選擇,下面對 caffe cudnn conv 進行介紹。

  還是先來看一下頭,在 include/caffe/layers/cudnn_conv_layer.hpp:

```c++

ifdef USE_CUDNN

template class CuDNNConvolutionLayer : public ConvolutionLayer { public: explicit CuDNNConvolutionLayer(const LayerParameter& param) : ConvolutionLayer(param), handles_setup_(false) {} virtual void LayerSetUp(const vector*>& bottom, const vector*>& top); virtual void Reshape(const vector*>& bottom, const vector*>& top); virtual ~CuDNNConvolutionLayer();

protected: virtual void Forward_gpu(const vector*>& bottom, const vector*>& top); virtual void Backward_gpu(const vector*>& top, const vector& propagate_down, const vector*>& bottom);

bool handles_setup_; cudnnHandle_t handle_; cudaStream_t stream_;

// algorithms for forward and backwards convolutions cudnnConvolutionFwdAlgo_t fwd_algo_; cudnnConvolutionBwdFilterAlgo_t bwd_filter_algo_; cudnnConvolutionBwdDataAlgo_t *bwd_data_algo_;

vector bottom_descs_, top_descs_; cudnnTensorDescriptor_t bias_desc_; cudnnFilterDescriptor_t filter_desc_; vector conv_descs_; int bottom_offset_, top_offset_, bias_offset_;

size_t workspace_fwd_sizes_; size_t workspace_bwd_data_sizes_; size_t workspace_bwd_filter_sizes_; size_t workspaceSizeInBytes; // size of underlying storage void workspaceData; // underlying storage void **workspace; // aliases into workspaceData };

endif

```

  可以看到 cudnn conv 必須是在構建與 caffe_cudnn 的基礎上,這裡理所當然的,所以也只有 Forward_gpu,沒有 cpu 版本了。其實現在 src/caffe/layers/cudnn_conv_layer.cu:

```c++ template void CuDNNConvolutionLayer::Forward_gpu( const vector*>& bottom, const vector*>& top) { const Dtype weight = this->blobs_[0]->gpu_data(); for (int i = 0; i < bottom.size(); ++i) { const Dtype bottom_data = bottom[i]->gpu_data(); Dtype* top_data = top[i]->mutable_gpu_data();

// Forward through cuDNN in parallel over groups.
for (int g = 0; g < this->group_; g++) {
  // Filters.
  CUDNN_CHECK(cudnnConvolutionForward(handle_[g],
        cudnn::dataType<Dtype>::one,
        bottom_descs_[i], bottom_data + bottom_offset_ * g,
        filter_desc_, weight + this->weight_offset_ * g,
        conv_descs_[i],
        fwd_algo_[i], workspace[g], workspace_fwd_sizes_[i],
        cudnn::dataType<Dtype>::zero,
        top_descs_[i], top_data + top_offset_ * g));

  // Bias.
  if (this->bias_term_) {
    const Dtype* bias_data = this->blobs_[1]->gpu_data();
    CUDNN_CHECK(cudnnAddTensor(handle_[g],
          cudnn::dataType<Dtype>::one,
          bias_desc_, bias_data + bias_offset_ * g,
          cudnn::dataType<Dtype>::one,
          top_descs_[i], top_data + top_offset_ * g));
  }
}

// Synchronize the work across groups, each of which went into its own
// stream, by launching an empty kernel into the default (null) stream.
// NOLINT_NEXT_LINE(whitespace/operators)
sync_conv_groups<<<1, 1>>>();

} } ```

  由於卷積的計算方式為 y = w * x + b,所以這裡的實現也是有兩個過程,即 乘w 和 加b。其中乘 w 呼叫了 cudnn 中的 cudnnConvolutionForward 介面來實現;加 b 呼叫了 cudnn 中的 cudnnAddTensor 介面來實現。


   小結一下,以上介紹了 caffe 中關於 conv 前向推理的兩個實現方式,可以看到到最後其實都是呼叫了高效能的推理庫:

  (1) img2col + gemm + bias:先做 img2col 展開為平行計算做基礎,gemm 和 bias 的計算在 cpu 時最後都呼叫了 MKL 中的 cblas_sgemm 介面,gpu 時最後都呼叫了 cublas 中的 cublasSgemm 介面;

  (2) cudnn conv:權重矩陣乘呼叫了 cudnn 中的 cudnnConvolutionForward 介面來實現,加偏置呼叫了 cudnn 中的 cudnnAddTensor 來實現。


   關於 caffe conv 的介紹就這樣了,有問題歡迎討論~ 以上分享了 caffe 的 conv 運算元,希望我的分享能對你的學習有一點幫助。


 【公眾號傳送】

《【模型推理】談談 caffe 的 conv 運算元》


logo_show.gif