explicit 修饰constructor

比较区别：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

// 无explicit
class foo{
public:
    foo(int n){ num_ = n; }
private:
    int num_;
};
// 有explicit
class bar{
public:
    explicit bar(int n){ num_ = n; }
private:
    int num_;
};

int main() {
    foo f1(2);   // correct
    foo* f2 = new foo(2);  // correct
    foo f3 = 2;  // correct 
    

    bar b1(3);   // correct
    bar* b2 = new bar(3); // correct
    bar b3 = 3; // incorrect
}

bar b3 = 3含义是，3的类型int正好是bar单参数构造函数的参数类型int，这时候编译器就自动调用这个构造函数，创建一个bar的对象。这在explicit修饰下是不可行的，只能显式地以3为参数传给构造函数。

shared_ptr<>

理解shared_ptr可以这样理解：
对于std::shared_ptr<string> ptr(new string)，这里有两个对象，“我”和“你”（均为指针）。“我”是ptr，而“你”对应了new string。这时“我”和“你”就有了关系，“我”管理“你”。当new string即“你”不存在时，我不管理任何人，即我管理一个空指针。

定义并初始化

定义一个空共享指针，指向类型为int的对象, owns no pointer（“我”没有管理任何指针）, 由于ptr1是空指针，所以为ptr1所指向的对象赋值是违法的：segmentation fault：

1
2
3
4
5

std::shared_ptr<string> ptr1; 
if (!ptr1)
    cout<<"ptr1==NULL"<<endl;
// 违法，因为ptr1没有指向任何对象，所以给对象赋值是逻辑错误
*ptr1 = "string ptr1";

初始化一个shared_ptr ptr2，非空指针，指向“”。此后，ptr2 就可以像string* 类型的指针一样使用。ptr2 可以管理使用new 关键字分配出来的指针。（new 出来的一定是个指针）

1
2
3
4

// ptr2 管理一个指向string的指针
std::shared_ptr<string> ptr2(new string);
*ptr2 = "string ptr2";
cout<<*ptr2<<endl;

常用的初始化方法:

1
2
3
4
5
6
7

// 一般的初始化一个shared_ptr方法
std::shared_ptr<string> ptr3(new string("string ptr3"));
cout<<*ptr3<<endl;

// 推荐的初始化一个shared_ptr
std::shared_ptr<string> ptr4 = make_shared<string>("string ptr4");
cout<<*ptr4<<endl;

成员方法

ptr.get()方法，它返回一个指针，这个指针指向ptr所指向对象：

1
2

string* p = ptr4.get();
cout<<*p<<endl;

ptr.reset()方法，改变ptr所指向内容。这里有个好例子：

1
2
3
4
5
6
7
8
9
10
11
12
13
14

std::shared_ptr<int> sp;  // empty 

// 管理一个指针
sp.reset (new int);       
*sp=10;
std::cout << *sp << '\n';

// 删除所管理的指针, 干礼一个新指针
sp.reset (new int);       
*sp=20;
std::cout << *sp << '\n';

// 删除所管理的指针
sp.reset();

以下情况相同：

1
2
3
4
5
6

ptr4.reset(new string("string reset ptr4"));
cout<<*ptr4<<endl;

string* pp = new string("fuck you");
ptr4.reset(pp);
cout<<*ptr4<<endl;

如果某个值需要保持不变，那么你就应该说出来，这样编译器能确保这写约束不被违反。const可以放于很多位置。

1. 作用于指针

1
2
3
4
5

char greeting[] = "hello";
char* p = greeting;
const char* p = greeting;   // 1
char* const p = greeting;   // 2
const char* const p = greeting;  // 3

const的出现位置[左数右针]：

1. const *，数据是常量，表示被指向的对象是常量，而指针是可以变的。
2. * const，表示指针是常量。这个指针不能指向不同的东西，但它指向东西的数据可以改变。
3. const * const，表示指针和被指向的对象都是常量

对于在*左侧，下面两种表达意义相同，都在左侧：

1
2

const char* greeting;
char const* greeting;

2. 作用于特殊指针，iterator

对于STL中的迭代器，加入我希望迭代器所指向的东西不可被修改（即对数据只读），那么就需要一个const_iterator：

1
2
3
4
5

vector<int> vec;
// 数据const
vector<int>::const_iterator cItr = vec.begain();
*cItr = 10;   // 错，所指向的东西不能改
cItr++；      // 对，指针可改

而另一种情况，指针不能改变其指向，但是可改变其所指向数据：

1
2
3
4

// 指针const
const vector<int>::iterator itr = vec.begain();
*itr = 10;   // 对，所指向的东西可改
itr++；      // 错，指针指向不可改

3. 作用与函数

1. 函数返回值为const

   将函数返回值声明为const，这个返回值不能被修改。降低了因为用户的错误而造成的意外。

2. 函数参数为const

   除非你要改动局部变量（函数参数），否则将参数声明为const。

4. 作用与类成员函数

即const位于成员函数()和{}之间，如：

1
2
3
4
5
6
7
8
9
10
11
12
13

class foo {
public:
    inline void area(int x, int y) {
        length_ = x;
        width_ = y;
    }
    // length和width，不能被这个函数修改。
    // 限制对象成员函数 改变对象数据成员，保护了程序的健壮性和稳定性
    void print() const {cout<<"area="<<length_ * width_<<endl;}
private:
    int length_;
    int width_;
};

print()成员函数，为const函数，表示这个成员函数不能修改成员变量。所以不能将area()成员函数设为const，因为这个函数的 目的就是为成员变量赋值。

const成员函数，和非const成员函数使得我们知道一个类中哪些函数可以修改成员属性，哪些不可以。

tip:
提升c++程序效率的一个方法是以pass by referrence-to-const的方式或pass by pointer-to-const传递对象：

1
2

// 此为以成员函数
const int opt( const foo& x) const { return x.length_ * x.width_}

关于如何使用，在网上找了半天，没有一篇博客说清楚。之后才发现，下载好的caffe的examples目录中就是我想找的。通过examples，caffe的工作流程基本明了。一般由4步：

1. 数据准备： 将原始数据转化为caffe-format
2. 定义网络结构和参数，
3. 定义训练超参数solver，
4. 训练

给个caffe自带的mnist例子

1.数据准备

回到caffe的根目录后执行下面脚本：

1
2

$./data/mnist/get_mnist.sh
$./examples/mnist/create_mnist.sh

上述代码所做的事情是：先下载解压mnist数据集，后将源数据通过build/examples/mnist/convert_mnist_data.bin写入将原始mnist数据写成LMDB格式，所以会有两个文件生成：mnist_train_lmdb和mnist_test_lmdb。

2.定义模型结构和参数

我使用caffe为我们定义好的LeNet网络，它的定义在文件examples/mnist/lenet_train_test.prototxt中。定义的格式是Google Protobuff，解析.prorotxt文件的　统一规则由文件/src/caffe/proto/caffe.proto.提供。

打开文件lenet_train_test.prototxt，是lenet的网络结构：

1
2
3
4
5
6
7
8
9

name: "LeNet"
layer {
// ...layer definition...
include: { phase: TRAIN }
}
layer {
}
layer {
}

