http://www.hyperparameter.cn/science-high-performance-computation/cuda-basic-model/
CUDA简介
CUDA(Compute Unified Device Architecture,统一计算架构)是由NVIDIA所推出的一种集成技术,通过这个技术,用户可以使用显卡中的资源进行大规模并行计算。
为了后续CUDA编程的展开,该系列的第一篇首先从N系显卡的物理结构以及CUDA编程中的基本模型开始。
物理结构
Nvidia公司开发的GPU系列现在已经有Tesla、Fermi、Kepler、Maxwell、Pascal、Volta等多种GPU架构。下图就是Pascal架构中GP100(关于显卡核心的命名,如果是Kepler就是GK+数字,如果是Pascal就是GP+数字)的基本构造,其中包括GPC(图形处理簇,Graphics Processing Clusters)、TPC(纹理处理簇,Texture Processing Clusters)、SM(流多处理器,Stream Multiprocessors)以及内存控制器。SM中的一个个小格子就代表一个个可调用计算的线程(绿色代表单精度,黄色代表多精度),这里就可以充分展现GPU与CPU之间的架构差异,一个是做大量重复性工作,另一个则是专注于逻辑计算。
下图对于单个的SM结构描述得更清楚。
编程模型
编程模型中,以主机-设备(显卡)的形式,通过主机调用核函数启用显卡中的资源进行计算,
KernelFunction<<<dimGrid, dimBlock>>>(param1, param2,...);
这就牵涉到如何组织线程进行计算了,CUDA里面主要以网格-线程块-线程进行组织。其中,
dim3 dimGrid(x, y, z);
dim3 dimBlock(x, y, z);
dimGrid声明了线程块的三维组织方式(x, y, z)以及线程块数量xyz,dimBlock则声明了线程块中线程的组织方式(x, y, z)以及每个线程块中的线程总数xyz,下图是该组织方式的体现。
下面介绍这个系列里的第一个程序,对于当前主机上的CUDA资源的统计展示。
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
//用于查看显卡信息
int main(int argc, char ** argv)
{
cudaError_t cudaStatus;
int num = 0;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&num);//查看总共有多少张显卡
for (int i = 0; i<num; i++)
{
cudaGetDeviceProperties(&prop, i);//读出每张显卡的资源参数
}
return 0;
}
具体输出我贴在最后面了,略多。这里主要需要解释CUDA相关的几个资源参数。
- maxThreadsDim 表示(x, y, z)各个维度上的最大线程数
- maxGridSize 表示(x, y, z)各个维度上的最大块数量
- maxThreadsPerBlock 表示每个数据块中的最大线程数
- multiProcessorCount 表示SM的数量
- sharedMemPerBlock 表示每个数据块中的共享内存大小
- warpSize 表示一个warp单元中划分的线程数量,其中warp表示硬件实现中的线程调度单元,一个warp包含threadIdx连续的线程
关于warp,CUDA中按照SIMD模式执行一个warp,也就是任何时刻,warp中所有线程只能取一条指令执行,任意时刻SM只能执行所有warp中的一部分。之所以要设计warp这个部分,是考虑到当一个warp中的指令需要等待先前启动的长延迟操作的结果时,就不会选择该warp。,因此需要调度其他warp来执行掩盖延迟。有点类似于银行办手续,先让你在一边填写表格,之后再办理业务。
因此,在CUDA编程中如何合理安排线程是一个首先需要考虑的问题。假定一个设备支持8个线程块、每个线程块支持1024个线程,每个线程块支持512个线程。对于矩阵乘法中,如何抉择8*8,16*16以及32*32大小的线程块呢。如果选择8*8,则分配的线程块数量会超过,如果选择32*32大小这回超过线程块的线程最大量。
这里再贴一张图,可以比较清楚地看出各代显卡核心之间计算能力的区别(我的笔记本显卡太渣了~ ~)。
cuda_Device_Properties
名称 | 值 | 类型 |
---|---|---|
name | “GeForce GT 640M” | char[256] |
totalGlobalMem | 2147483648 | unsigned int |
sharedMemPerBlock | 49152 | unsigned int |
regsPerBlock | 65536 | int |
warpSize | 32 | int |
memPitch | 2147483647 | unsigned int |
maxThreadsPerBlock | 1024 | int |
maxThreadsDim | 0x0020f8c0 {1024, 1024, 64} | int[3] |
maxGridSize | 0x0020f8cc {2147483647, 65535, 65535} | int[3] |
clockRate | 708500 | int |
totalConstMem | 65536 | unsigned int |
major | 3 | int |
minor | 0 | int |
textureAlignment | 512 | unsigned int |
texturePitchAlignment | 32 | unsigned int |
deviceOverlap | 1 | int |
multiProcessorCount | 2 | int |
kernelExecTimeoutEnabled | 1 | int |
integrated | 0 | int |
canMapHostMemory | 1 | int |
computeMode | 0 | int |
maxTexture1D | 65536 | int |
maxTexture1DMipmap | 16384 | int |
maxTexture1DLinear | 134217728 | int |
maxTexture2D | 0x0020f914 {65536, 65536} | int[2] |
maxTexture2DMipmap | 0x0020f91c {16384, 16384} | int[2] |
maxTexture2DLinear | 0x0020f924 {65000, 65000, 1048544} | int[3] |
maxTexture2DGather | 0x0020f930 {16384, 16384} | int[2] |
maxTexture3D | 0x0020f938 {4096, 4096, 4096} | int[3] |
maxTexture3DAlt | 0x0020f944 {2048, 2048, 16384} | int[3] |
maxTextureCubemap | 16384 | int |
maxTexture1DLayered | 0x0020f954 {16384, 2048} | int[2] |
maxTexture2DLayered | 0x0020f95c {16384, 16384, 2048} | int[3] |
maxTextureCubemapLayered | 0x0020f968 {16384, 2046} | int[2] |
maxSurface1D | 65536 | int |
maxSurface2D | 0x0020f974 {65536, 32768} | int[2] |
maxSurface3D | 0x0020f97c {65536, 32768, 2048} | int[3] |
maxSurface1DLayered | 0x0020f988 {65536, 2048} | int[2] |
maxSurface2DLayered | 0x0020f990 {65536, 32768, 2048} | int[3] |
maxSurfaceCubemap | 32768 | int |
maxSurfaceCubemapLayered | 0x0020f9a0 {32768, 2046} | int[2] |
surfaceAlignment | 512 | unsigned int |
concurrentKernels | 1 | int |
ECCEnabled | 0 | int |
pciBusID | 1 | int |
pciDeviceID | 0 | int |
pciDomainID | 0 | int |
tccDriver | 0 | int |
asyncEngineCount | 1 | int |
unifiedAddressing | 0 | int |
memoryClockRate | 900000 | int |
memoryBusWidth | 128 | int |
l2CacheSize | 262144 | int |
maxThreadsPerMultiProcessor | 2048 | int |
streamPrioritiesSupported | 0 | int |
globalL1CacheSupported | 0 | int |
localL1CacheSupported | 1 | int |
sharedMemPerMultiprocessor | 49152 | unsigned int |
regsPerMultiprocessor | 65536 | int |
managedMemory | 0 | int |
isMultiGpuBoard | 0 | int |
multiGpuBoardGroupID | 0 | int |
hostNativeAtomicSupported | 0 | int |
singleToDoublePrecisionPerfRatio | 24 | int |
pageableMemoryAccess | 0 | int |
concurrentManagedAccess | 0 | int |
——————
资料来源:
[1]. pascal-architecture-whitepaper
[2]. 《大规模并行处理器编程实战》
本作品采用知识共享署名-非商业性使用-相同方式共享 3.0 中国大陆许可协议进行许可。