GPU笔记

​ GPU全称是Graphics Processing Unit，图形处理单元，起初主要用来绘制图像和处理图元数据，后来逐渐加入了很多其他的功能。GPU是显卡的最主要的部分，显卡除了GPU还有散热器、通讯元件、各类插槽等其他部件。GPU的主要功能有：图形绘制、物理模拟、海量计算、AI运算、其他计算等。显卡不能独立工作，需要装载在主板上，结合CPU、内存、显存、显示器等硬件，组成完整的PC机。

GPU结构

NVidia Teala架构

Tesla的微观总览图如上：

• 拥有7组TPC（Texture Processor Cluster，纹理处理簇）
• 每个TPC拥有两组SM（Stream Multiprocessor，流多处理器）
• 每个SM包括：
  • 6个SP（Streaming Processor，流处理器）
  • 2个SFU（Special Function Unit，特殊函数单元）
  • L1缓存、MT Issue（多线程指令获取）、C-Cache（常量缓存）、共享内存
• 除了TPC核心单元，还有与显存、CPU、系统内存交互的各种部件。

NVidia Fermi架构

Fermi架构如上图，它的特性如下：

• 拥有16个SM
• 每个SM：
  • 2个Warp（线程束）
  • 两组共32个Core
  • 16组加载存储单元（LD/ST）
  • 4个特殊函数单元（SFU）
• 每个Warp：
  • 16个Core
  • Warp编排器（Warp Scheduler）
  • 分发单元（Dispatch Unit）
• 每个Core：
  • 1个FPU（浮点数单元）
  • 1个ALU（逻辑运算单元）

NVidia Maxwell架构

采用了Maxwell的GM204，拥有4个GPC，每个GPC有4个SM，对比Tesla架构来说，在处理单元上有了很大的提升。

NVidia Kepler架构

Kepler除了在硬件有了提升，有了更多处理单元之外，还将SM升级到了SMX。SMX是改进的架构，支持动态创建渲染线程（下图），以降低延迟。

NVidia Turing架构

上图是采纳了Turing架构的TU102 GPU，它的特点如下：

• 6 GPC（图形处理簇）
• 36 TPC（纹理处理簇）
• 72 SM（流多处理器）
• 每个GPC有6个TPC，每个TPC有2个SM
• 4,608 CUDA核
• 72 RT核
• 576 Tensor核
• 288 纹理单元
• 12x32位 GDDR6内存控制器 (共384位)

单个SM的结构图如下：

每个SM包含：

• 64 CUDA核
• 8 Tensor核
• 256 KB寄存器文件

GPU运行机制

下图是Fermi架构总览：

​ 从Fermi开始NVIDIA使用类似的原理架构，使用一个Giga Thread Engine来管理所有正在运行的工作，GPU被划分为多个GPCs（Graphics Processing Cluster），每个GPC拥有多个SM（SMX、SMM）和一个光栅化引擎（Raster Engine），它们其中有很多的连接，最显著的是CrossBar，它可以连接GPCs和其他功能性模块（例如ROP或其他子系统）。

​ 程序员编写的shader（着色器）是在SM上完成的。每个SM包含许多为线程执行数学运算的Core（核心）。例如，一个线程可以是顶点或像素着色器调用。这些Core和其他单元由Warp Scheduler驱动，Warp Scheduler管理一组32个线程作为Warp（线程束），并将要执行的指令移交给Dispatch Units。

如上图，对于某些GPU（如Fermi部分型号）的单个SM，包含：

• 32个运算核心 （Core，也叫流处理器Stream Processor）
• 16个LD/ST（load/store）模块来加载和存储数据
• 4个SFU（Special function units）执行特殊数学运算（sin、cos、log等）
• 128KB寄存器（Register File）
• 64KB L1缓存
• 全局内存缓存（Uniform Cache）
• 纹理读取单元
• 纹理缓存（Texture Cache）
• PolyMorph Engine：多边形引擎负责属性装配（attribute Setup）、顶点拉取(VertexFetch)、曲面细分、栅格化（这个模块可以理解专门处理顶点相关的东西）。
• 2个Warp Schedulers：这个模块负责warp调度，一个warp由32个线程组成，warp调度器的指令通过Dispatch Units送到Core执行。
• 指令缓存（Instruction Cache）
• 内部链接网络（Interconnect Network）

GPU逻辑管线

下面以Fermi家族的SM为例，进行逻辑管线的详细说明：

1. 程序通过图形API（DX、GL、WEBGL）发出drawcall指令，指令会被推送到驱动程序，驱动会检查指令的合法性，然后把指令放到GPU可以读取的Pushbuffer中。

2. 经过一段时间或者显式调用flush指令后，驱动程序把Pushbuffer的内容发送给GPU，GPU通过主机接口（Host Interface）接受这些命令，并通过前端（Front End）处理这些命令。

3. 在图元分配器（Primitive Distributor）中开始工作分配，处理indexBuffer中的顶点三角形并分成批次（batches），发送给多个GPCs，这一步就是将提交上来的n个三角形，分配给这几个GPCs同时处理。

   

4. 在GPC中，每个SM中的Poly Morph Engine负责通过三角形索引（triangle indexes）取出三角形的数据（vertex data），即图中的Vertex Fetch模块。

5. 在获取数据之后，在SM中以32个线程为一组的线程束（Warp）来调度，开始处理顶点数据。Warp是典型的单指令多线程（SIMT，SIMD单指令多数据的升级）的实现，也就是32个线程同时执行的指令是一模一样的，只是线程的数据不一样，这样的好处是一个warp只需要一套逻辑对指令进行解码和执行就可以了，芯片可以做得更小更快，所以可以说GPU处理的任务天然是并行的。

6. SM的warp调度器会按照顺序分发指令给整个warp，单个warp中的线程会锁步（lock-step）执行各自的指令，如果线程碰到不激活执行的情况也会被遮掩（be masked out）。被遮掩的原因有很多，例如当前的指令是if(true)的分支，但是当前线程的数据是false，或者循环的次数不一样（比如for循环次数n不是常量，或者被break了，但是别的还在走），因此在shader（着色器）中的分支会显著增加时间消耗，在一个warp中的分支除非32个线程都走到if或者else里面，否则相当于所有的分支都走了一遍。线程不能独立执行指令，而是以这些warp为单位，这些warp之间才是独立的。

7. warp中的指令可以被一次执行，也可以经过多次调度执行。例如通常SM中的LD/ST（读写）单元数量明显少于基础数学操作单元。

8. 由于某些指令比其他指令需要更长的时间才能完成，特别是内存加载，warp调度器可能会简单地切换到另一个没有内存等待的warp。为了切换更快，调度器管理的所有warp在寄存器文件中都有自己的寄存器。

   画图处理流程如下：

   

9. 一旦warp完成了vertex-shader的所有指令，运算结果会被Viewport Transform模块处理，三角形会被裁剪然后进行栅格化，GPU会使用L1和L2缓存来进行vertex-shader和pixel-shader的数据通信。

   