• 若干层（Layer）堆叠在一起，构成了一个网络（Net）。lenet网络每层定义的细节间附录。
• 这个文件被称作network definition protobuf file。其中top表示这层输出，bottom表示这层出入。caffe中模型的逻辑图结构与其.prototxt文件的对应。
• 网络结构可以通过可视化工具绘制（详见后）。

3.指明训练超参数：solver.prototxt

从这里查看训练超参数examples/mnist/lenet_solver.prototxt。在caffe中，设定训练超参数（非训练参数）被称作solver，这个文件被称作solver protobuf file，也是Google protobuff 文件格式。如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23

# 网络结构定义
net: "examples/mnist/lenet_train_test.prototxt"
# Test阶段的迭代次数。batch_size也是100，所以由10,000测试数据
test_iter: 100
# 每500次训练后执行1次测试
test_interval: 500
# The base learning rate, momentum and the weight decay of the network.
base_lr: 0.01
momentum: 0.9
weight_decay: 0.0005
# 设定学习率的策略
lr_policy: "inv"
gamma: 0.0001
power: 0.75
# 每100次迭代后，打印
display: 100
# 最大迭代次数
max_iter: 10000
# 每5000次记录一次快照，中间结果
snapshot: 5000
snapshot_prefix: "examples/mnist/lenet"
# 用什么计算
solver_mode: GPU

这个文件在程序中按照caffe.proto的协议解析到一个Caffe::SolverParameter 对象中，进而将解析后的参数使用到网络中。用于解析的文件是 caffe.pb.h 和 caffe.pb.cc，其位于build/src/caffe/proto/。

默认使用GPU。可以使用同目录的其他的参数配置文件（不同优化方法）：

lenet_multistep_solver.prototxt
lenet_solver_adam.prototxt
lenet_solver_rmsprop.prototxt

4.训练与测试

有了caffe-format的数据；有了网络结构；有了确定的超参数，就可以执行下面脚本开始训练：

1

./examples/mnist/train_lenet.sh

train_lenet.sh文件包含一行代码：

1
2
3
4

#!/usr/bin/env sh
set -e
 
./build/tools/caffe train --solver=examples/mnist/lenet_solver.prototxt $@

