CUDA系列学习（一）An Introduction to GPU and CUDA

本文从软硬件层面讲一下CUDA的结构，应用，逻辑和接口。分为以下章节：

（一）、GPU与CPU

（二）、CUDA硬件层面

（三）、CUDA安装

（四）、CUDA 结构与接口

4.1 Kernels

4.2 Thread，Block, Grid

4.3 Memory

4.4 Execution

（五）、码HelloWorld——数组求和

希望感兴趣的同学可以一起讨论。

(一)、GPU与CPU

对于浮点数操作能力，CPU与GPU的能力相差在GPU更适用于计算强度高，多并行的计算中。因此，GPU拥有更多晶体管，而不是像CPU一样的数据Cache和流程控制器。这样的设计是因为多并行计算的时候每个数据单元执行相同程序，不需要那么繁琐的流程控制，而更需要高计算能力，这也不需要大cache。

（二）、CUDA硬件层面：

Nvidia于2006年引入CUDA，一个GPU内嵌通用并行计算平台。CUDA支持C, C++, Fortran, Java, Python等语言。

那么一个多线程CUDA程序如何执行的呢？

GPU建立在一组多处理器（SMX，Streaming Multiprocessors）附近。

一个SMX的配置：

192 cores（都是SIMT cores（Single Instruction Multiple Threads） and 64k registers（如下图所示）

GPU中的SIMT对应于CPU中的SIMD（Single Instruction Multiple Data）

64KB of shared memory / L1 cache
8KB cache for constants
48KB texture cache for read-only arrays
up to 2K threads per SMX

不同显卡有不同配置（即SMX数量不同），几个例子：

每个multi-thread程序的execution kernel instance（kernel定义见下一节,instance指block）在一个SMX上执行，一个多线程程序会分配到blocks of threads（每个block中负责一部分线程）中独立执行。所以GPU中的处理器越多执行越快（因为如果SMX不够给每个kernel instance分配一个，就要几个kernel抢一个SMX了）。具体来讲，如果SMX上有足够寄存器和内存（后面会讲到，shared memory），就多个kernel instance在一个SMX上执行，否则放到队列里等。

图：表示不同SM数带来的执行速度差异。

GPU工作原理：首先通过主接口读取中央处理器指令，GigaThread引擎从系统内存中获取特定的数据并拷贝到显存中，为显存控制器提供数据存取所需的高带宽。GigaThread引擎随后为各个SMX创建和分派线程块（warp, 详细介绍见SIMT架构或者CUDA系列学习（二）），SMX则将多个Warp调度到各CUDA核心以及其他执行单元。在图形流水线出现工作超载的时候，GigaThread引擎还负责进行工作的重新分配。

（三）、CUDA安装

装CUDA主要装以下3个组建：

1. driver

low-level software that controls the graphics card

2. toolkit

nvcc CUDA compiler
profiling and debugging tools
several libraries

3. SDK

lots of demonstration examples
some error-checking utilities
not officially supported by NVIDIA
almost no documentation

详情请见CUDA 安装与配置

（四）、CUDA 结构与接口

4.1 Kernels

CUDA C 中可通过定义kernel，每次被调用就在N个CUDA thread中并行执行。

kernel的定义：

声明 __global__
配置kernel_routine<<<gridDim, Blockdim>>>(args)

其中gridDim和Blockdim变量可以是int或dim3（<=3维）类型的变量。gridDim表示每个grid中block结构(the number of instances(blocks) of the kernel)，Blockdim表示每个block中thread结构。那么。。thread，block，grid又是啥？往下看。。。见4.2节

每个执行该kernel的thread都会通过被分配到一个unique thread ID，就是built-in变量：threadIdx

4.2 Thread，Block，Grid

很多threads组成1维，2维or3维的thread block. 为了标记thread在block中的位置（index），我们可以用上面讲的threadIdx。threadIdx是一个维度<=3的vector。还可以用thread index（一个标量）表示这个位置。

thread的index与threadIdx的关系：

	Thread index
1	T
2	T.x + T.y * Dx
3	T.x+T.yDx+zDx*Dy

其中T表示变量threadIdx。（Dx, Dy, Dz）为block的size（每一维有多少threads）。