10. 接下来这些三角形会被分割，再分配給多个GPC，三角形的范围决定着它被分配到哪个光栅引擎（raster engines），每个光栅引擎覆盖了多个屏幕的tile，这等于把三角形的渲染分配到多个tile上面。也就是像素阶段就把三角形按像素划分了。

    

11. SM上的Atrribute Setup保证了vertex-shader来的数据经过插值之后是pixel-shader可读的。

12. GPC上的光栅引擎（raster engines）在它接收到的三角形上工作，负责这些三角形的像素信息的生成（同时会裁剪、背面剔除和Early-Z剔除）。

13. 32个像素线程被分成一组，或者说8个2*2的像素块，这是像素着色器上面的最小的工作单元，在这个像素线程内，如果没有被三角形覆盖就会被遮掩，SM中的warp调度器会管理像素着色器的任务。

14. 接下来的阶段和vertex-shader中的步骤完全一致，但是变成了在像素着色器线程中执行，由于不耗费任何性能可以获取一个像素内的值，导致锁步执行非常便利，所有的线程的指令可以保证在同一点。

    

15. 最后一步，在像素着色器已经完成了颜色的计算和深度值的计算之后，在这个点上，我们必须考虑三角形的原始API的顺序，然后再将数据移交给ROP（render output unit，渲染输入单元），一个ROP内部有很多ROP单元，在ROP单元中处理深度测试，和framebuffer的混合，深度和颜色的设置必须是原子操作，否则两个三角形在同一像素点就会有冲突和错误。

图的基本知识

稀疏矩阵的存储方式

Coordinate（COO）

​ 这是最简单的一种存储格式，每一个元素需要使用一个三元组来表示，分别是（行号，列号，数值），一个例子，如下：

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

6, 0, 8, 0
0, 7, 0, 0
0, 0, 0, 0
0, 0, 0, 5
对应的存储矩阵如下
0, 0, 6
3, 3, 5
1, 1, 7
0, 2, 8
即row = [0, 3, 1, 0]
  col = [0, 3, 1, 2]
  data = [6, 5, 7, 8]

Compressed Sparse Column（CSC）

​ CSC是按列进行存储的，其中indptr中的数据代表矩阵中每一列所存的数据在data中的开始和结束的索引，例如下面的indptr为[0, 2, 3, 6]，即表示在data中，索引[0, 2)为第一列的数据，索引[2, 3)为第二列的数据，索引[3, 6)为第三列的数据。而indices中的数据代表data数据对应的行（len(data) = len(indices)）。例如，这里的indices为[0, 2, 2, 0, 1, 2]，代表data中数据1在第0行，数据2在第2行，数据3在第2行，数据4在第0行，数据5在第一行，数据6在第2行。从而建立起一个稀疏矩阵。

1
2
3
4
5
6
7

>>> indptr = np.array([0, 2, 3, 6])
>>> indices = np.array([0, 2, 2, 0, 1, 2])
>>> data = np.array([1, 2, 3, 4, 5, 6])
>>> sparse.csc_matrix((data, indices, indptr), shape=(3, 3)).toarray()
array([[1, 0, 4],
       [0, 0, 5],
       [2, 3, 6]])

Compressed Sparse Row（CSR）

​ CSR是按行来进行存储稀疏矩阵的，其原理和CSC类似。indptr中的数据表示矩阵中的每一行在data中的开始和结束索引，而indices表示data数据所在的列数。一个例子如下：

1
2
3
4
5
6
7

>>> indptr = np.array([0, 2, 3, 6])
>>> indices = np.array([0, 2, 2, 0, 1, 2])
>>> data = np.array([1, 2, 3, 4, 5, 6])
>>> sparse.csr_matrix((data, indices, indptr), shape=(3, 3)).toarray()
array([[1, 0, 2],
       [0, 0, 3],
       [4, 5, 6]])

同构图和异构图

​ 在同构图中，所有节点表示同一类型的实体，所有边表示同一类型的关系。例如，社交网络的图由表示同一实体类型的人及其相互之间的社交关系组成。在异构图中，节点和边的类型可以不同。例如，市场的图可以有表示“顾客”、“商家”和“商品”的节点，它们通过“想购买”、“已经购买”、“是顾客”和“正在销售”的边进行相互连接。二分图是一种特殊的、常用的异构图，其中的边连接两类不同类型的节点。例如，在推荐系统中，可以使用二分图表示“用户”和“物品”之间的关系。

图的一些属性

特征向量中心性

考虑了自己的度以及它连接节点的度：$A*x = \lambda *x$ 。

中介中心性

$$
Betweenness = \frac{经过该节点的最短路径}{其余两两节点的最短路径}
$$

中介中心性和介数类似，只不过上下都是总数

连接中心性

$$
Closeness = \frac{n-1}{节点到其他节点最短路径之和}
$$

图的一些算法

PageRank

节点a的PageRank值计算如下：
$$
PR(a){i+1}=\sum{i=0}^n\frac{PR(Ti)_i}{L(Ti)}
$$

• $PR(Ti)$代表的是其他节点（指向a节点的节点）的PR值
• $L(Ti)$代表的是其他节点（指向a节点的节点）的出度数
• $i$代表的是循环次数

下面是PageRank的一个实例，PageRank在多次迭代之后会逐渐收敛。

Make命令教程

​ 代码变成可执行文件，叫做编译，先编译这个，还是先编译那个（即编译的顺序安排），叫做构建（build）。make就是最常用的构建（相当于批处理）工具，主要用于C语言的项目，只要某个文件有变化，都要重新构建项目。

Make的概念

​ make的意思就是“制作”，make命令就是要做出某个文件，比如，要做出文件a.txt，就可以执行下面的命令。

1

make a.txt

​ 但是，如果真的输入这个命令，并不会有任何作用，因为make命令本身并不知道如何作出a.txt，需要有人告诉它，如何调用其他命令完成这个目标。

​ 比如，假设文件a.txt依赖于b.txt和c.txt，是两个文件连接（cat）的产物，那么，make命令需要知道下面的规则。

1
2

a.txt: b.txt c.txt    #表示a依赖于b和c
	cat b.txt c.txt > a.txt		# 利用b和c生成a

也就是说，make a.txt这条命令的背后，实际上分为两步：第一步，确认b.txt和c.txt必须已经存在，第二步使用cat命令将这两个文件合并，输出为新文件。

像这样的规则，都写在一个叫做Makefile的文件中，Make命令依赖于这个文件进行构建，可使用命令指定规则的文件名，如下，指定为rules.txt文件作为规则。

1
2
3

make -f rules.txt
# 或者
make --file=rules.txt

Makefile文件的格式

概述

Makefile文件由一系列规则（rules）构成，每条规则的形式如下：

1
2

