caffe源码分析--poolinger_layer.cpp

本文详细解析了Caffe框架中的池化层实现原理,包括最大池化与平均池化的算法流程,并探讨了激活函数在Caffe中的使用方式。

摘要生成于 C知道 ,由 DeepSeek-R1 满血版支持, 前往体验 >


caffe源码分析--poolinger_layer.cpp


对于采样层,cafffe里实现了最大采样和平均采样的算法。

最大采样,给定一个扫描窗口,找最大值,

平均采样,扫描窗口内所有值的平均值。


其实对于caffe的实现一直有个疑问,

就是每一层貌似没有绑定一个激活函数?

看ufldl教程,感觉激活函数是必要存在的。

这怎么解释呢?


看到源码中,看到一些激活函数,比如sigmoid_layer.cpp和sigmoid_layer.cu。

也就是说,激活函数作为layer层面来实现了。当然,还有tanh_layer和relu_layer。


那,这个意思是说,让我们建立网络的时候更加随意,可自由搭配激活函数吗?

但是,我看了caffe自带的那些例子,貌似很少见到用了激活函数layer的,顶多看到用了relu_layer,其他的没见过。

这意思是说,激活函数不重要吗?真是费解啊。


// Copyright 2013 Yangqing Jia

#include <algorithm>
#include <cfloat>
#include <vector>

#include "caffe/layer.hpp"
#include "caffe/vision_layers.hpp"
#include "caffe/util/math_functions.hpp"

using std::max;
using std::min;

namespace caffe {

template <typename Dtype>
void PoolingLayer<Dtype>::SetUp(const vector<Blob<Dtype>*>& bottom,
      vector<Blob<Dtype>*>* top) {
  CHECK_EQ(bottom.size(), 1) << "PoolingLayer takes a single blob as input.";
  CHECK_EQ(top->size(), 1) << "PoolingLayer takes a single blob as output.";
  KSIZE_ = this->layer_param_.kernelsize();//核大小
  STRIDE_ = this->layer_param_.stride();//步长
  CHANNELS_ = bottom[0]->channels();//通道
  HEIGHT_ = bottom[0]->height();//高
  WIDTH_ = bottom[0]->width();//宽
  POOLED_HEIGHT_ = static_cast<int>(
      ceil(static_cast<float>(HEIGHT_ - KSIZE_) / STRIDE_)) + 1;//计算采样之后的高
  POOLED_WIDTH_ = static_cast<int>(
      ceil(static_cast<float>(WIDTH_ - KSIZE_) / STRIDE_)) + 1;//计算采样之后的宽
  (*top)[0]->Reshape(bottom[0]->num(), CHANNELS_, POOLED_HEIGHT_,//采样之后大小
      POOLED_WIDTH_);
  // If stochastic pooling, we will initialize the random index part.
  if (this->layer_param_.pool() == LayerParameter_PoolMethod_STOCHASTIC) {
    rand_idx_.Reshape(bottom[0]->num(), CHANNELS_, POOLED_HEIGHT_,
      POOLED_WIDTH_);
  }
}

// TODO(Yangqing): Is there a faster way to do pooling in the channel-first
// case?
template <typename Dtype>
void PoolingLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
      vector<Blob<Dtype>*>* top) {
  const Dtype* bottom_data = bottom[0]->cpu_data();//采样层输入
  Dtype* top_data = (*top)[0]->mutable_cpu_data();//采样层输出
  // Different pooling methods. We explicitly do the switch outside the for
  // loop to save time, although this results in more codes.
  int top_count = (*top)[0]->count();
  switch (this->layer_param_.pool()) {
  case LayerParameter_PoolMethod_MAX://最大采样方法
    // Initialize
    for (int i = 0; i < top_count; ++i) {
      top_data[i] = -FLT_MAX;
    }
    // The main loop
    for (int n = 0; n < bottom[0]->num(); ++n) {
      for (int c = 0; c < CHANNELS_; ++c) {
        for (int ph = 0; ph < POOLED_HEIGHT_; ++ph) {
          for (int pw = 0; pw < POOLED_WIDTH_; ++pw) {
            int hstart = ph * STRIDE_;
            int wstart = pw * STRIDE_;
            int hend = min(hstart + KSIZE_, HEIGHT_);
            int wend = min(wstart + KSIZE_, WIDTH_);
            for (int h = hstart; h < hend; ++h) {//找出核范围内最大
              for (int w = wstart; w < wend; ++w) {
                top_data[ph * POOLED_WIDTH_ + pw] =
                  max(top_data[ph * POOLED_WIDTH_ + pw],
                      bottom_data[h * WIDTH_ + w]);
              }
            }
          }
        }
        // compute offset 指针移动到下一个channel。注意代码这里的位置。采样是针对每个channel的。
        bottom_data += bottom[0]->offset(0, 1);
        top_data += (*top)[0]->offset(0, 1);
      }
    }
    break;
  case LayerParameter_PoolMethod_AVE:
    for (int i = 0; i < top_count; ++i) {
      top_data[i] = 0;
    }
    // The main loop
    for (int n = 0; n < bottom[0]->num(); ++n) {
      for (int c = 0; c < CHANNELS_; ++c) {
        for (int ph = 0; ph < POOLED_HEIGHT_; ++ph) {
          for (int pw = 0; pw < POOLED_WIDTH_; ++pw) {
            int hstart = ph * STRIDE_;
            int wstart = pw * STRIDE_;
            int hend = min(hstart + KSIZE_, HEIGHT_);
            int wend = min(wstart + KSIZE_, WIDTH_);
            for (int h = hstart; h < hend; ++h) {//核范围内算平均
              for (int w = wstart; w < wend; ++w) {
                top_data[ph * POOLED_WIDTH_ + pw] +=
                    bottom_data[h * WIDTH_ + w];
              }
            }
            top_data[ph * POOLED_WIDTH_ + pw] /=
                (hend - hstart) * (wend - wstart);
          }
        }
        // compute offset
        bottom_data += bottom[0]->offset(0, 1);
        top_data += (*top)[0]->offset(0, 1);
      }
    }
    break;
  case LayerParameter_PoolMethod_STOCHASTIC://这种算法这里未实现
    NOT_IMPLEMENTED;
    break;
  default:
    LOG(FATAL) << "Unknown pooling method.";
  }
}

