CUDA的基本概念：



CUDA是NVIDIA公司推出的并行计算框架，只能基于自家GPU的硬件平台进行性能加速运算，所以使用CUDA编程的前提是必须要有NVIDIA的显卡。
主流的深度学习框架也都是基于CUDA进行GPU并行加速的，几乎无一例外。还有一个叫做cudnn，是针对深度卷积神经网络的加速库。
与之相类似的并行计算框架还有苹果公司推出的OpenCL，OpenCL其优势在于跨平台性和通用性，更像是一个开放标准，按理说OpenCL会更加受欢迎，但是由于NVIDIA公司在GPU显卡领域一家独大，市场份额独占鳌头，使得CUDA的生态支持等方面更加完善，因此在短期以及未来的相当一段时间，并行加速领域还是CUDA更受欢迎一些。

学习内容：

1.CUDA的组成：

1.1、CPU、GPU与kernel
虽然CUDA平台的并行计算是在GPU上完成的，但是并不仅仅依赖于GPU，还需要和CPU分工配合才能实现。在这里CPU是（Host）主机端，GPU属于(Device)设备端，两者通过kernel（核函数）连接，在CUDA程序构架中，主程序还是由CPU来执行，而当遇到数据并行处理的部分（即kernel），CUDA 就会将程序上传送到GPU端运行。kernel用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行及结构。核函数返回类型是void。

__global__ void hello_world_from_gpu(void)
{
	printf("Hello World from GPU\n");
	return;
}

1.2、Grid、Block与Thread

Grid、Block与Thread这些概念简单来说就是GPU在运算执行过程中计算资源分配的单位，大白话说就是你打算用多少GPU的资源去执行并行计算。所谓网格 (grid)，其实就是线程块（Block）的组合体，而Block中又包含了很多的Thread线程，按包含关系来说Grid>Block>Thread，结构如下图。

体现在kerne函数中，就是指在CPU调用kernel函数时，在kernel函数后面紧跟着的<<<grid, block>>>中指定Grid、Block的分配情况。

比如前面定义了一个kernel函数hello_world_from_gpu，在CPU调用时需要写成hello_world_from_gpu <<<Grid, Block >>> ();Grid和Block就是你要分配的资源线程数，具体怎么分配GPU的线程是很有讲究的，不是说随便制定一个数字都行的，下面会提到注意事项。

每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个 thread 则有共享的一份 share memory。此外，所有的 thread (包括不同 block 的 thread) 都共享一份 global memory。不同的 grid 则有各自的 global memory。

在cuda训练的权重不能放到cpu上吗 cuda重要吗_cuda

1.3、SM与Warp
当一个 kernel 被执行时，grid 中的线程块被分配到 SM (多核处理器) 上，SM是GPU上真正的物理执行单元，一个线程块的 thread 只能在一个SM 上调度，SM 一般可以调度多个线程块，大量的 thread 可能被分到不同的 SM 上。
warp (线程束) 是最基本的执行单元。一个 warp 包含32个并行 thread，这些 thread 以不同数据资源执行相同的指令。所以warp 本质上是线程在 GPU 上运行的最小单元。由于warp的大小为32，所以block所含的thread的大小一般要设置为32的倍数。

2.CUDA程序的运行步骤：

2.1：从主机 (host) 端申请 device memory，把要拷贝的内容从 host memory 拷贝到申请的 device memory 里面。

在CPU上申请显存的函数 malloc()：
float* a = (float*)malloc(sizeof(float) * size * size);
在Device端申请显存的函数 cudaMalloc()：
cudaMalloc((void**)&a_cuda, sizeof(float) * size * size);
把要拷贝的内容从 host memory 拷贝到申请的 device memory 里面:
cudaMemcpy(a_cuda, a, sizeof(float) * size * size, cudaMemcpyHostToDevice);

2.2：设备端的核函数对拷贝进来的东西进行计算，来得到和实现运算的结果，Kernel 就是指在 GPU 上运行的函数。

//核函数
__global__ void kernel(int a, int b, int* c)
{
	*c = a * (b+1);
}

2.3：把结果从 device memory 拷贝到申请的 host memory 里面，并且释放设备端的显存和内存

cudaMemcpy(&c, device_c, sizeof(int), cudaMemcpyDeviceToHost);   //将设备端内存拷贝到主机端

因此一个完整的CUDA程序实例如下（.cu格式的文件）:

#include <iostream>
#include "cuda_runtime.h"
using namespace std;

//核函数
__global__ void kernel(int a, int b, int* c)
{
	*c = a * (b+1);
}

int main()
{
	int c;
	int* device_c;
	cudaMalloc((void**)&device_c, sizeof(int));  //GPU设备端分配和申请内存
	kernel <<< 1, 1 >> > (2, 7, device_c);	//kernel函数在设备端运算
	cudaMemcpy(&c, device_c, sizeof(int), cudaMemcpyDeviceToHost);   //将设备端内存拷贝到主机端
	printf("a * (b+1)=%d\n", c);
	cudaFree(device_c);
	system("pause");
	return 0;
}

从上面完整的CUDA程序可以看出其格式结构和cpp文件基本一致，只不过这是一个.cu的格式，并且多了类似cuda_runtime.h这样的头文件。定义的kernel函数在main函数中进行调用。

kernel <<< 1, 1 >> > (2, 7, device_c);

注意核函数只能在主机端调用，调用时必须申明执行参数。调用形式如下：

Kernel<<<grid,block, Ns, S>>>(param list);

调用通过<<<grid,block>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。
参数 Ns 是一个可选参数，用于设置每个 block 除了静态分配的 shared Memory 以外，最多能动态分配的shared memory 大小，单位为 byte。不需要动态分配时该值为0或省略不写，一般可默认不写。
参数 S 是一个 cudaStream_t 类型的可选参数，初始值为零，表示该核函数处在哪个流之中，一般可默认不写。
由于Grid和Block都可以从1维定义到3维，那么在不同维度下Thread的索引ID怎么唯一确定呢？建议仔细阅读一下这篇文章，看完绝对豁然开朗！
【CUDA】grid、block、thread的关系及thread索引的计算

更多详情，请参考官方CUDA编程指南