<target>: <prerequisites> # target叫做目标，后面的叫做前置条件
[tab]  <commands> # 第二行必须由一个tab键起首，后面跟着命令

target是必须的，不可省略，前置条件和命令两个必须存在一个。

每条规则实际就是在明确两件事：构建目标的前置条件是什么，以及如何构建

目标（target）

一个目标就构成一条规则，目标通常是文件名，指明Make命令所要构建的对象，比如上文中的a.txt，目标可以是一个文件名，也可以是多个空格分隔的文件名。除了文件名，目标还可以是某个操作的名字，这称为“伪目标”

1
2

clean:
	rm  *.o

上面代码的目标是clean，它不是文件名，而是一个操作的名字，属于“伪目标”，作用是删除相应的文件。但是，如果当前目录中已经存在一个名为clean的文件，make clean就不会被执行，为避免这种情况，可以明确声明clean是伪目标，写法如下：

1
2
3

.PHONY: clean
clean: 
	rm *.o temp

因为声明为伪目标后，make就不会去检查是否存在一个叫做clean的文件，而是每次运行都执行对应的命令。

如果make 命令运行时没有指定目标，默认会执行Makefile文件的第一个目标(make)

前置条件（prerequistites）

​ 前置条件通常是一组文件名，之间用空格分隔。它指定了“目标”是否重新构建的判断标准：只要有一个前置文件不存在，或者有过更新（前置文件的last-modification时间戳比目标的时间戳新），“目标”就需要重新构建。

​ 一个例子如下，构建result.txt的前置条件是source.txt，如果当前目录中，source.txt已经存在，那么make result.txt 就可以正常运行，否则，必须再写一条规则来生成source.txt。

1
2

result.txt: source.txt
	cp source.txt  result.txt

无source的情况生成其如下：

1
2

source.txt:
	echo "this is the source" > source.txt

命令（commands）

​ 命令（commands）表示如何更新目标文件，由一行或多行的shell命令组成，它是构建“目标”的具体指令，它的运行结果通常是生成目标文件。

​ 每行命令之前必须有tab键，如果想用其他键，可以用内置变量.RECIPEPREFIX声明。

1
2
3

.RECIPEREFIX = >		#使用>来代替tab
all:
> echo Hello,world		#实际使用来代替

需要注意的是，每行命令在一个单独的shell中执行，这些shell之间没有继承关系，即行与行之间的shell变量不能传递，一个例子如下：

1
2
3

var-lost: 
	export foo=bar
	echo "foo=[$$foo]" #获取不到foo变量的值

改进方法是将其写在一行，中间用分号分隔或者用反斜杠转义。

1
2
3
4
5

var-kept:
	export foo=bar; echo "foo=[$$foo]"
var-kept2:
	export foo=bar; \
	echo "foo=[$$foo]"

还有一个方法是加上.ONSHELL:命令。

1
2
3
4

.ONSHELL
var-kept:
	export foo=bar;
	echo "foo=[$$foo]"

Makefile文件的语法

注释

井号（#）在Makefile中表示注释

回声（echoing）

正常情况下，make会打印每条命令，然后再执行，这就叫回声（echoing），在命令前加上@，就可以关闭会上。

判断和循环

makefile只用Bash语法，完成判断和循环。

1
2
3
4
5

ifeq ($(CC), gcc)
	libs=$(libs_for_gcc)
else
	libs=$(normal_libs)
endif

上面代码判断当前编译器是否是gcc，然后指定不同的库文件。

又一个例子，如下：

1
2
3
4
5
6
7
8
9
10

LIST = one two three
all:
	for i in $(LIST); do \
		echo $$i; \
	done
# 等同于
all:
	for i in one two three; do \
		echo $$i; \
	done

上面的运行结果如下：

1
2
3

one
two
three

CUDA笔记

CUDA函数前缀

​ CUDA使用cu作为文件类型后缀，在CUDA中，有三种常见的前缀__device__、__global__、__host__，分别代表CUDA的三种运行场景，如下表所示：

注意，一个函数可以同时被多个前缀所修饰，如CUDA 10浮点数的转换：

1

__host__ __device__ __half__ __float2half( const float a ) throw()

上面这个函数可以在host端被调用，也可以在device端被调用，我们可以通过一个函数前缀判断这个函数的运行环境。

CUDA shared memory

​ CUDA内存模型如上，最底层的DRAM代表global memory。在global memory部分，数据对齐仍然是很重要的话题，当使用L1（L1和Shared Memory是共享的）的时候，对齐问题可以忽略，但是非连续内存仍然会降低性能。在某些情况下，非连续的访问不可避免，使用shared memory 可以提高系统的性能。

GPU上的内存有两种：

• On-board memory(板载显存)：主要包括全局内存（global memory）、本地内存（local memory）、常量内存（constant memory）、纹理内存（texture memory）等
• On-chip memory（片上内存）：主要包括寄存器（register）和共享内存（shared memory）

片上内存往往比板载内存快很多，因此我们可以使用shared来进行编程，其主要作用有：

• 线程间交流通道
• 可编程的cache
• 转换数据的临时存储器，以减少全局内存访问

shared memory(SMEM)是GPU的重要组成之一，物理上，每个SM包含一个当前正在执行的block中所有thread共享的低延迟的内存池。SMEM使得同一个block中的thread可以相互合作，重用on-chip数据，并且能够减少kernel需要的global memory带宽。

​ 由于shared memory和L1要比L2和global memory更接近SM，shared memory的延迟要比global memory低20到30倍，带宽大约高10倍。

​ 当一个block开始执行时，GPU会分配其一定数量的shared memory，这个shared memory的地址空间会由block的所有thread共享。所以，使用越多的shared memory，能够并行的active就越少。

Shared Memory Allocation（SMEM分配）

​ 我们可以动态或者静态地分配shared memory，其声明可以在kernel内部，也可以作为全局变量。其标识符为__shared__。下面这句声明了一个2D的浮点型数组：

1

__shared__ float tile[size_x][size_y];

如果在kernel中声明的话，其作用域就是在kenel内，否则就是对所有kernel有效。如果shared memory的大小在编译期未知的话，可以使用extern关键字修饰，如下面声明一个未知大小的1D数组：

1

extern __shared__ int tile[];

然后在每个kernel调用时需要使用kernel<<<grid, size*sizeof(int)>>>(...)，而且，只有一维数组才能这样使用。

Shared Memory Banks and Access Mode（共享内存的带宽和访问模式）

pinned memory(固定内存)

在CUDA编程中，内存拷贝是一个非常费时的一个动作。

从上图我们可以看出：

1. CPU和GPU之间的总线是PCIe，双向传输
2. CPU和GPU之间的数据拷贝是使用DMA机制来实现