使用caffe命令（这个caffe是tools中的命令，这是caffe提供的命令行接口cmdcaffe）， 指明是trian，指明训练超参数配置，即使用哪个*solver.prototxt文件。(而网络结构超参数的地址，在*solver.prototxt中指明）

训练结束后，在examples/mnist/路径下生成两个文件lenet_iter_10000.caffemodel和lenet_iter_10000.solverstate。

1. .caffemodel文件是最终模型，二进制文件。
2. .solverstate是保存的训练状态（snapshot），可以从此继续开始训练。

训练结束

lenet_iter_10000.caffemodel 是binary protobuf 文件，就是训练10000次的最终模型，用于对真实样本的推理。

MNIST数据集是小数据集，所以使用GPU的效果并不好，原因是GPU的计算核心与存储的通讯开销。所以对于复杂模型和数据集GPU的速度提升是显著的。

使用预训练的model对新样本预测

训练结束后生成的.caffemodel是可以直接拿来使用的模型（其实就是所有训练好的参数），用作推理。如何使用？

如果要使用预训练的模型，从这里下载预训练的模型caffe model index，如bvlc_reference_caffenet.caffemodel， 并且将其放到caffe-root/models/bvlc_reference_caffenet/中。

使用预训练模型：

1
2
3
4
5

# 指明 -weights 关键字，提供预训练模型
./build/tools/caffe \
train \
--solver examples/finetuning_on_flickr_style/solver.prototxt \
--weights models/bvlc_reference_caffenet/bvlc_reference_caffenet.caffemodel

上面有问题

总结

生成制定格式数据{数据&模型之数据}：

1
2

$./data/mnist/get_mnist.sh
$./examples/mnist/create_mnist.sh

在GPU上训练{数据&模型之模型}：

1
2

./build/tools/caffe train \
--solver=examples/mnist/lenet_solver.prototxt -gpu 0

关键字为train，需要提供*solver.prototxt文件，和网络结构文件，后者的路径在*solver.prototxt文件中指明。有了结构参数和超参数，就可以进行学习。

提供快照文件.solverstate可以接着训练：

1
2
3

./build/tools/caffe train \
-solver examples/mnist/lenet_solver.prototxt \
-snapshot examples/mnist/lenet_iter_5000.solverstate

附录 lenet模型定义

说明，以下是模型结构参数文件的内容，其包括了TRAIN网络和TEST网络，所需的每个Layer。其实在运行Log中打印的两个网络才更直观。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168

name: "LeNet"
layer {
  name: "mnist"
  type: "Data"
  top: "data"
  top: "label"
  include {
    phase: TRAIN  # 本层只用于训练过程中
  }
  transform_param {  # 数据变换使用的方所因子
    scale: 0.00390625
  }
  data_param {     # 数据层参数，包括二进制数据的路径
    source: "examples/mnist/mnist_train_lmdb"
    batch_size: 64
    backend: LMDB
  }
}
layer {
  name: "mnist"
  type: "Data"
  top: "data"
  top: "label"
  include {
    phase: TEST   # 同样的数据层，只用于测试阶段
  }
  transform_param {
    scale: 0.00390625
  }
  data_param {
    source: "examples/mnist/mnist_test_lmdb"
    batch_size: 100
    backend: LMDB
  }
}
layer {
  name: "conv1"
  type: "Convolution"
  bottom: "data"
  top: "conv1"
  param {
    lr_mult: 1  # 权值学习速率因子，1表示保持与全局参数一致
  }
  param {
    lr_mult: 2  # bias学习因子，2表示是全局的2倍
  }
  convolution_param {  # 卷基层参数
    num_output: 20
    kernel_size: 5
    stride: 1
    weight_filler {
      type: "xavier"  # 权值使用xavier方法初始化？
    }
    bias_filler {
      type: "constant"  # bias初始化为常量？
    }
  }
}
layer {
  name: "pool1"
  type: "Pooling"
  bottom: "conv1"
  top: "pool1"
  pooling_param {
    pool: MAX
    kernel_size: 2
    stride: 2
  }
}
layer {
  name: "conv2"
  type: "Convolution"
  bottom: "pool1"
  top: "conv2"
  param {
    lr_mult: 1
  }
  param {
    lr_mult: 2
  }
  convolution_param {
    num_output: 50
    kernel_size: 5
    stride: 1
    weight_filler {
      type: "xavier"
    }
    bias_filler {
      type: "constant"
    }
  }
}
layer {
  name: "pool2"
  type: "Pooling"
  bottom: "conv2"
  top: "pool2"
  pooling_param {
    pool: MAX
    kernel_size: 2
    stride: 2
  }
}
layer {
  name: "ip1"
  type: "InnerProduct"
  bottom: "pool2"
  top: "ip1"
  param {
    lr_mult: 1
  }
  param {
    lr_mult: 2
  }
  inner_product_param { # 全连接层InnerProduct
    num_output: 500   
    weight_filler {
      type: "xavier"
    }
    bias_filler {
      type: "constant"
    }
  }
}
layer {   # 非现行变换层
  name: "relu1"
  type: "ReLU"
  bottom: "ip1"
  top: "ip1"
}
layer {  
  name: "ip2"
  type: "InnerProduct"
  bottom: "ip1"
  top: "ip2"
  param {
    lr_mult: 1
  }
  param {
    lr_mult: 2
  }
  inner_product_param {
    num_output: 10
    weight_filler {
      type: "xavier"
    }
    bias_filler {
      type: "constant"
    }
  }
}
layer {
  name: "accuracy"   # 计算accuracy 只用于Test阶段
  type: "Accuracy"
  bottom: "ip2"
  bottom: "label"
  top: "accuracy"
  include {
    phase: TEST
  }
}
layer {
  name: "loss"   # loss层
  type: "SoftmaxWithLoss"
  bottom: "ip2"
  bottom: "label"
  top: "loss"
}

Layer和Net

caffe中的每个层对象，都是一个模型的基本计算单元。一个层对象包括 “…filters, pool, take inner products, apply nonlinearities like rectified-linear and sigmoid and other elementwise transformations, normalize, load data, and compute losses like softmax and hinge”。见官方文档

几乎所有的层类都会继承自Layer类：class caffe::Layer<Dtype> 这些层类必须实现一个 Forward(), 它接受一个input Blob（bottom）计算输出的Blob（top）。还要实现一个 Backward()，它使用给定其输出Blob的误差梯度，计算相对于其输入Blob的误差梯度。（实现反向传播）

每个层对象一定包含三个基本方法：setup()，forward()，backward()。

1. setup()：一些计算前的操作，用来初始化层。
2. forward()：前向计算。多个Layer构成一个Net，一个Net中连续的前向传播由Net::Forward()实现。
3. backward()：官方文档解释的很好：“given the gradient w.r.t. the top output, compute the gradient w.r.t. to the input and send to the bottom. A layer with parameters computes the gradient w.r.t. to its parameters and stores it internally.”。其中w.r.t.表示“with respect to”，中文表示“对…（求梯度等操作）”。这就是反向对参数求梯度的过程。当多个层链接到一起，就会形成反向传播的链条（对应了链式法则）。这个链式法则在Net::Backward()中实现

前向计算和后向计算都有CPU和GPU版本的实现。实现自己的层也是不难的，只需要定义好上述三个关键方法。

这里是管方文档中所有的层。

这里是caffe中的前向传播和后向传播的官方描述。

一个Net由多个Layer构成，如下面一个计算图（有向无环图）：

其中蓝色矩形表示layer，里边的小写表示层的name，大写表示层的type；黄色多边形表示在图中游走的数据Blob，数据移动方向从下向上，下为bottom，上为top。这个图对应的.prototxt如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27

name: "LogReg"
layer {
  name: "mnist"
  type: "Data"
  top: "data"   #输出
  top: "label"  #输出
  data_param {
    source: "input_leveldb"
    batch_size: 64
  }
}
layer {
  name: "ip"
  type: "InnerProduct"
  bottom: "data"  #输入
  top: "ip"       #输出
  inner_product_param {
    num_output: 2
  }
}
layer {
  name: "loss"
  type: "SoftmaxWithLoss"
  bottom: "ip"     #输入
  bottom: "label"   #输入
  top: "loss"
}

这个Net的初始化通过执行Net::Init()，其作用是生成Blob和Layers，并调用Layer的setup函数。终端的输出信息表示caffe还有记录的工作被执行。

用Caffe::mode()和Caffe::set_mode()指定使用CPU或GPU执行.CPU和GPU的切换是无缝的，与模型的定义是无关的。

模型的定义在机器中是以protocol buffer schema (prototxt)形式存在的。对应训练好的模型是以binary protocol buffer (binaryproto)存在与磁盘的，就是相应路径中的.caffemodel文件。

模型在caffe中的格式。prototxt由caffe.proto解析。细节见这里本地路径：caffe-ROOT/src/caffe/proto/caffe.proto。所使用的技术是Google Protocol Buffer通过科学上网查看。

为什么要使用protocal buffer来格式化模型？
因为，你哪里见过一边砌墙一边烧砖的。使用 protobuff 将整个流程分开，实现每一个组件，之后在将所需的组件拼接起来，最后训练。

Blob

Blob为模型提供数据和数据载体，在模型的正向反向传播中移动。它还提供了CPU和GPU间的同步机制。Blob中的数据可以是一批图像，模型参数，中间计算结果。在Blob中不同类型的数据，其大小是不同的。

1. 图像数据Blob，其中的数据是4维的：（N, K, H, W）分别表示batch size，channel数量，长和宽。当一个Blob对象中的数据变化时，变化优先级从右向左。所以索引为（n, k, h, w）的数值在物理存储中的索引是((n×K + k)×H + h)×W + w。

2. 对于非图像的数据，使用2D Blob，（N, D），此时通常与InnerProductLayer一同使用。

3. 对于参数Blob，其维度要具体问题具体分析了。比如对于conv层，由64个conv kernel，一个kernel的大小是11×11，输入通道数为3，那么此情况的参数Blob大小为（96,3,11,11）。而对于全连接层或上述的InnerProductLayer，如果输出1000维，输入1024维，那么此情况的参数Blob大小为（1000,1024）。

Blob数据一旦定以好后，caffe的模块化就可以做剩下的工作了。

Blob细节

通常我们关心Blob中的数值（一般的数据，比如一批输入图像）和梯度值。所以一个Blob由两块儿存储：data和diff。

更进一步，Blob中的数值有两种访问的方式，cont和mutable，前者数值不变，后者数值可改变。比如

1
2

const Dtype* cpu_data() const;
Dtype* mutable_cpu_data();

同样的，对于cpu_diff和gpu_data, gpu_diff 也有两种访问方式。

Blob为什么要如此设计？

官方文档给出如下解释：

1. Blob使用SyncedMem类实现CPU与GPU间的数据拷贝，目的是隐藏延时。

2. 如果不想更改数值，那么始终使用const方式调用，而且永远不要将指针存储在自己的对象中。每次处理Blob时，都要调用函数来获取指针，因为SyncedMem类需要这个函数来确定何时复制数据。

3. 在实际中，调用设备内核来执行GPU计算的同时，CPU将数据从磁盘加载到Blob，并将Blob转移到下一层，这是GPU和CPU延时隐藏的基本技巧。因为大部分层都有GPU实现，所以所有的中间数据和梯度都将保留在GPU的内存中。

其他原因如，Blob与其他深度学习框架的数据结构相对应，tensorflow中的tensor，cuda-convnet中的NVMatrix，等，这使得框架之间的转换更容易。

Blob中数据在CPU和GPU之间拷贝行为

假设数据在CPU（Host）上被初始化，存在于一个Blob中。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

// 定义一个const访问指针，一个mutable访问指针
const Dtype* foo;   
Dtype* bar;
// GPU上没有数据数据
foo = blob.gpu_data(); // data copied cpu->gpu.
// 上下两句间没有任何操作，所以下一句并不会发生数据拷贝
foo = blob.cpu_data(); // no data copied since both have up-to-date contents.
bar = blob.mutable_gpu_data(); // 一样，没有数据拷贝.
//
// 对数据进行操作...
//
bar = blob.mutable_gpu_data(); // 当前位于GPU上，没有数据拷贝到gpu.   ？？？？？？
foo = blob.cpu_data(); // 因为数据被操作，所以数据拷贝从GPU到CPU发生。
foo = blob.gpu_data(); // 因为GPU和CPU都是最新数据，所以没有数据拷贝发生。

// ？？？？？？
bar = blob.mutable_cpu_data(); // still no data copied.
bar = blob.mutable_gpu_data(); // data copied cpu->gpu.
bar = blob.mutable_cpu_data(); // data copied gpu->cpu.

1. 安装依赖库

执行以下命令：

1
2
3
4

$sudo apt-get install libprotobuf-dev libleveldb-dev libsnappy-dev libopencv-dev libhdf5-serial-dev protobuf-compiler
$sudo apt-get install --no-install-recommends libboost-all-dev
$sudo apt-get install libgflags-dev libgoogle-glog-dev liblmdb-dev
$sudo apt-get install libatlas-base-dev

安装第一个就出问题：

1

sudo apt-get install libprotobuf-dev

返回错误：

由于解决问题后，将terminal的输出被clear掉了，所以用以下相似的问题来说明解决过程

1
2

The following packages have unmet dependencies:
 libxml2-dev : Depends: libxml2 (= 2.7.8.dfsg-5.1ubuntu4) but 2.7.8.dfsg-5.1ubuntu4.6 is to be installed

解决方法：

1. 安装aptitude:

1
2

$sudo apt-get install aptitude
$aptitude why-not libxml2

2. 找到已经安装的对应的libxml2包：

1

$dpkg -l | grep libxml2

3. 删除libxml2，并删除其他所有依赖包，--force-all参数不能少：

1

$sudo dpkg --purge --force-all libxml2

4. 更正错误：

1

$sudo apt-get -f install

5. 安装欠缺的包

1

$sudo apt-get install libxml2-dev

遇到相同的问题，将libxml2-dev和libxml2，替换成自己问题中的对应库。

2. 安装NVIDIA驱动和CUDA，openCV

早已安装完成

3. 下载caffe源码

从这里下载到home/XXX。

修改Makefile.config文件。先复制一份：

1

cp Makefile.config.example Makefile.config

后修改配置文件Makefile.config，下面是我的修改：

1. hdf5

   将hdf5的路径添加到INCLUDE_DIRS和LIBRARY_DIRS之后：

   1 2	INCLUDE_DIRS := $(PYTHON_INCLUDE) /usr/local/include /usr/include/hdf5/serial LIBRARY_DIRS := $(PYTHON_LIB) /usr/local/lib /usr/lib /usr/lib/x86_64-linux-gnu/hdf5/serial

2. sm_version

   由于我的机器的CUDA版本是9.0，所以要将下面的sm_20, sm_21注释掉，低版本的一些指令对于高版本的不适用。如下：

   1 2 3 4 5 6 7 8 9

   CUDA_ARCH := -gencode arch=compute_30,code=sm_30 \ -gencode arch=compute_35,code=sm_35 \ -gencode arch=compute_50,code=sm_50 \ -gencode arch=compute_52,code=sm_52 \ -gencode arch=compute_60,code=sm_60 \ -gencode arch=compute_61,code=sm_61 \ -gencode arch=compute_61,code=compute_61 \ # -gencode arch=compute_20,code=sm_20 \ # -gencode arch=compute_20,code=sm_21 \

3. OpenCV版本

   去掉下面的注释，表示使用openCV 版本3：

   1	OPENCV_VERSION := 3

4. CUDA路径

   检查机器CUDA路径是否正确：

   1	CUDA_DIR := /usr/local/cuda

4. 编译测试

配置完成之后编译，-j4表示用4个核心执行任务：

1
2
3

$sudo make all -j4 
$sudo make test -j4
$sudo make runtest -j4

期待不出错。

如果有错，比如编译期间又修改了Makefile.config文件，返回类似错误：

1

undefined reference to `cv::imread(cv::String const&, int)

可是OpenCV已经修改过了，此时的处理情况：

将编译了一半的build文件夹删除，后重新编译：

1
2

$sudo rm -rf ./build/*
$sudo make all -j4

其中build/文件夹是编译后得到的，

最终期望的是sudo make runtest -j4运行无误，如果terminal最后返回：

表示caffe 安装配置成功完成（如果要使用python接口，还要其他操作，我使用CXX）。

5. 编译python接口

caffe使用python2.

1. 第O步，使用linux自带python2，安装依赖库：

   1 2	sudo apt-get install python-pip sudo apt-get install python-numpy

   在caffe根目录的python文件夹下，有一个requirements.txt的清单文件，上面列出了需要的依赖库，需要按照这个清单安装：

   1 2 3

   sudo apt-get install gfortran cd ./python sudo cat python/requirements.txt|xargs -L 1 sudo pip install -i https://pypi.tuna.tsinghua.edu.cn/simple

   检查

   1	sudo pip install -r python/requirements.txt

   编译成功的会显示Requirement already satisfied，没有编译成功的会继续安装。

2. 第一步，执行命令：

   1	$ make pycaffe -j4

3. 第二步，在~/.bashrc或/etc/profile中添加caffe中的python 路径：

   1	export PYTHONPATH=/home/junhui/caffe-master/python:$PYTHONPATH

4. 第三步，source使之生效：

   1 2	$ source /etc/profile $ source ~/.bashrc

   现在在python脚本中就可以import caffe了。

可能的出错：

1. numpy/arrayobject.h: missing

   1 2 3 4 5 6 7 8

   junhui@gnome:~/caffe$ sudo make pytest –j4 CXX/LD -o python/caffe/_caffe.so python/caffe/_caffe.cpp python/caffe/_caffe.cpp:10:31: fatal error: numpy/arrayobject.h: No such file or directory #include <numpy/arrayobject.h> ^ compilation terminated. Makefile:517: recipe for target 'python/caffe/_caffe.so' failed make: *** [python/caffe/_caffe.so] Error 1

   解决方法：

   在Makefile.config文件中的这个地方：

   1 2 3

   # We need to be able to find Python.h and numpy/arrayobject.h. PYTHON_INCLUDE := /usr/include/python2.7 \ /usr/lib/python2.7/dist-packages/numpy/core/include

   添加一行：

   1 2 3 4

   # We need to be able to find Python.h and numpy/arrayobject.h. PYTHON_INCLUDE := /usr/include/python2.7 \ /usr/lib/python2.7/dist-packages/numpy/core/include \ /usr/local/lib/python2.7/dist-packages/numpy/core/include

   再次编译成功

   1	$ make pycaffe -j4

2. Python.h: No suchfile or directory

   caffe 使用python2，需要在./bashrc中添加python2的路径：

   如果不是使用anaconda2中的python2，则添加如下第一句；如果使用用的anaconda2中的python2，则添加如下第二句

   1 2	export CPLUS_INCLUDE_PATH=/usr/include/python2.7:$CPLUS_INCLUDE_PATH export CPLUS_INCLUDE_PATH=/home/junhui/anaconda2/include/python2.7/:$CPLUS_INCLUDE_PATHort

3. protobuf出错

   使用anaconda时出的错：

   1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18

   CXX .build_release/src/caffe/proto/caffe.pb.cc In file included from .build_release/src/caffe/proto/caffe.pb.cc:5:0: .build_release/src/caffe/proto/caffe.pb.h:12:2: error: #error This file was generated by a newer version of protoc which is #error This file was generated by a newer version of protoc which is ^ .build_release/src/caffe/proto/caffe.pb.h:13:2: error: #error incompatible with your Protocol Buffer headers. Please update #error incompatible with your Protocol Buffer headers. Please update ^ .build_release/src/caffe/proto/caffe.pb.h:14:2: error: #error your headers. #error your headers. ^ In file included from .build_release/src/caffe/proto/caffe.pb.cc:5:0: .build_release/src/caffe/proto/caffe.pb.h:26:55: fatal error: google/protobuf/generated_enum_reflection.h: No such file or directory #include <google/protobuf/generated_enum_reflection.h> compilation terminated. make: *** [.build_release/src/caffe/proto/caffe.pb.o] Error 1

   解决方法，将机器上的anaconda3删除，使用linux自带python2，并安装以来库。（其实可以使用anaconda2）.

4. make clean

   如果make中间出错，改错后，先make clean，后重新make

import caffe 成功后测试pycaffe

在mnist数据集上运行lenet

1
2
3

sudo ./data/mnist/get_mnist.sh
sudo ./examples/mnist/create_mnist.sh
sudo ./examples/mnist/train_lenet.sh

从源文件到可执行文件

• 预处理（pre-processing）E
  从.c到.i。编译器将C源代码中的包含的头文件如stdio.h编译进来，替换宏。
• 编译（Compiling）S
  从.i到.s。gcc首先要检查代码的规范性、是否有语法错误等，以确定代码的实际要做的工作，在检查无误后，gcc把代码翻译成汇编语言。
• 汇编（Assembling） c
  把编译阶段生成的.s文件转成二进制目标代码.o文件。
• 链接 （Linking）
  链接到库中，生成可执行文件。

绘制这里的图(https://blog.csdn.net/Meteor_s/article/details/85208589)

这里更详细

1
2
3
4

gcc -E hello.c -o hello.i
gcc –S hello.i –o hello.s
gcc –c hello.s –o hello.o
gcc hello.o –o hello

生成.o文件，可以重定向，方便其他应用使用。

一条命令从源文件到可执行文件：

1

gcc hello.c –o hello

1 虚析构函数

补充内容（tip 7）：

任何时候都应该为待多态性质的基类（父类）声明virtual析构函数。如果一个class含有任何virual函数，就一定要有一个virtual析构函数。

而如果这个class的设计不是作为基类使用，或者不准备让这个class具备多态性，就不该声明virtual析构函数。

含有多态性质的base class的设计目的是为了通过base class的接口来使用derived class。所以通常使用base class类的指针指向derived class的对象。

并不是所有的base class都是为了多态用途，比如STL容器的设计不是用作base class的。

virtual同名函数的目的是允许derived class的实现可以客制化。

2 C++帮你实现了的函数

（tip 5）

当定义了一个class，又没有实现任何成员函数时，编译器会为这个class声明4个函数：

• 一个default构造函数
• 一个copy构造函数
• 一个析构函数
• 一个copy操作符

这些函数都是public和inline的。如下：

1
2
3
4
5
6
7
8
9
10
11
12
13

class Myclass{
public:
    Myclass(){...}
    Myclass(const Myclass& rhs){...}
    ~Myclass(){...}
    Myclass& operator=(const Myclass& rhs){...}
};

// 实例化这个class，并使用
Myclass e1;       //default 构造函数
Myclass e2(e1);   //copy构造函数
e2=e1;            //copy符号操作符
                  //析构函数

其中

1. 编译器给出的析构函数是个non-virtual函数，除非这个class的base class含有virtual析构函数。此时这个class的虚拟性来自其base class。
2. 对于构造函数，如果我声明了一个构造函数，那么编译器不再创建default构造函数
3. copy构造函数，和copy操作符只是单纯的将来源对象的每一个non-static成员变量拷贝到目标对象。所以:

   1	Myclass e2(e1)

   是使用e1的成员变量来初始化e2的成员变量。

3 为classes实现赋值操作符

使用赋值操作时，通常可以写成如下形式：

1
2

int x, y, z;
x=y=z=2;

上述实际上的赋值行为是

1

x=(y=(z=2));

其行为是将2赋值给z，再将z赋值给y，最后将y赋值给x。

为了实现上述的连锁赋值，赋值操作符必须返回一个引用指向操作符的左侧：

1
2
3
4
5
6
7

class Myclass{
public:
    Myclass& operator=(const Myclass& rhs){
        ...
        return* this;
    }
};

同样的上述的形式也适用于所有相关的操作符，如+=，-=，×=。

这是一条协议，被所有内置类型和标准库函数共同遵守，所以，遵守它吧。

4 将成员变量声明为private

为什么要把classes的成员变量声明为private：因为这体现了封装。

类成员变量应该只被这个类的成员方法可见，protected成员变量同public成员变量一样缺乏封装性。

从封装的角度讲，访问权限只有提供封装（private）和不提供封装两种。

5 使用non-member non-friend函数而非再定义一个member函数

假设由一个类含有若干个清理函数，只是清理的对象不同：

1
2
3
4
5
6
7
8

class WebBrowser{
public:
    ...
    void clearHistory();
    void clearCache();
    void clearCookies();
    ...
};

其实很多时候，用户想要一次性执行这些动作，如何一次性完成，两种方案：

1. 给这个class中添加一个member函数clearAll()，它调用上述三个清理函数：

   1 2 3 4 5 6

   class WebBrowser{ public: ... void clearAll(); ... }；

2. 使用一个non-member，non-friend函数：

   1 2 3 4 5

   void clearBrowser(WebBrowser& wb){ wb.clearHistory(); wb.clearCache(); wb.clearCookies(); }

   pro tip告诉你使用后者。后者更体系那封装性。后者保护了类的包裹性，进而有较低的编译依赖，增强了类的可延展性。

对于封装性，可以访问private成员变量的只有member函数和friend函数。上述两种方式提供了相同的功能，而后者提供了较强的封装性，因为它不能增加访问private变量的能力。

P.S. friend函数和member函数对于类private成员的访问能力相同。

在C++ 中自然的做法是将类WebBrowser和函数clearBrowser()放置于同一个namespace中：

1
2
3
4
5

namespace Web{
    class WebBrowser{...};
    void clearBrowser(WebBrowser& wb);
    ...
}

namespace 可以跨越多个源码文件，而class不能。就是说相同的namespace中可以有多个头文件，这正是C++标准库的组织方式：std::vector, std::sort, std::map, … C++ 将不同的部分，不同container放在不同的头文件中，每个声明了std的某部分功能，这样当使用vector，时只用include <vector>，而不必将所有的std内容include进来。就是说，用户只用对所用的部分进行编译。而class必须整体定义，不能分割。

想要扩充这个namespace的能力，只需要在这个namespace中添加更多的non-member，non-friend函数。即增强了功能，又没有破坏封装性，由没有增加编译依赖。

敲黑板两个角度：

• 增强封装性
• 减小编译依赖

卷积操作应用于许多领域，而其特点：计算量大，每个输出元素的计算都是相互独立的。这两个特点是并行计算期望处理的。卷积中对于边界的处理影响着分块算法的效率。

一维卷积基本形式

假设一维数组N是被操作对象，一维数组M是卷积核，一维数组P为卷积结果。假如对于边界元素（幽灵元素）的处理是赋值为0，也就是说幽灵元素与对应的M元素的乘积为0.

若每个thread负责P数组中的一个元素，那么一维卷积的基本形式如下图：

一个thread得到结果数组P中一个元素。实现：

1
2
3
4
5
6
7
8
9
10
11
12

__global__ void conv1D(float* N, float* M, float* P, 
                        int kernelSize, int NLength){
    int tid = threadIdx.x + blockDim.x * blockIdx.x;
    float Pvalue = 0.0f;
    int start = tid-kernelSize/2;
    for (int j=0; j<kernelSize; j++){
        if(start + j >= 0 && start + j <= NLength){
            Pvalue += N[start + j] * M[j]; 	
        }
    }
    P[tid] = Pvalue;
}

一般卷积核的长度是奇数，这样计算过程是对称的。

上述实现的缺点：

• 代码中这句 Pvalue += N[start + j] * M[j]包含两次对Global 的访存，和两次算术运算（一个乘法，一个加法）。运算访存比仅为1.0。性能有限。
• 对边界的处理出现Divergence。

Constant Memory优化的一维卷积

卷积操作有三个特点：

1. 考虑到卷积计算的过程，以及实际中的卷积操作，比如在Google Inception Net分组卷积，MobileNet深度可分离卷积，ResNet中的残差学习单元，中的描述，卷积核大小都是3x3, 1x1, 1x7, 7x1，5x5很小的卷积核。
2. 卷积核的内容是不变的。
3. 所有的threads都访问卷积核（假设就一个卷积核），并且访问卷积核中元素的顺序是一样的。

根据这些特点结合Ray-Tracing这篇笔记中Constant Memory
描述的Constant Memory的特性。考虑将卷积核放入Constant Memory。

这里补充一点，Constant Memory中的内容对于所有Blocks可见。使用.totalConstMem可以查询其大小。

与Ray-Tracing中的使用一样：

1
2
3
4
5
6
7
8
9

#define MAX_KERNEL_LENGTH 10;
__constant__ float M[MAX_KERNEL_LENGTH];

...

// 在main函数中使用cudaMemcpy的特殊版本函数
cudaMemcpyToSymbol(M, h_M, kernel_length*sizeof(float));

...

kernel函数与基本形式一样，除了参数列表中不需要再传入卷积核数组M。

虽然Constant Memory在Global中，但是cuda知道Constant Memory中的内容是不变的。所以直接将其放入高速缓存中（L1缓存）。

补充:
高速缓存的层次结构。从DRAM中访问一个变量需要数百上千个时钟周期，而且处理器处理数据要比访存快得多。DRAM的长延迟和有限带宽成了现代处理器的性能瓶颈，称为存储器墙。现在的处理器都会有片内的高速缓存，来减少DRAM的访存次数。

缓存一致性:
使用高速缓存的一个重要设计问题是缓存一致性，当一个或多个处理器核心试图修改缓存中的数据时，问题出现了。每一个处理器核心有自己的L1缓存，如果这个被修改了，而其他核心的L1缓存不变，缓存中的内容就不一致了。在并行处理器中处理缓存一致性开销很大（就相当于实现全局线程同步开销很大一样）。所以GPU中没有缓存一致性的机制。

将Constant memory放到L1高速缓存中。WHY

• Constant memory中的内容在kernel执行期间不会被改变，因此没有缓存一致性的问题干扰。所以硬件可以直接将Constant memory放到L1高速缓存中。
• 在处理器的缓存设计中通常优化了线程的广播。所以当一个warp中的线程访问同一个Constant memory中的（相同的）变量时，高速缓存能为threads所需要的数据提供巨大的带宽。

Shared memory优化的分块并行卷积

分析卷积过程，以一维卷积为例，假设卷积核的长度是5，N（Global）中有100个元素，其中每一个元素要平均被访问5次，那么就需要500次的Global memory的访存。显然，N中每个元素被重复访问了。所以考虑是用分块Shared memory来降低Global的访存。

假设N长度为15，M长度为5，block大小为4，Shared memory大小为4，输出到P。过程如下图：

<>pic2<>

每一个thread处理一个P中元素，对于每一个block，分两步操作：

1. 将数据从N中写入blocks对应的Shared 中。
2. 执行卷积计算。

而读入Shared的操作分为左中右，三部分

实现：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30

#define KERNEL_LENGTH 5;
__constant__ float M[KERNEL_LENGTH]; // 卷积核

__global__ void Convolution(float* N, float* P, int N_size){
    __shared__ float N_ds[TILE_SIZE + KERNEL_LENGTH-1];
    int tid = threadIdx.x + blockDim.x * blockIdx.x;
    int n = KERNEL_LENGTH/2;

    // 读左
    int left = (blockIdx.x -1)*blockDim.x + threadIdx.x;
    if(threadIdx.x >= blockDim.x - n){
        N_ds[threadIdx - (blockDim.x - n)] = (left<0) ? 0 : N[left];
    }
    // 读中
    N_ds[n+threadIdx.x] = N[tid];

    // 读右
    int right = (blockIdx.x + 1)*blockDim.x + threadIdx.x;
    if(threadIdx.x < n){
        N_ds[n+blockDim.x+threadIdx.x] = (left>=N_size) ? 0 : N[right];
    }

    __syncthreads();
    // 卷积操作
    float Pvalue = 0.0f;
    for(int i=0;i<KERNEL_LENGTH; i++){
        Pvalue += N_ds[threadIdx.x + i] * M[i];
    }
    P[tid] = Pvalue;
}

L2缓存了的分块并行卷积

新知
上个实现的代码复杂度集中在了将左右元素加载到Shared 中，注意，一个block的左右元素，同时又是相邻block的内部元素，因此很有可能在block #1需要这些左右元素时，它们已经由于block #2的访问而存储到了L2缓存上。这样对于左右元素的访问从代码上是访问了global，实际上却是转化为对L2缓存的访问。依然达到减少Global访存的目的。看下图：

<>pic3<> 图8-10

如此一来，shared memory 的大小变成如下，与block的大小一致了：

1

__shared__ float N_ds[TILE_SIZE];

从而Shared memory的加载就更简单了：

1
2

int tid = threadIdx.x + blockDim.x * blockIdx.x;
N_ds[threadIdx.x] = N[tid];

之后就是对于每一个thread遍历卷积核中元素，乘累加。只是对于中间元素，已经存在于自己的Shared中了，访问自己的Shared；对于左右元素，访问N，实际上是访问L2 缓存。

实现与基本形式相似：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30

#define KERNEL_LENGTH 5;
__constant__ float M[KERNEL_LENGTH]; // 卷积核

__global__ void Convolution(float* N, float* P, int N_size){
    __shared__ float N_ds[TILE_SIZE];

    int tid = threadIdx.x + blockDim.x * blockIdx.x;
    N_ds[threadIdx.x] = N[tid];

    __syncthreads();
    int this_block_start = blockIdx.x * blockDim.x;
    int next_block_start = (blockIdx.x+1) * blockDim.x;

    int n_start = tid-(KERNEL_LENGTH/2);
    float Pvalue=0.0f;
    // 遍历卷积核：
    for(int j=0;j<KERNEL_LENGTH; j++){
        int index = n_start + j;
        // 对于有效Index：
        if(index>=0 && index<=N_size){
            // 对于中间元素：
            if((index >= this_block_start) && (index < next_block_start)){
                Pvalue += N_ds[threadIdx.x + j-KERNEL_LENGTH/2]*M[j];
            }else{  // 对于左右元素：
                Pvalue += N[index]*M[j];
            }
        }
    }
    P[tid] = Pvalue;
}

敲黑板

• 所有这些技术都可以用在二维，三维卷积，只不过是对存储访问的索引计算更加复杂。
• 存在对Global的重复访问，考虑使用Shared。
• shared memory在L1缓存中。
• 考虑元素是否被缓存到L2.
• 缓存：作为名词是指在芯片内部的高速存储区域；作为动词指将其他存储空间中的内容放入高速缓存空间。[待求证]

20, 1021, 1019, 155，921，

#20 valid parenthese

• 问题描述

  对于string S，判断是否是合法的括号对儿。

  如：

  1 2	{[()]()}() 合法 {()] 不合法

• 思路

  遍历s：

  1. 如果s[i]是左括号“(”，“[”, “{”，则入栈stack.push(s[i])。

  2. 如果是右括号，如果栈为空，表示第一个就不合法，结束。否则，记录栈顶元素curr=stack.top()，后出栈顶元素stack.pop()。

  3. 根据匹配规则，得到curr应该是什么，用match表示。然后匹配：看实际curr是否与应该match相同。若相同，继续遍历，否则不合法，结束。

     关键：判断当前元素应该是什么，与实际是什么，是否一致。

• 实现

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28

  bool matchParenthese(string S){ stack<char> stack; // 遍历每一个字符 for(int i=0; i<S.size(); i++){ // 如果是'(', 入栈 if(S[i] == '(' || S[i] == '[' || S[i] == '{' ) stack.push(S[i]); // 如果是')'： else { // 如果栈开始就为空，则不合法 if(stack.size() == 0) return false; // 得到栈顶元素 “实际”是啥 char curr = stack.top(); stack.pop(); // 匹配原则 得到“应该”是啥 char match; if (S[i] == ')') match = '('; else if(S[i] == ']') match = '['; else if(S[i] == '}') match = '{'; // 这时才开始匹配 “实际”与“应该”一致否 if(curr != match) return false; } } if(stack.size() != 0) return false; return true; }

这个问题是stack 的经典应用。

#1021 remove outermost parentheses

• 问题描述

  括号对儿是和法的。对于输入，返回输出：

  1 2 3 4 5 6 7 8

  Input: "(()())(())" Ouput: "()()()" Explanation: The input string is "(()())(())", with primitive decomposition "(()())" + "(())". After removing outer parentheses of each part, this is "()()" + "()" = "()()()". Input: "()()" Output: ""

• 思路

  1. 第一步，遍历S，将每一对儿最外层的括号的右括号的Index记录在arr中，（初始化的arr中有个-1）。例：

     1 2	S: ( ( ( ) ) ) ( ( ) ( ) ) index: 0 1 2 3 4 5 6 7 8 9 10 11

     根据描述，返回arr: [-1, 5, 11]。因为0~5是第一对儿最外层括号，6~11是第二对儿最外层括号。

  2. 第二步，对arr中的前n-1个元素，取arr[i]+2和arr[i+1]-1之间的 所有对应的S中的元素。结束。

• 实现

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28

  string removeOutermostParentheses(string S){ vector<int> arr; arr.push_back(-1); // the 1st step stack<char> stack; for(int i=0;i<S.size(); i++){ if(S[i] == '(') stack.push(S[i]); else{ stack.pop(); // 当stack为空，表示一组匹配结束，记录结束的index if(stack.size()==0) arr.push_back(i); } } // the 2nd step string resS; for(int i=0; i<arr.size()-1; i++){ int start = arr[i]+2; int end = arr[i+1]-1; if(start>=end) continue; // 将outermost括号内的所有内用append到resS中 for(int j=start; j<=end; j++){ resS+=S[j]; } } return resS; }

#1019 next greater nodes in linked list

• 问题描述

  看下面这个例子：

  1 2 3 4 5 6 7 8 9 10

  Input: [2,7,4,3,5] Output: [7,0,5,5,0] 解释： 对于2，其后第一个比他大的是7 对于7，气候第一个比他大的没有，0 对于4，其后第一个比他大的是5 对于3，其后第一个比他大的是5 对于5，为最后一个元素，其后没有比他大的，0 所以返回[7,0,5,5,0]

• 思路

• 实现

#155 Min Stack

• 描述

  1	Design a stack that supports push(), pop(), top(), and retrieving the minimum element in constant time: getMin().

• 逻辑
  在这个类中，使用两个stack。一个mystack正常工作，另一个min记录栈mystack目前的最小值，保证min中的栈顶元素是最小的，而且min从栈顶到栈底元素大小递增。比如min： 底|3|2|1 ... |顶。并且始终保持min的性质始终不变。

  特点：使用两个stack，一个正常工作，引入另一个使得整个stack类有了其他功能。

• 实现

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36

  class MinStack{ private: stack<int> mystack; // 提供stack正常的功能。 stack<int> min; // 记录最小值。 public: MinStack(){}; ~MinStack(){}; void push(int x){ mystack.push(x); // 只要待处理元素小于min的top元素， // min也要push进这个元素， // 保证min的top为当前mystack的最小值 if(min.empty() || x<=min.top()) min.push(x); } int top(){ return mystack.top(); } int getMin(){ return min.top(); } bool empty(){ return mystack.empty(); } // 当mystack.top() == min.top()时，两个stack都要pop() // 保证min的top为当前mystack的最小值。 void pop(){ if(mystack.top() == min.top()) min.pop(); mystack.pop(); } };

#921 Minimum Add to Make Parentheses Valid

• 描述

  添加最少的括号，使得括号字符串有效，例子：

  1 2 3 4 5 6 7 8 9 10

  Input: "())" Output: 1 Input: "(((" Output: 3 Input: "()" Output: 0

• 逻辑

  使用栈：

  遍历S，遇到'('入站，遇到')'出站，表示有一对是匹配的，并将记录匹配对数的count+1。遍历完后，只需要：S.length()-2*count。的结果。

  还可以不使用栈。其实栈的作用只是记录'('，所以可以只是用一个int变量来记录。如此空间复杂度从O(N)变为O(1)。

  所有思路中一定有一个是适用于所有情况的，所以如果一个思路中有很多情况不能满足，思路不对。

• 实现

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34

  // use stack int minAddToMakeValid(string S){ stack<int> stack; int count=0; for(int i=0;i<S.size(); i++){ if(S[i] == '(') stack.push(S[i]); else if(S[i] == ')'){ if(!stack.empty()){ stack.pop(); count++; } } } return S.length()-2*count; } // No stack int minAddToMakeValid(string S){ int left_p=0; int count=0; for(int i=0;i<S.size(); i++){ if(S[i] == '(') left_p++; else if(S[i] == ')'){ if(left_p!=0){ left_p--;; count++; } } } return S.length()-2*count; }

1. 二维数组在存储时会被扁平化处理，所以在程序中，最好是将多为数组扁平化为等价的一维数组。CUDA对于多为数组也是按行优先存储。

2. 算法与硬件越是契合，算法的执行效率，硬件的使用率越高。所以有两个选择

   • 可以为当前算法设计特殊的处理架构，DSA（Domain Specific Architecture）如Google TPU。
   • 根据当前硬件精心设计算法。

3. 向量数据类型

   同时适用于Host和Device，通过make_<type name>来构造，如

   int2 i2 = make_int2(3,4)：i2向量含有3和4两个元素。
   float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f)：f4是含有4个元素的数组。

   访问方式如下：

   int x = i2.x; int y = i2.y;

4. CUDA程序调试和开发

   • 可以使用ssh登路远程含有CUDA enabled GPU的服务器。
   • 通常使用双GPU的系统开发CUDA程序，一个GPU负责显示，另一个负责计算，可以使用Nsight等工具。
   • 只有一个GPU时，在Linux系统中可以关闭桌面环境（释放桌面环境对GPU的占用），只在命令行中使用CUDA-gdb调试。（实际上，实验阶段不关闭桌面环境，也是可以正确执行的）

5. CUDA开发的任务

   有效的数据并行算法 + 针对GPU架构特性的优化 = 获得并行最优性能

6. OpenCL

   OpenCL使用起来繁琐，而且运行速度远远低于CUDA运行速度。OpenCL与CUDA的主要功能有着十分相似之处，一个CUDA程序员很容易掌握OpenCL编程。

7. half-warp

   截图

8. Streams-流

   任务并行（Task Parallelism），不同于在大量的数据上执行相同的任务（SIMD），而是同时执行多个不同的任务。

9. CUDA数据并行原语

   啥是原语，就是这个领域的基本操作。CUDA并行原语库（CUDA Data Parallel Primitive Library, CUDPP）. 含有并行前缀和，并行排序，并行规约等。这些原语为许多数据并行算法提供了重要基础。如果你正在编写某个复杂算法，那么CUDPP很可能已经提供了这个算法。

   CUDPP 下载地址：http://code.google.com/p/cudpp

   for more information：CUDA By Example 178页

26. 想清楚一个问题

    如果使用shared memory，每个block对应自己的shared memory，当这个block中threads的ID更新后，这个block并没有再被分配新的shared memory. 也就是说如果使用两个blocks, 两个blocks有两段shared memory, block中threads 的ID更新后, 所计算结果会写入这个thread所在block的shared memory中. thread ID变了, 但所属的block不变.

27. 理解SIMD，SIMT

    Single Instruction Multiple Data（Thread）.

    • Single：相同的操作，kernel函数只有一个
    • Instruction：kernel函数所做的事情
    • Multiple：所处理的数据量大，要拆分为一批一批
    • Data：大的数据量
    • Threads：一个SP上的大量thread超快速切换，获得延时隐藏

28. warp中的divergence

    已经知道，在一个warp中，所有的threads执行相同的指令，但是如果指令中含有条件分支语句，很大程度上会发生divergence。比如优化并行归约 中的描述：一个宿舍的6个学生可以是一个warp，今天有的想先上厕所，后吃饭，而有的不需要上厕所，此时所有的同学都会一起先上厕所，后一起吃饭。

    总之，分支发散使得性能明显下降。但是，注意divergence只发生在一个warp中。

    可以综合算法的上下文，将divergence的粒度变为32（warp的大小）的倍数，从而避免warp内的分支发散。比如：

    1 2 3 4 5 6 7 8 9 10

    __global__ void kernel(float* c){ int tid = threadIdx.x + blockDim.x * blockIdx.x; float a=0; float b=0; if (tid%2 == 0){ c[tid] = 100; }else{ c[tid] = 200; } }

    偶数id的thread把100写入偶数编号的地址，奇数id的thread将200写入奇数编号的地址。会发生分支发散。如果将分支的粒度定为32，则没有了divergence [代码中tid%32 == 0有问题]：

    1 2 3 4 5 6 7 8 9 10

    __global__ void kernel(float* c){ int tid = threadIdx.x + blockDim.x * blockIdx.x; float a=0; float b=0; if ( tid%32 == 0 ){ c[tid] = 100; }else{ c[tid] = 200; } }

    但是两者的结果的是不同的，实际中，需要根据算法上下文考虑结合此方法。

    使用brand_efficiency指标来查看divergence的情况：

    $ nvprof --mereics brand_efficiency ./out, 但是CUDA编译器会进行优化，将短的，有条件的代码段的断定指令取代了分支指令。所以，会看到虽然代码中有分支，却显示分支效率100%。

41. 数据预读取?

    在一次读取global memory的操作和使用这个数据之间，插入独立于以上数据的操作，可以隐藏访问延迟。如：

    1 2 3

    float m = dev_a[i]; // 1. 从global memory中读取 float f = a * b; // 2. 与m无关的操作 float f2 = m*f; // 3. 使用m

    分析：在warp切换中，

42. 指令优化

    GPU中执行指令时很快速的，所以通常不用太在意指令的优化。优化顺序一般是存储优化，后执行配置优化，最后可以考虑指令优化。

    比如：除以2^n，使用移位操作>>n，以2^n求模，使用&(2^n - 1)；避免从double到float的自动转换，float a = 0.0，使用float a = 0.0f。

    还比如，两种运行时数学库函数的取舍，精度高的速度低，精度低的速度高。使用-use-fast-math编译选项后，强制将速度慢的func()转化为速度快的__func()。

    还比如循环展开。

扫描算法

Scan，是并行编程的一个重要原语，作为基本模块使用与很多不同的算法。Scan做的是什么事呢？看下例：

1
2
3
4
5

input:
[a0, a1, a2, a3, ..., an]

output:
[a0, (a0+a1), (a0+a1+a2), ..., (a0+a1+a2+...+an)]

其特点是，输出的每一个值有前缀依赖性，就是说每个输出依赖前面的所有输入。两种基本的扫描的过程如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17

template<class T>
T scan1(T* out, T* in, size_t N){
    T sum = 0;
    for int i=0;i<N;i++){
        sum += in[i];
        out[i] = sum;
    }
}

template<class T>
T scan2(T* out, T* in, size_t N){
    T sum=0; 
    for(int i=0;i<N;i++){
        out[i] = sum;
        sum += in[i];
    }
}

不多解释。

Blelloch并行扫描算法

这个方法在前元素后依赖的情况下，并行实现。Blelloch算法分了两段执行。过程如下图：

两个阶段的伪代码：

阶段一

1
2
3
4
5
6

(int* a, int N){
    int d;
    d 从0到（logN-1）:
        并行：i从0到（N-1），步长2**(d+1):
            a[i+2**(d+1) -1] += a[i+2**d -1];
}

阶段二：

1
2
3
4
5
6
7
8
9

(int* a, int N){
    a[N-1]=0
    int d;
    d 从（logN-1）到0:
        并行：i从0到（N-1），步长2**(d+1):
            t = a[i+2**d -1];
            a[i+2**d -1] = a[i+2**(d+1) -1];
            a[i+2**(d+1) -1] += t;
}

将得到的数组与原数组相加的最终结果。

每个阶段的操作虽然容易想到，但是，这个框架是很值得借鉴的。图中实例，每两行之间的操作是并行的，而且这两行之间的并行操作必须全部执行完毕，才能向下执行，所以需要同步操作，尤其是当存在线程ID更新时。