template <typename Dtype>
Dtype PoolingLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
      const bool propagate_down, vector<Blob<Dtype>*>* bottom) {
  if (!propagate_down) {
    return Dtype(0.);
  }
  const Dtype* top_diff = top[0]->cpu_diff();
  const Dtype* top_data = top[0]->cpu_data();
  const Dtype* bottom_data = (*bottom)[0]->cpu_data();
  Dtype* bottom_diff = (*bottom)[0]->mutable_cpu_diff();
  // Different pooling methods. We explicitly do the switch outside the for
  // loop to save time, although this results in more codes.
  memset(bottom_diff, 0, (*bottom)[0]->count() * sizeof(Dtype));
  switch (this->layer_param_.pool()) {
  case LayerParameter_PoolMethod_MAX:
    // The main loop
    for (int n = 0; n < top[0]->num(); ++n) {
      for (int c = 0; c < CHANNELS_; ++c) {
        for (int ph = 0; ph < POOLED_HEIGHT_; ++ph) {
          for (int pw = 0; pw < POOLED_WIDTH_; ++pw) {
            int hstart = ph * STRIDE_;
            int wstart = pw * STRIDE_;
            int hend = min(hstart + KSIZE_, HEIGHT_);
            int wend = min(wstart + KSIZE_, WIDTH_);
            for (int h = hstart; h < hend; ++h) {
              for (int w = wstart; w < wend; ++w) {
                bottom_diff[h * WIDTH_ + w] +=//采样层输出的残传播给输入。由于是最大采样方法,输出存的都是输入范围内最大的值,所以残差传播的时候也只有范围内最大的值受影响
                    top_diff[ph * POOLED_WIDTH_ + pw] *
                    (bottom_data[h * WIDTH_ + w] ==
                        top_data[ph * POOLED_WIDTH_ + pw]);
              }
            }
          }
        }
        // offset  移动到下一个channel
        bottom_data += (*bottom)[0]->offset(0, 1);
        top_data += top[0]->offset(0, 1);
        bottom_diff += (*bottom)[0]->offset(0, 1);
        top_diff += top[0]->offset(0, 1);
      }
    }
    break;
  case LayerParameter_PoolMethod_AVE:
    // The main loop
    for (int n = 0; n < top[0]->num(); ++n) {
      for (int c = 0; c < CHANNELS_; ++c) {
        for (int ph = 0; ph < POOLED_HEIGHT_; ++ph) {
          for (int pw = 0; pw < POOLED_WIDTH_; ++pw) {
            int hstart = ph * STRIDE_;
            int wstart = pw * STRIDE_;
            int hend = min(hstart + KSIZE_, HEIGHT_);
            int wend = min(wstart + KSIZE_, WIDTH_);
            int poolsize = (hend - hstart) * (wend - wstart);
            for (int h = hstart; h < hend; ++h) {
              for (int w = wstart; w < wend; ++w) {
                bottom_diff[h * WIDTH_ + w] +=//采样层输出的残差传播给输入,由于是平均采样,所以权重都是1 / poolsize。
                  top_diff[ph * POOLED_WIDTH_ + pw] / poolsize;
              }
            }
          }
        }
        // offset
        bottom_data += (*bottom)[0]->offset(0, 1);
        top_data += top[0]->offset(0, 1);
        bottom_diff += (*bottom)[0]->offset(0, 1);
        top_diff += top[0]->offset(0, 1);
      }
    }
    break;
  case LayerParameter_PoolMethod_STOCHASTIC:
    NOT_IMPLEMENTED;
    break;
  default:
    LOG(FATAL) << "Unknown pooling method.";
  }
  return Dtype(0.);
}