我们使用cudaMalloc为GPU分配内存，malloc为CPU分配内存，除此之外，CUDA还提供了自己独有的机制来分配host内存：cudaHostAlloc()，它与malloc不同的是，malloc分配的是可分页的主机内存，而cudaHostAlloc分配的是页锁定的主机内存，也称作固定内存（pinned memory）、不可分页内存。它的一个重要特点是操作系统不会对这块内存进行分页并交换到磁盘上，从而保证了内存始终驻留在物理内存中。因此，操作系统能够安全地使某个应用程序访问该内存的物理地址，因为其不会被破坏或重新定位。

由于GPU知道内存的物理地址，因此就可以使用DMA技术来在GPU和CPU之间复制数据。当使用可分页的内存进行复制时（使用malloc），CUDA驱动程序仍会通过dram把数据传给GPU，但这时数据会执行两遍：

1. 从可分页内存复制一块到临时的页锁定内存
2. 从这个页锁定内存复制到GPU上

而页锁定内存只需执行后面这一步，因而速度提高了一倍。

当我们在调用cudaMemcpy(dst, src, …)时，程序会自动检测dest或者src是否为Pinned Memory，若不是，则会先将内容拷进一不可见的Pinned Memory中，然后再进行传输。可以手动指定Pinned Memory，对应的API为cudaHostAlloc(address, size, option)分配地址，cudaFreeHost(pointer)释放地址。注意，分配的内存都是在Host端。

不过，把所有的malloc都替换为cudaHostAlloc()是不对的。固定内存是一把双刃剑，当使用固定内存时，虚拟内存的功能就会失去，尤其是在应用程序中都使用固定内存时，会导致系统内存很快被耗尽，影响自身以及其他应用程序的执行。

所以，建议针对cudaMemcpy()调用中的源内存或者目标内存，才使用页锁定内存，并且在不再使用他们的时候立即释放，而不是关闭应用程序的时候才释放。

零拷贝内存

通常来说，主机不能直接访问设备变量，同时设备变量也不能直接访问主机变量。但是有一个例外：零拷贝内存，主机和设备都可以访问零拷贝内存。它允许你将主机内存直接映射到GPU内存空间上。因此，当你对GPU上的内存解引用时，如果它是基于GPU的，那么你就将获得了全局内存的高速带宽（200GB/s），如果是零拷贝内存的，它会提交一个PCI-E读取事务，很长时间之后，主机会通过PCI-E总线返回数据，带宽降低到16GB/s。

当程序是计算密集型时，零拷贝可能是一项非常有用的技术，它节省了设备显式传输的时间，事实上，它可以将计算和数据传输流水化，而且无需执行显式的内存管理。

当使用零拷贝内存来共享主机和设备间的数据时，必须同步主机和设备间的内存访问，同时更改主机和设备的零拷贝内存中的数据将导致不可预料的后果。

零拷贝内存是固定（不可分页）内存，该内存映射到设备地址空间中，使用以下函数创建一个到固定内存的映射（和锁页内存类似）：

1

cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);

这个函数分配了count字节的主机内存，该内存是页面锁定的内存且设备可访问的，用这个函数分配的内存必须用cudaFreeHost函数释放。flags参数可以对以分配的特殊属性进一步进行配置：

• cudaHostAllocDefault：使cudaHostAlloc函数的行为与cudaMallocHost函数一致。
• cudaHostAllocPortable：可以返回被CUDA上下文所使用的固定内存，而不仅仅是执行内存分配。
• cudaHostAllocWriteCombined：返回写结合内存，该内存可以在某些系统配置上通过PCIe总线更快地传输，但是它在大多数主机上不能被有效地读取。因此，写结合内存对缓冲区来说是一个很好的选择，该内存通过设备使用映射的固定内存或主机到设备的传输。
• cudaHostAllocMapped：零拷贝内存，该标志返回可以实现主机写入和设备读取被映射到设备地址空间的主机内存。

分配好零拷贝内存之后，就可以使用下列函数获取映射到固定内存的设备指针了：

1

cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

该函数返回一个在pDevice中的设备指针，该指针可以在设备上被引用以访问映射得到的固定主机内存。如果设备不支持映射得到的固定内存，则该函数失效。flag保留，始终为0.

在频繁进行读写操作时，使用零拷贝内存作为设备内存的补充将显著降低性能。因为每一次映射到零拷贝内存的传输必须经过PCIe总线，与全局内存相比，延迟也显著增加。

在使用零拷贝内存时，需要检查设备是否支持固定内存映射，cudaDeviceProp的canMapHostMemory成员是一个bool类型值，true表示支持固定内存映射。

一个使用例子如下：

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