因为一个block内的所有threads会在同一处理器内核上共享内存资源，所以block内有多少threads是有限制的。目前GPU限制每个 block最多有1024个threads。但是一个kernel可以在多个相同shape的block上执行，效果等效于在一个有N*#thread per block个thread的block上执行。

Block又被组织成grid。同样，grid中block也可以被组织成1维，2维or3维。一个grid中的block数量由系统中处理器个数或待处理的数据量决定。

和threadIdx类似，对于block有built-in变量blockDim（block dimension）和blockIdx（block index）。

回过头来看4.1中的configureation，举个栗子，假设A，B，C都是大小[N][N]的二维矩阵，kernel MatAdd目的将A，B对应位置元素加和给C的对应位置。

声明：

[cpp] view plain copy

// Kernel definition
__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 invocation
dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}

这里threadsPerBlock(16,16)一般是标配。例子中，我们假定grid中block足够多，保证N/threadsPerBlock不会超限。

4.3 Memory

前面提到了Block中的threads共享内存，那么怎样同步呢？在kernel中调用内部__synthreads()函数，其作用是block内的所有threads必须全部执行完，程序才能继续往下走。那么thread到底怎样使用memory呢？

每个thread有private local memory
每个block有shared memory
所有thread都能访问到相同的一块global memory
所有thread都能访问两块read-only memory：constant & texture array(通常放查找表)

其中，global，constant，texture memory伴随kernel生死。

CUDA程序执行的时候，GPU就像一个独立的设备一样，kernel部分由GPU执行，其余部分CPU执行。于是memory就被分为host memory（for CPU）& device memory（for GPU）。因此，一个程序需要在CUDA运行时管理device memory的分配，释放和device & host memory之间的data transfer。

4.4 Execution

从执行角度看，程序经过了以下步骤：

1. initialises card
2. allocates memory in host and on device
3. copies data from host to device memory
4. launches multiple instances of execution “kernel” on device
5. copies data from device memory to host
6. repeats 3-5 as needed
7. de-allocates all memory and terminates

总结：每个kernel放在一个grid上执行，1个kernel有多个instance，每个instance在一个block上执行，每个block只能在一个SM上执行，如果block数>SM数，多个block抢SM用。kernel的一个instance在SMX上通过一组进程来执行。如下图所示：

总结：

CUDA的3个key abstraction：thread groups, shared memories, 和barrier synchronization

CUDA中的built-in变量：gridDim, blockDim, blockIdx（block在grid中的index）, threadIdx, warpSize（threads的warp size）

（五）、码HelloWorld

kernel code很像MPI，从单线程的角度coding
需要think about每个变量放在哪块内存

这里我们以数组对应元素相加为例，看下Code ：

[cpp] view plain copy

