CUDA——基本模型

点击量:35

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个线程。对于矩阵乘法中,如何抉择88,1616以及3232大小的线程块呢。如果选择88,则分配的线程块数量会超过,如果选择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 中国大陆许可协议进行许可。

发表评论

电子邮件地址不会被公开。 必填项已用*标注