#include <numeric>
#include <stdio.h>
#include <stdlib.h>
void checkCUDAError(const char *msg) {
    cudaError_t err = cudaGetLastError();
    if (cudaSuccess != err) {
        fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
}

__global__ void sumNum( int *data) {

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if(i<1000000000){
     data[i]=10;
    }
}
int main(void) {
    size_t size = 1*1000000000 * sizeof(int);//4G
    //1.启用零复制
    cudaSetDeviceFlags (cudaDeviceMapHost);
    int* data;
    //2.分配主机内存
    cudaHostAlloc((void**) &data, size,
            cudaHostAllocWriteCombined | cudaHostAllocMapped);
     checkCUDAError("cudaHostAlloc data");
    
     memset(data, 0, 1*1000000000 * sizeof(int));
    int *gpudata;
    //3.将常规的主机指针转换成指向设备内存空间的指针
    cudaHostGetDevicePointer(&gpudata, data, 0);
    checkCUDAError("cudaHostGetDevicePointer");
    sumNum<<<1000000000/1024+1023, 1024>>>(gpudata);
    //注意！！因为下面要打印出来测试，所以要先同步数据，这个函数可以保证cpu等待gpu的kernel函数结束才往下运行。如果数据暂时用不到，可以在整体结束以后再加这句话。明显等待kernel函数结束会占用程序进行的时间。
    cudaDeviceSynchronize();
    for (int i = 99999999; i < 1000000000; i=i+100000000) {
        printf("%d \n", data[i]);
    }
    //记得零拷贝的free是这个函数
    cudaFreeHost(data);
    return 0;
}

注意零拷贝内存还有另一种实现方式，就是直接对malloc申请的内存使用cudaHostRegister进行标志位设置（设置成固定内存），然后使用cudaHostGetDevicePointer获取其设备指针。代码例子如下：

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

#include "arrayTools.h"
#include "cudaCode.h"
#include "MyTimer.h"

template<class T>
__global__ void warmup(T *A, T *B, T *C, const int n)
{
    int tid = blockDim.x*blockIdx.x + threadIdx.x;
    if (tid >= n)return;
    C[tid] = A[tid] + B[tid];
}

template<class T>
__global__ void sumVecOnDeviceZeroCopy(T *A, T *B, T *C, const int n)
{
    int tid = blockDim.x*blockIdx.x + threadIdx.x;
    if (tid >= n)return;
    C[tid] = A[tid] + B[tid];
}

template<class T>
__global__ void sumVecOnDeviceZeroCopyRegister(T *A, T *B, T *C, const int n)
{
    int tid = blockDim.x*blockIdx.x + threadIdx.x;
    if (tid >= n)return;
    C[tid] = A[tid] + B[tid];
}
int main()
{
    int nElem=1<<20;
    int nBytes=nElem*sizeof(float);
    float * d_A, * d_B, * d_C;
    float * h_A, * h_B, * gpuRes;
    h_A=(float*)malloc(nBytes);
    h_B=(float*)malloc(nBytes);
    initialInt(h_A, nElem);
    initialInt(h_B, nElem);
    gpuRes=(float*)malloc(nBytes);
    
    dim3 block(512, 1);
    dim3 grid((nElem + block.x - 1) / block.x);

    //将h_A、h_B、h_C三块内存锁定并获取对应的设备地址，进行计算，这样避免了从一开始就得cudaHostAlloc申请零拷贝内存的不便
    CHECK(cudaHostRegister(h_A, nBytes, cudaHostRegisterMapped));
    CHECK(cudaHostGetDevicePointer((void **)&d_A, h_A, 0));
    CHECK(cudaHostRegister(h_B, nBytes, cudaHostRegisterMapped));
    CHECK(cudaHostGetDevicePointer((void **)&d_B, h_B, 0));
    CHECK(cudaHostRegister(gpuRes, nBytes, cudaHostRegisterMapped));
    CHECK(cudaHostGetDevicePointer((void **)&d_C, gpuRes, 0));
    memset(gpuRes, 0, nBytes);
    sumVecOnDeviceZeroCopyRegister<float> << <grid, block >> > (d_A, d_B, d_C, nElem);
    CHECK(cudaDeviceSynchronize());
    
    compareVec(gpuRes, cpuRes, 0, nElem);
    return 0;
}

集成架构和离散架构

这两种架构是常见的异构计算系统架构。

集成架构：cpu和gpu集成在一个芯片上，并且在物理内存上共享内存，在这种架构中，由于无需在PCIe总线上备份，所以零拷贝内存在性能和可编程性方面更好一些。

离散架构：cpu和gpu是分离的，物理内存上也是分离的。数据需要通过PCIe总线进行交互，数据拷贝的耗时和延迟的代价通常也比较大。因此在这种架构下，零拷贝只在特殊情况下有优势。

由于通过映射的固定内存在主机和设备之间是共享的，所以必须同步内存访问来避免任何潜在的数据冲突，这种数据冲突一般是有多线程异步访问相同的内存而引起的。

CUDA编程流程

CUDA编程模型基础

CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。在CUDA中，host和device是两个重要的概念。host指代CPU及其内存，device指代GPU及其内存。CUDA程序既包括host程序，又包括device程序，它们分别在CPU和GPU上运行。同时，host和device之间可以进行通信，这样他们之间可以进行数据拷贝。典型的CUD程序执行流程如下：

1. 分配host内存，并进行数据初始化
2. 分配device内存，并从host拷贝数据到device上
3. 调用CUDA的核函数在device上完成指定的运算
4. 将device上的运算结果拷贝到host上
5. 释放host和device上分配的内存

上面最重要的一个过程是调用CUDA的核函数来执行并行计算，kernel是CUDA中一个重要的概念，kernel是在device上线程中并行执行的函数。核函数使用__global__来声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量，并且每个线程会分配一个唯一的thread ID，这个ID值可以通过核函数内置的变量threadidx来获得。

在CUDA中使用不同的限定词来区分host和device上的函数，三个前缀在前面已经有了说明。

kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一网格上的线程共享相同的全局内存空间，grid是线程的第一层次。同一个网格上的线程又可以分为很多线程块（block），一个线程块里包含很多线程，这是第二层次。两层结构如下图所示：

1
2
3
4

// 对应的代码
dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<<grid, block>>>(params...);

这是一个grid和block均为2-dim的线程组织。grid和block都定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x, y, z）成员的结构体变量。在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，对于图中结构（水平方向为x轴），展示如上。kernel在调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。

所以，一个线程需要两个内置的坐标变量（blockidx， trheadidx）来唯一标识，它们都是dim3类型变量，其中blockidx指明线程所在的grid中的位置，而threadidx指明线程所在block中的位置。如图中的Thread（1, 1）满足：

1
2
3
4

threadIdx.x = 1;
threadIdx.y = 1;
blockIdx.x = 1;
blockIdx.y = 1;

​ 一个线程块的线程是放在同一个流式多处理器（SM）上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块可支持的线程数达1024个。有时候，我们要知道一个线程在block中的全局ID，就必须知道block的组织结构，这是通过blockDim来获得的。它获取线程块各个维度的大小。对于一个2-dim的block(Dx, Dy)，线程（x，y）的ID值为(x + y * Dx)，此外，如果是3-dim的block(Dx, Dy, Dz)，线程(x, y, z)的ID值为（x + y * Dx + z * Dx * Dy）。另外还有内置变量gridDim，用于获取网格块各个维度的大小。

​ kernel的这种线程组织结构天然适合vector，matrix等运算，如下我们将利用上图2-dim结构实现两个矩阵的加法，每个线程负责处理两个位置的元素相加。线程块大小为(16, 16)，然后将N×N大小的矩阵均分为不同的线程块来执行加法运算。

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

//Kernel定义
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if(i < N && j < N){
        C[i][j] = A[i][j] + B[i][j];
    }
}

