第三章 简介

缓冲与缓存的区别，核函数，内置参数，参数传递

GitHub

建议下载下来用Typora软件阅读markdown文件

作者github:littlebearsama 原文链接

(建议下载Typora来浏览markdown文件)

缓冲区(buffer)与缓存区(cache)

一、缓冲

缓冲区(buffer)，它是内存空间的一部分。也就是说，在内存空间中预留了一定的存储空间，这些存储空间用来缓冲输入或输出的数据，这部分预留的空间就叫做缓冲区，显然缓冲区是具有一定大小的。

缓冲区根据其对应的是输入设备还是输出设备，分为输入缓冲区和输出缓冲区。

二、缓存

1. CPU的Cache，它中文名称是高速缓冲存储器，读写速度很快，几乎与CPU一样。由于CPU的运算速度太快，内存的数据存取速度无法跟上CPU的速度，所以在cpu与内存间设置了cache为cpu的数据快取区。当计算机执行程序时，数据与地址管理部件会预测可能要用到的数据和指令，并将这些数据和指令预先从内存中读出送到Cache。一旦需要时，先检查Cache，若有就从Cache中读取，若无再访问内存，现在的CPU还有一级cache，二级cache。简单来说，Cache就是用来解决CPU与内存之间速度不匹配的问题，避免内存与辅助内存频繁存取数据，这样就提高了系统的执行效率。
2. 磁盘也有cache,硬盘的cache作用就类似于CPU的cache，它解决了总线接口的高速需求和读写硬盘的矛盾以及对某些扇区的反复读取。

三、缓存（cache）与缓冲(buffer)的主要区别

Buffer的核心作用是用来缓冲，缓和冲击。比如你每秒要写100次硬盘，对系统冲击很大，浪费了大量时间在忙着处理开始写和结束写这两件事嘛。用个buffer暂存起来，变成每10秒写一次硬盘，对系统的冲击就很小，写入效率高了，日子过得爽了。极大缓和了冲击。

Cache的核心作用是加快取用的速度。比如你一个很复杂的计算做完了，下次还要用结果，就把结果放手边一个好拿的地方存着，下次不用再算了。加快了数据取用的速度。

简单来说就是buffer偏重于写，而cache偏重于读。

第三章 简介

• 将CPU即系统的内存称为主机（host），而将GPU及其内存称为设备（device）

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

#include<stdio.h>

__global__ void add(int a,int b,int *c){
  *c = a + b;
}
int main(){
  int c;
  int *dev_c;
  cudaMalloc((void**)&dev_c,sizeof(int));
  add<<<1,1>>>(2,7,dev_c);
  cudaMemcpy(&c,
             dev_c,
             sizeof(int),
             cudaMemcpyDeviceToHost);
  printf("2 + 7 = %d",c);
  cudafree（dev_c）;
  return 0;
}

1.核函数调用

• 1.函数的定义带有了__global__这个标签，表示这个函数是在GPU上运行。函数add()将被交给“编译设备代码的编译器”。需要指出的是尽管是在GPU上执行，但是仍然是由CPU端发起调用的。
• 在每个启动线程中都被调用一遍。
• 2.主机代码发送给一个编译器，而将设备代码发送给另外一个编译器（CUDA编译器）,CUDA编译器运行时将负责实现从主机代码中调用设备代码。
• 3.核函数相对于CPU代码是异步的，也就是控制会在核函数执行完成之前就返回，这样CPU就可以不用等待核函数的完成而继续执行后面的CPU代码
• 4.核函数内部只能访问device内存。因为核函数是执行在设备端，所以只能访问设备端内存。所以要使用cudaMalloc在GPU的内存(全局内存)里开辟一片空间。用来存放结果dev_c。再通过*cudaMemcpy这个函数把内容从GPU**复制出来。

函数部分前缀：

限制：

1. __host__：

   限定符无法一起使用

2. __gobal__：

   限定符无法一起使用；

   函数不支持递归；

   函数的函数体内无法声明静态变量；

   函数不得有数量可变的参数；

   支持函数指针；

   函数的返回类型必须为空；

   函数的调用是异步的，也就是说它会在设备执行完成之前返回；

   函数参数将同时通过共享存储器传递给设备，且限制为 256 字节；

3. __device__：

   函数不支持递归；

   函数的函数体内无法声明静态变量；

   函数不得有数量可变的参数；

   函数的地址无法获取

4. 之前说了__host__和__gobal__限定符无法和其他限定符使用，但与 __device__限定符不是

__constant__ 限定符可选择与 __device__限定符一起使用，所声明的变量具有以下特征：

1.位于固定存储器空间中；2. 与应用程序具有相同的生命周期；3.可通过网格内的所有线程访问，也可通过“运行时库”从主机访问。

__shared__ 限定符可选择与 __device__限定符一起使用，所声明的变量具有以下特征：1.位于线程块的共享存储器空间中；2. 与块具有相同的生命周期；3.尽可通过块内的所有线程访问。只有在_syncthreads()_的执行写入之后，才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量，否则只要之前的语句完成，编译器即可随意优化共享存储器的读写操作。

2.参数传递

• <<<>>>尖括号表示要将一些参数传递给运行时系统，这些参数并不是传递给设备代码的参数，而是告诉运行时如何启动设备代码。传递给设备代码本身的参数是放在圆括号中传递的。

  尖括号作用？线程配置。

  <<<Dg, Db, Ns, S>>>

  1. Dg 的类型为 dim3，指定网格的维度和大小，Dg.x * Dg.y 等于所启动的块数量，Dg.z =1无用，目前还不支持三维的线程格；如果指定Dg=256，那么将有256个线程块在GPU上运行。
  2. Db 的类型为 dim3，指定各块的维度和大小，Db.x * Db.y * Db.z 等于各块的线程数量；
  3. Ns 的类型为 size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为外部数组的其他任何变量使用，Ns 是一个可选参数，默认值为 0；
  4. S 的类型为 cudaStream_t，指定相关流；S 是一个可选参数，默认值为 0。

• 核函数内部可以调用CUDA内置变量，比如threadIdx，blockDim等。下下章将具体谈到线程索引。

• 参数传递和普通函数一样，通过括号内的形参传递。