INSTANTIATE_CLASS(PoolingLayer);


}  // namespace caffe

本文作者:linger
本文链接:http://blog.youkuaiyun.com/lingerlanlan/article/details/38294169

-- Build files have been written to: /root/Workspace/openpose/build/caffe/src/openpose_lib-build [ 75%] Performing build step for 'openpose_lib' [ 1%] Running C++/Python protocol buffer compiler on /root/Workspace/openpose/3rdparty/caffe/src/caffe/proto/caffe.proto [ 1%] Building CXX object src/caffe/CMakeFiles/caffeproto.dir/__/__/include/caffe/proto/caffe.pb.cc.o [ 1%] Linking CXX static library ../../lib/libcaffeproto.a [ 1%] Built target caffeproto [ 1%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/util/cuda_compile_1_generated_math_functions.cu.o [ 1%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/layers/cuda_compile_1_generated_absval_layer.cu.o [ 2%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/layers/cuda_compile_1_generated_base_data_layer.cu.o [ 2%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/layers/cuda_compile_1_generated_accuracy_layer.cu.o [ 2%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/layers/cuda_compile_1_generated_batch_norm_layer.cu.o [ 2%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/layers/cuda_compile_1_generated_batch_reindex_layer.cu.o [ 4%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/layers/cuda_compile_1_generated_bnll_layer.cu.o [ 4%] Building NVCC (Device) object src/caffe/CMakeFiles/cuda_compile_1.dir/layers/cuda_compile_1_generated_bias_layer.cu.o In file included from /root/Workspace/openpose/3rdparty/caffe/src/caffe/util/math_functions.cu:1: /usr/local/cuda-11.8/include/math_functions.h:54:2: warning: #warning "math_functions.h is an internal header file and must not be used directly. This file will be removed in a future CUDA release. Please use cuda_runtime_api.h or cuda_runtime.h instead." [-Wcpp] 54 | #warning "math_functions.h is an internal header file and must not be used directly. This file will be removed in a future CUDA release.
最新发布
03-15
<think>好的,我现在要解决用户提到的CUDA编译时出现的math_functions.h警告问题,并寻找替代方案,比如使用cuda_runtime_api.h或cuda_runtime.h。首先,我需要回忆一下CUDA头文件的结构和用途。 用户提到的警告可能是因为math_functions.h被标记为内部头文件,不建议直接引用。根据CUDA的文档,开发人员应该使用更高级别的头文件,比如cuda_runtime.h或cuda_runtime_api.h,这些头文件会自动包含必要的内部文件。所以,用户需要检查他们的代码中是否有直接包含math_functions.h的情况,并替换为官方推荐的头文件。 接下来,用户可能还需要检查编译环境是否正确配置。例如,确保CUDA安装路径正确,并在编译时用-I指定包含路径,如-I/usr/local/cuda/include。此外,交叉编译时可能需要安装对应的工具链,如用户提到的aarch64-linux-gnu工具链,这可能涉及到引用[4]中的内容。 另外,用户提供的引用[3]中的代码涉及到CUDA和OpenGL的互操作,这可能与他们的项目相关,但当前问题主要集中在头文件引用上。需要确认在链接阶段是否正确链接了CUDA库,比如使用-lcudart。 此外,用户提到的警告信息可能类似于“math_functions.h is an internal header file and should not be used directly.”这时,建议用户检查所有源文件中是否有直接包含该头文件的地方,并替换为cuda_runtime.h,因为后者已经包含了必要的数学函数声明。 还可能要考虑代码中使用的数学函数是否有替代方案。例如,CUDA的运行时API提供了一些数学函数,或者使用标准库函数,但需要注意设备代码和主机代码的区别。例如,在设备代码中应使用__sin、__cos等内置函数,而不是标准库的sin、cos,这些可能在math_functions.h中定义,但现在应该通过包含正确的头文件来访问。 最后,用户需要确保编译命令正确,包括正确的CUDA版本和路径设置。如果用户遇到找不到头文件的错误,可能需要检查CUDA安装是否正确,或者参考引用[2]中的解决方案,安装相应的开发包,如nvidia-cuda-dev。</think>### 解决CUDA编译时`math_functions.h`警告的替代方案 #### 1. **问题分析** CUDA的`math_functions.h`是内部头文件,直接包含会触发警告: ```bash warning: math_functions.h is an internal header file and should not be used directly. ``` 官方建议使用`cuda_runtime.h`或`cuda_runtime_api.h`替代,因为它们已封装了必要的数学函数接口[^1]。 --- #### 2. **替代方案** ##### (1) 使用`cuda_runtime.h`替代 - **作用**:包含CUDA运行时API的核心接口和数学函数(如`sin`、`cos`、`sqrt`等)。 - **修改代码**: ```c // 原代码(不推荐) #include <math_functions.h> // 修改为(推荐) #include <cuda_runtime.h> ``` ##### (2) 使用标准库函数(仅限主机代码) - **场景**:在主机代码(CPU端)中使用标准数学函数。 - **示例**: ```c #include <cmath> // 标准C++数学库 // 主机代码中使用std::sin、std::sqrt等 ``` ##### (3) 使用CUDA内置函数(设备端代码) - **场景**:在设备端(GPU核函数)中直接调用CUDA内置数学函数。 - **示例**: ```cuda __global__ void kernel(float* data) { int idx = threadIdx.x; data[idx] = __sinf(data[idx]); // 使用CUDA内置单精度正弦函数 } ``` --- #### 3. **编译配置调整** ##### (1) 确保CUDA头文件路径正确 在编译命令中指定CUDA头文件路径: ```bash nvcc -I/usr/local/cuda/include your_code.cu -o output ``` ##### (2) 链接CUDA运行时库 添加`-lcudart`参数: ```bash nvcc your_code.cu -o output -lcudart ``` ##### (3) 交叉编译配置(如适用) 若需交叉编译(如ARM架构),安装对应工具链: ```bash sudo apt-get install gcc-aarch64-linux-gnu g++-aarch64-linux-gnu # 引用[4] ``` --- #### 4. **验证代码示例** ```cuda #include <cuda_runtime.h> // 替代math_functions.h #include <cmath> // 主机端标准数学库 __global__ void compute(float* d_data) { int idx = blockIdx.x * blockDim.x + threadIdx.x; d_data[idx] = __expf(d_data[idx]); // 设备端使用内置指数函数 } int main() { float* d_data; cudaMalloc(&d_data, 100 * sizeof(float)); compute<<<1, 100>>>(d_data); cudaFree(d_data); return 0; } ``` --- #### 5. **常见问题解决** - **头文件缺失错误**:确保已安装CUDA开发包(如`nvidia-cuda-dev`)[^2]。 - **OpenGL互操作问题**:若涉及CUDA与OpenGL交互,需正确注册缓冲区(参考引用[3])。 ---
评论 6
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值