int main(){
    ...;
    // Kernel线程配置
    dim3 threadPerBlock(16, 16);
    dim3 numBlocks(N / threadPerBlock.x, N / threadPerBlock.y);
    //Kernel调用
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

​ 再次强调CUDA内存模型如下，每个线程都有自己的私有本地内存（Local Memory），而每个线程块都有包含共享内存（Shared Memory），可以被县城快中所有的线程共享，其生命周期和线程块一致。此外，所有的线程都可以访问全局内存（Global Memory），以及一些只读内存块：常量内存（Constant Memory）和纹理内存（Texture Memory）。

​ GPU有很多CUDA核心，可以充分发挥GPU的并行能力。GPU硬件的一个核心组件是SM（Streaming Multiprocess，流式多处理器）。SM的核心组件包括CUDA核心，共享内存，寄存器等，SM可以并发地执行数百个线程，并发能力就取决于SM所拥有的资源数。当一个kernel被执行时，它的grid中的线程块被分配到SM上。一个线程块只能在一个SM上被调度。SM一般可以调度多个线程块。一个kernel上的各个线程可能被分配给多个SM，所以grid只是逻辑层，SM才是物理层。SM采用的是SIMT（Single-Instruction，Multiple-Thread，单指令多线程）架构，基本的执行单元是线程束（wraps），线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。所以尽管线程束中的线程同时从统一程序地址执行，但是可能有不同的行为，比如遇到了分支结构，一些线程可能进入这个分支，但是另外一些有可能不执行，它们只能死等。因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。

​ 当线程块被划分到某个SM上时，它将进一步划分为多个线程束，因为这才是SM的基本执行单元，但是一个SM同时并发的线程束数是有限的，这是因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程分配独立的寄存器。所以SM的配置会影响其所支持的线程块和线程束的并发数量。总之，就是网格和线程块只是逻辑划分，一个kernel的所有线程其实在物理层不一定同时并发。所以kernel的grid和block的配置的不同，性能会出现差异。由于SM的基本执行单元是包含32个线程的线程束，所以block一般要设置为32的倍数。

​ 在进行CUDA编程前，可以先检查自己的GPU的硬件配置，通过下面的程序可以获得GPU的配置属性：

1
2
3
4
5
6
7
8
9

int dev = 0;
cudaDeviceProp devProp;
CHECK(cudaGetDeviceProperties(&devProp, dev));
std::cout << "使用GPU device " << dev << ": " << devProp.name << std::endl;
std::cout << "SM的数量：" << devProp.multiProcessorCount << std::endl;
std::cout << "每个线程块的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
std::cout << "每个线程块的最大线程数：" << devProp.maxThreadsPerBlock << std::endl;
std::cout << "每个EM的最大线程数：" << devProp.maxThreadsPerMultiProcessor << std::endl;
std::cout << "每个SM的最大线程束数：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;

向量加法实例

下面介绍一下CUDA编程中的内存管理API，首先是在device上分配内存的cudaMalloc函数:

1

cudaError_t cudaMalloc(void** devPtr, size_t size);

这个函数和C语言中的malloc类似，在device上申请一定字节大小的显存，其中devPtr是指向所分配内存的指针。同时释放内存使用cudaFree函数，这和C语言中的free函数所对应。另外一个重要的函数是负责host和device之间数据通信的cudaMemcpy函数：

1

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

其中src指向数据源，dst是目标区域，count是复制的字节数，kind控制复制的方向：cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost及cudaMemcpyDeviceToDevice，如cudaMemcpyHostToDevice将host上数据拷贝到device上。

以下是一个向量加法实例，这里的grid和block都设计为1-dim，首先定义kernel如下：

1
2
3
4
5
6
7
8
9
10

// 两个向量加法，grid和block均为一维
__global__ void add(float* x, float* y, float* z, int n){
    // 获取全局索引，避免线程重复处理同一个元素，相当于把二维数组一维化
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    // 步长，每次有多少个线程在处理
    int stride = blockDim.x * gridDim.x;
    for(int i = index; i < n; i += stride){
        z[i] = x[i] + y[i];
    }
}

其中stride是整个grid的线程数，有时候向量的元素数很多，这时可以在一个线程实现多个元素的加法（元素总数/线程总数的加法），相当于使用了多个grid来处理，这是一种grid-strip loop的方式。不过下面的例子一个线程只处理一个元素，所以kernel里面的循环是不执行的。下面实现具体的向量加法：

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

int main(){
    int N = 1 << 20;
    int nBytes = N * sizeof(float);
    // 申请host内存
    float *x, *y, *z;
    x = (float*)malloc(nBytes);
    y = (float*)malloc(nBytes);
    z = (float*)malloc(nBytes);
    
    // 初始化数据
    for(int i = 0; i < N; ++i){
        x[i] = 10.0;
        y[i] = 20.0;
    }
    
    // 申请device上的内存
    float *d_x, *d_y, *d_z;
    cudaMalloc((void**)&d_x, nBytes);
    cudaMalloc((void**)&d_y, nBytes);
    cudaMalloc((void**)&d_z, nBytes);
    
    // 将host数据拷贝到device上
    cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);
    
    //定义kernel的执行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    
    //执行kernel
    add<<<gridSize, blockSize>>>(d_x, d_y, d_z, N);
    
    // 将device结果拷贝到host
    cudaMemcpy((void*)z, (void *)d_z, nBytes, cudaMemcpyDeviceToHost);
    
    // 检查执行结果
    float maxError = 0.0;
    for(int i = 0; i < N; i++){
        maxError = fmax(maxError, fabs(z[i] - 30.0));
    }
    std::cout << "最大误差为：" << maxError << std::endl;
    
    // 释放device内存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    
    // 释放host内存
    free(x);
    free(y);
    free(z);
    
    return 0;
}

这里我们向量的大小是1<<20，而block的大小为256,那么grid大小为（2^20 / 256）= 4096，kernel的线程层级如下图所示：

同时，block不是越大越好，要适当选择。

在上面的实现中，我们需要单独在host和device上进行内存分配，并且进行数据拷贝，这是很容易出错的，在CUDA 6.0引入了统一内存(Unified Memory)来避免这种麻烦，简单地说就是统一内存使用一个托管内存来共同管理host和device的内存，并且自动在host和device中进行数据传输。CUDA中使用cudaMallocManaged函数分配托管内存：

1

cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flag = 0);

利用统一内存，可以将上面的程序简化如下：

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

int main(){
    int N = 1 << 20;
    int nBytes = N * sizeof(float);
    
    // 申请托管内存
    float *x, *y, *z;
    cudaMallocManaged((void**)&x, nBytes);
    cudaMallocManaged((void**)&y, nBytes);
    cudaMallocManaged((void**)&z, nBytes);
    
    // 初始化数据
    for(int i = 0; i < N; ++i){
        x[i] = 10.0;
        y[i] = 20.0;
    }
    
    // 定义kernel的执行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    
    // 执行kernel
    add<<<gridSize, blockSize>>>(x, y, z, N);
    
    // 同步device，保证结果可以正确访问
    cudaDeviceSynchronize();
    
    // 检查执行结果
    float maxError = 0.0;
    for(int i = 0; i < N; i++){
        maxError = fmax(maxError, fabs(z[i] - 30.0));
    }
    std::cout << "最大误差为：" << maxError << std::endl;
    
    // 释放内存
    cudaFree(x);
    cudaFree(y);
    cudaFree(z);
    
    return 0;
}

需要注意的是kernel执行和host是异步的，由于托管内存自动进行数据传输，这里要用cudaDeviceSynchronize()函数保证device和host同步，这样后面才可以正确访问kernel计算的结果。

矩阵乘法实例

设输入矩阵为A和B，得到的矩阵为C = A×B。实现的思路就是每个线程计算C的一个元素的值Ci,j，对于矩阵运算，应该选用grid和block为2-D的，首先定义矩阵的结构体：

1
2
3
4
5
6

// 矩阵类型，行优先，M(row, col) = *(M.elements + row * M.width + col)
struct Matrix {
    int width;
    int height;
    float *elements;
}

矩阵乘法的实现图如下：

然后实现矩阵的核函数，这里我们定义了两个辅助的__device__函数分别用于获取矩阵的元素值和为矩阵元素赋值，具体代码如下：

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

// 获取矩阵A的（row，col）的元素
__device__ float getElement(Matrix *A, int row, int col) {
    return A->elements[row * A->width + col];
}

