CUDA基础理解

1、什么是CUDA

简单的理解与我们普通在CPU上编程的区别就是将能够并行加速的部分放到GPU上进行加速，再将加速后的数据传给CPU进行后续处理。 传统CPU只能串行执行，GPU可并行执行。GPU加速部分原理如下图所示

                                    

2、CUDA架构

线程块及线程的组织结构如下图：

                                                           

以线程格（Grid）的形式组织，每个线程格由若干个线程块（block）组成，而每个线程块又由若干个线程（thread）组成。

Grid：线程格 (gridDim.x表示x方向线程格中每一维的线程块数量）
block：线程块（blockDim.x表示x方向线程格中每一维的线程数量）（blockIdx.x 表示当前线程块在当前线程格中x方向的索引值）
thread：线程（threadIdx.x表示当前线程在当前线程块中x方向的索引值）
tid = threadIdx.x + blockIdx.x*blockDim.x（当前线程是索引 = 线程块中线程的索引 + 当前线程块索引*每个线程块中线程的数量）

3、CUDA编程基本步骤

1）分配内存（CPU）空间和显存（GPU）空间，并初始化内存空间

2）将要计算的数据从内存复制到显存上

3）执行GPU上并行处理部分的kernel函数（在GPU上，每个线程都会执行一次核函数）

4）将计算后显存上的数据复制到内存上，对复制到内存上的数据进行后续处理。

4、核函数举例

__global__ void dot(float *a, float *b, float *c)

{

__shared__ float cache[threadsPerBlock];

int tid = threadIdx.x + blockIdx.x * blockDim.x;

int cacheIndex = threadIdx.x;

float temp = 0;

while (tid < N)

{

temp += a[tid] * b[tid];

tid += blockDim.x * gridDim.x;

}

//设置cache中相应位置上的值

cache[cacheIndex] = temp;

//对线程块中的线程进行同步

__syncthreads();

}

dot函数表示的是a(x1, x2,.....xn)*b(y1, y2,.........yn)=x1*y1+x2*y2+.....+xn*yn。若此时线程为第一个线程中第一个线程块的第一个元素。则tid += blockIdx.x*blockDim.x表示第二个线程格中第一个线程块的第一个元素。以此类推，每个线程格都相当于同样的表格。执行完上段程序后，所有的值将放在第一个线程格中，再对这个线程格中的所有值进行求和。这是《GPU高性能编程CUDA实战里》里第五章的代码，完整代码就不给出了。

5、基本函数

cudaMalloc((void**)&dev_a,sizeof(int)); 在设备上分配内存，后面参数为分配内存大小

cudaFree(dev_a); 释放设备上的内存。

__syncthreads();对线程块中的所有线程进行同步，职业线程块中所有线程执行了它，才会执行下一条语句，否则将一直等待。

__shared__,变量保存在共享内存缓存区。为每个线程块生成共享变量的一个副本，根据每个线程块中线程的数量分配内存。

__global__,被修饰的函数在GPU上执行，在CPU上被调用。

__device__,被修饰的函数在GPU上执行，可在__global__和__device__函数中被调用