#include<cutil_inline.h>
#include<iostream>
using namespace std;
#define N 32
// Kernel definition
__global__ void MatAdd(float A[N], float B[N], float* C)
{
int i = blockIdx.x * blockDim.x + threadIdx.x; //get thread index by built-in variables
if (i < N)
C[i] = A[i] + B[i];
}
int main()
{
float A[N],B[N]; // host variable
float *dA, *dB; // device variable, to have same value with A,B
float *device_res, *host_res; // device and host result, to be device and host variable respectively
// initialize host variable
memset(A,0,sizeof(A));
memset(B,0,sizeof(B));
A[0] = 1;
B[0] = 2;
// allocate for device variable and set value to them
cudaMalloc((void**) &dA,N*sizeof(float));
cudaMalloc((void**) &dB,N*sizeof(float));
cudaMemcpy(dA, A, N*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(dB, B, N*sizeof(float),cudaMemcpyHostToDevice);
//malloc for host and device variable
host_res = (float*) malloc(N*sizeof(float));
cudaMalloc((void**)&device_res, N*sizeof(float));
// Kernel invocation
int threadsPerBlock = 16;
int numBlocks = N/threadsPerBlock;
MatAdd<<<numBlocks, threadsPerBlock>>>(dA, dB, device_res);
cudaMemcpy(host_res, device_res, N*sizeof(float),cudaMemcpyDeviceToHost); //copy from device to host
// validate
int i;
float sum = 0;
for(i=0;i<N;i++)
sum += host_res[i];
cout<<sum<<endl;
//free variables
cudaFree(dA);
cudaFree(dB);
cudaFree(device_res);
free(host_res);
}

编译：

nvcc -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc/ Matadd.cu
运行结果：

OK，大功告成。

这里注意kernel部分的code，所有变量都必须是device variable，即需要通过cudaMalloc分配过memory的。之前我忘记将A，B数组cudaMemcpy到dA,dB，而直接传入MatAdd kernel就出现了运行一次过后卡住的问题。

参考：

1. CUDA C Programming Guide

2. An Introduction to CUDA

3. CUDA 安装与配置

4. CUDA调试工具——CUDA GDB

5. GPU工作方式

6. Fermi 架构白皮书（GPU继承了Fermi的很多架构特点）

7. GTX460架构

from: http://blog.csdn.net/abcjennifer/article/details/42436727

CUDA系列学习（一）An Introduction to GPU and CUDA相关推荐

[tensorflow]各个tensorflow版本和CUDA版本对应，以及各个GPU版本CUDA和cuDNN对应
目录各个CPU版本tensorflow对应的环境要求各个GPU版本tensorflow对应的CUDA版本各个版本的CUDA和英伟达显卡驱动对应表缺失cudnn64_7.dll文件查看本地CU ...
CUDA系列学习（三）GPU设计与结构QA coding练习
啥?你把CUDA系列学习(一),(二)都看完了还不知道為什麼要用GPU提速? 是啊..经微博上的反馈我默默感觉到提出这样问题的小伙伴不在少数,但是更多小伙伴应该是看了(一)就感觉离自己太远所以赶紧撤粉 ...
DL之IDE：深度学习环境安装之CUDA的简介(显卡GPU/驱动/CUDA间的关系)、安装(根据本地电脑的NVIDIA显卡驱动版本去正确匹配CUDA版本)之详细攻略
DL之IDE:深度学习环境安装之CUDA的简介(显卡GPU/驱动/CUDA间的关系).安装(根据本地电脑的NVIDIA显卡驱动版本去正确匹配CUDA版本)之详细攻略目录 CUDA的简介 1.显卡GP ...
CUDA系列学习（五）GPU基础算法: Reduce, Scan, Histogram
喵~不知不觉到了CUDA系列学习第五讲,前几讲中我们主要介绍了基础GPU中的软硬件结构,内存管理,task类型等:这一讲中我们将介绍3个基础的GPU算法:reduce,scan,histogram,它 ...
CUDA系列学习（四）Parallel Task类型与 Memory Allocation
本文为CUDA系列学习第四讲,首先介绍了Parallel communication patterns的几种形式(map, gather, scatter, stencil, transpose), ...
目的：使用CUDA环境变量CUDA_VISIBLE_DEVICES来限定CUDA程序所能使用的GPU
目的:使用CUDA环境变量CUDA_VISIBLE_DEVICES来限定CUDA程序所能使用的GPU 目的:使用CUDA环境变量CUDA_VISIBLE_DEVICES来限定CUDA程序所能使用的GP ...
conda创建虚拟环境和用conda创建GPU的cuda、cudnn使用环境
conda创建虚拟环境和用conda创建GPU的cuda.cudnn使用环境 1 conda在linux.windows上创建虚拟环境 1.1 首先在所在系统中安装Anaconda. 1.2 co ...
GPU（CUDA）学习日记（十一）------ 深入理解CUDA线程层次以及关于设置线程数的思考
GPU(CUDA)学习日记(十一)------ 深入理解CUDA线程层次以及关于设置线程数的思考标签: cuda存储线程结构网格 2012-12-07 16:30 6298人阅读评论(4)收藏举 ...
GPU（CUDA）学习日记（十三）------ CUDA内存简介
GPU(CUDA)学习日记(十三)------ CUDA内存简介标签: cuda存储线程结构 2012-12-07 16:53 2902人阅读评论(0)收藏举报分类: GPU(16) CUDA ...

CUDA系列学习（一）An Introduction to GPU and CUDA

CUDA系列学习（一）An Introduction to GPU and CUDA相关推荐

最新文章

热门文章