// 为矩阵A的（row，col）的元素赋值
__device__ void setElement(Matrix *A, int row, int col, float value) {
    A->elements[row * A->width + col] = value;
}

// 矩阵相乘的kernel，2-D，每个线程计算一个元素
__global__ void matMulKernel(Matrix *A, Matrix *B, Matrix *C) {
    float Cvalue = 0.0;
    int row = threadIdx.y + blockIdx.y * blockDim.y;
    int col = threadIdx.x + blockIdx,x * blockDim.x;
    for(int i = 0; i < A->width; ++i) {
        Cvalue += getElement(A, row, i) * getElement(B, i, col);
    }
    setElement(C, row, col, Cvalue);
}

最后我们使用统一内存来编写矩阵相乘的测试实例：

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

int main(){
    int width = 1 << 10;
    int height = 1 << 10;
    Matrix *A, *B, *C;
    
    // 申请托管内存
    cudaMallocManaged((void**)&A, sizeof(Matrix));
    cudaMallocManaged((void**)&B, sizeof(Matrix));
    cudaMallocManaged((void**)&C, sizeof(Matrix));
    int nBytes = width *height * sizeof(float);
    cudaMallocManaged((void**)&A->elements, nBytes);
    cudaMallocManaged((void**)&B->elements, nBytes);
    cudaMallocManaged((void**)&C->elements, nBytes);
    
    // 初始化数据
    A->height = height;
    A->width = width;
    B->height = height;
    B->width = width;
    C->height = height;
    C->width = width;
    for(int i = 0; i < width * height; i++){
        A->elements[i] = 1.0;
        B->elements[i] = 2.0;
    }
    
    // 定义kernel的执行配置
    dim3 blockSize(32, 32);
    dim3 gridSize((width + blockSize.x - 1) / blockSize.x, 
                 (height + blockSize.y - 1) / blockSize.y);
    
    // 执行kernel
    matMulKernel <<<gridSize, blockSize>>>(A, B, C);
    
    // 同步device，保证结果正确访问
    cudaDeviceSynchronize();
    
    //检查执行结果
    float maxError = 0.0;
    for(int i = 0; i < width * height; i++) {
        maxError = fmax(maxError, fabs(C->elements[i] - 2 * width));
    }
    std::cout << "最大误差为：" << maxError << std::endl;
    
    return 0;
}

这里设计的线程block为（32, 32），grid大小为（32, 32），最终测试结果如下：

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

nvprof cuda9.exe
==16304== NVPROF is profiling process 16304, command: cuda9.exe
最大误差: 0
==16304== Profiling application: cuda9.exe
==16304== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:  100.00%  1.32752s         1  1.32752s  1.32752s  1.32752s  matMulKernel(Matrix*, Matrix*, Matrix*)
      API calls:   83.11%  1.32762s         1  1.32762s  1.32762s  1.32762s  cudaDeviceSynchronize
                   13.99%  223.40ms         6  37.233ms  37.341us  217.66ms  cudaMallocManaged
                    2.81%  44.810ms         1  44.810ms  44.810ms  44.810ms  cudaLaunch
                    0.08%  1.3300ms        94  14.149us       0ns  884.64us  cuDeviceGetAttribute
                    0.01%  199.03us         1  199.03us  199.03us  199.03us  cuDeviceGetName
                    0.00%  10.009us         1  10.009us  10.009us  10.009us  cuDeviceTotalMem
                    0.00%  6.5440us         1  6.5440us  6.5440us  6.5440us  cudaConfigureCall
                    0.00%  3.0800us         3  1.0260us     385ns  1.5400us  cudaSetupArgument
                    0.00%  2.6940us         3     898ns     385ns  1.5390us  cuDeviceGetCount
                    0.00%  1.9250us         2     962ns     385ns  1.5400us  cuDeviceGet

==16304== Unified Memory profiling result:
Device "GeForce GT 730 (0)"
   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
    2051  4.0000KB  4.0000KB  4.0000KB  8.011719MB  21.20721ms  Host To Device
     270  45.570KB  4.0000KB  1.0000MB  12.01563MB  7.032508ms  Device To Host

数据库模式与视图

数据库模式

三级模式

​ 用户从三个层次来管理数据：外部层次（External Level）、概念层次（Conceptual Level）和内部层次（Internal Level）。其中外部层次的数据就是用户所看到的数据，所以又叫用户层次。概念层次的数据是DBMS中全局管理数据和数据之间的约束，所以又叫逻辑层次。内部层次的数据是存储在介质上的数据，包括存储路径、存储方式、索引方式等，所以又叫物理层次。

内模式

​ 又称物理模式，是数据存储方式和物理结构在数据库内部的组织方式。如记录存储方式是堆存储，还是按照某些属性值升降序存储，还是按章属性值聚簇存储。索引的组织方式，是B+树索引还是hash索引或其他等等。

1. 一个数据库只能有一个内模式，因为一个数据库对应的存储方式、路径只能有一个。

模式

​ 也称为逻辑模式或概念模式，是数据库中全体数据的逻辑结构和特征的描述，是所有用户的公共数据视图。实际应用中模式等同于程序员创建一个具体的数据库的全部操作，如：这是一个MySQL数据库，有两张表，每个表的名字，属性的名字、类型、取值范围，主键、外键、索引等。

1. 一个数据库只有一个模式，因为模式是DBMS对所有数据的全局性的结构描述，所以只能有一种形式，如不使用两张不同结构的学生表存放的都是所有学生的信息。
2. 定义模式时不仅要定义数据的逻辑结构（如类型、名字等），而且要定义与数据有关的安全性、完整性要求，定义这些数据之间的联系

外模式

​ 以视图的形式展示给用户，是保证数据库安全性的有力措施。是数据库的数据视图。

1. 一个数据库可有多个外模式，如不同用户可以对同一份数据的所需数据可能不同或用户的权限不同等。

使用概述

数据库产品中的schema并不对应三级模式中的模式，如postgre中的模式可以看作是一个表的集合，其可以包含视图、索引、数据类型、函数和操作符等。以下的话是针对类似postgre的模式的。

一个数据库可以有多个模式

一个模式下可以有多个表

表必须被放在模式中，一般默认放在public模式中

用户有自己的默认模式

可通过create schema创建模式

模式访问：数据库.模式名

数据库视图

​ 数据库视图是用户可以看见的虚关系，是从一个或几个表导出的虚表。数据库中只存放视图的定义（位于数据字典中），不存放视图对应的数据（数据位于其他表中），也称为动态窗口。

​ 视图也可以增删改查，但可能会有一定的限制。如加上check option的视图就要满足相应条件才能update，有两个基本表导出的视图不能更新，来自字段表达式（函数）或常数的视图不允许插入和修改，但可以删除。

​ 视图为数据库重构提供了一定的逻辑独立性，即用户->视图->表，这样如果表结构更改了可能也不用大规模修改。

leetcode: 两数相加

对于链表长度头部未知的，可在头部前多设一个节点来快速获取

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

class Solution {
    public ListNode addTwoNumbers(ListNode l1, ListNode l2) {
        ListNode temList1,temList2,newList,newTemList;
        temList1 = l1;
        temList2 = l2;
        int carryBit = 0, newBit = 0;
        newList = new ListNode(0);
        newTemList = newList;
        while(true){
            if(temList1 == null){
                while(temList2 != null){
                    newBit = temList2.val + carryBit;
                    if(newBit > 9){
                        carryBit = 1;
                        newBit = newBit - 10;
                    }else{
                        carryBit = 0;
                    }
                    newTemList.next = new ListNode(newBit);
                    newTemList = newTemList.next;
                    temList2 = temList2.next;
                }
                break;
            }
            if(temList2 == null){
                while(temList1 != null){
                    newBit = temList1.val + carryBit;
                    if(newBit > 9){
                        carryBit = 1;
                        newBit = newBit - 10;
                    }else{
                        carryBit = 0;
                    }
                    newTemList.next = new ListNode(newBit);
                    newTemList = newTemList.next;
                    temList1 = temList1.next;
                }
                break;
            }
            newBit = temList1.val + temList2.val + carryBit;
            if(newBit > 9){
                carryBit = 1;
                newBit = newBit - 10;
            }
            else{
                carryBit = 0;
            }
            newTemList.next = new ListNode(newBit);
            newTemList = newTemList.next;
            temList1 = temList1.next;
            temList2 = temList2.next;
        }
        if(carryBit == 1){
            newTemList.next = new ListNode(1);
        }
        return newList.next;
    }
    
}

Tex的一些笔记

宏语言

宏来源

​ 在早期使用汇编语言进行编程的时候，程序员发现程序中包含了很多重复的代码，于是他们创造了这种简单的方法来重用。

宏定义

​ 可简单理解为用较短指令表示较长指令，在编程语言中不常用，在文字处理中很常用。

Tex简介

​ Tex就是一种宏语言，主要应用于文字排版，百度百科就把tex说为一种排版系统，可以把这个程序看作为编译器。

​ Tex源文件的后缀是.tex,就是使用代码来规定文字样式，作用类似HTML，不过其功能更丰富，尤其是对数学公式的表示。一段tex文件如下：

\documentclass{article}

\usepackage{xeCJK} %调用 xeCJK 宏包 \setCJKmainfont{SimSun} %设置 CJK 主字体为 SimSun （宋体） \begin{document} 你好，world! \end{document}

最总就是显示相应样式的“你好，world！”

Tex的一些种类及命令

1. plain Tex： Tex中最基本的宏集合和基础语言构成的一种格式。
2. LaTex： Tex的一个宏集合，构成一种与plain Tex不一样的格式。
3. pdfTex： Tex的又一个实现，能将Tex语言直接编译成pdf文件。
4. XeTex： 适应了Unicode字符集。
5. xetex命令： 用于编译plain tex编写的文件，分两阶段处理，第一阶段输出DVI（xdv）文件，第二阶段由驱动生成pdf文件。
6. xelatex命令： 用于编译LaTex编写的文件，也是生成pdf文件。
7. 宏包：一般是在Tex或LaTex下面编写的宏，然后将其打印成包。

Tex的一些散装笔记

1. 使用Tex来写论文比word好，因为其有模板，可直接套用。
2. 可使用LaTex写源代码，然后用pdfTex编译为pdf
3. cls文件： 类文件，规定了tex源文件的排版格式，一般使用\documentclass{}导入
4. sty文件： 宏包文件，使用\usepackage{}导入
5. bst文件： 参考文件，使用\bibliographystyle{}导入
6. bib文件： 参考文献的库，使用\bibliography{}导入
7. aux文件： 辅助文件，不影响正常使用
8. dtx文件： 宏包重要部分
9. ins文件： 控制从dtx文件中释放宏包文件
10. cfg文件： 配置文件，可有以上两个文件生成
11. 对于有参考文献、目录排版的lex文件（如使用biblatex），需要用xelatex编译两次，第一次生成无文献、无目录的pdf文件，然后第二遍读入参考文献和目录并正确生成引用和目录表。

使用MikTex的一些问题

1. 直接使用使用xelatex命令时报dvipdfmx:fatal: File ended prematurely 错误，可能是环境不是UTF-8的问题，一般出现在中文环境中。执行以下命令即可解决：

   chcp 65001;

   xelatex --shell-escape tex文件

2. 出现! TeX capacity exceeded, sorry [main memory size=3000000]. 错误，一个解决方法如下：

   1. 运行cmd打开命令行窗口，输入：initexmf --edit-config-file=xelatex
   2. 在弹出的文件中输入以下内容：main_memory=5000000 extra_mem_bot=5000000 font_mem_size=5000000 pool_size=5000000 buf_size=5000000
   3. 在cmd窗口输入： initexmf --dump=pdflatex 更新latex格式文件

Windows下Hexo的一些问题记录

1. CMD直接使用hexo只出现三种选项，如下

​ 解决方法：将进入Hexo的文件夹里再使用相关命令

2. 为Hexo的文章创建分类

   1. hexo new page "Categories",然后就会在source文件夹下生成Categories文件夹，里面有index.md。这个Categories对应的是所有分类项的页面，我们不需要主动添加分类项进去，只要你的markdown里的categories属性写了分类，该分类就会被添加到Categories页面。

   2. 对index.md的填写如下，comment代表评论

      title: ' Categories ' date: 2020-05-20 22:06:04 type: "categories" comments: false

   3. 对于Next主题的，在hexo\themes\next_config.yml中去除categories前面的#号，并且检查home: /

      和/categories/ 后是否有空格，有则去掉，否则会出现Cannot find /Categories/%20的错误

      home: /|| home #about: /about/ || user #tags: /tags/ || tags categories: /categories/|| th archives: /archives/ || archive #schedule: /schedule/ || calendar #sitemap: /sitemap.xml || sitemap #commonweal: /404/ || heartbeat

   4. 执行hexo g 和hexo s 在localhost:4000即可看到结果，然后hexo d 将修改上传到网页。

3. hexo对语法的检查很严格，开头的title这些信息冒号后面都需一个空格，否则会有Error

4. 对于由于git大小写问题出现的404，如url：/Categories/在GitHub上变为：/categories，把public文件夹里的Categories改为 categories即可

5. Hexo图片引用问题，在source目录下新建一个images文件夹，然后把图片放在里面用路径/images/图片名 进行引用即可，这个方法的Hexo的主页和文章详细页都可见，但在本地不可见,可先放里面，然后deploy，接着使用https://toaoc.github.io/images/图片名 即可引用。

6. Hexo使用categories标签才会出现在分类页面中，而不是Tags标签

7. 博客的头部信息凑不用使用引号，即使是中文