CUDA——基本模型

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]
totalGlobalMem2147483648unsigned int
sharedMemPerBlock49152unsigned int
regsPerBlock65536int
warpSize32int
memPitch2147483647unsigned int
maxThreadsPerBlock1024int
maxThreadsDim0x0020f8c0 {1024, 1024, 64}int[3]
maxGridSize0x0020f8cc {2147483647, 65535, 65535}int[3]
clockRate708500int
totalConstMem65536unsigned int
major3int
minor0int
textureAlignment512unsigned int
texturePitchAlignment32unsigned int
deviceOverlap1int
multiProcessorCount2int
kernelExecTimeoutEnabled1int
integrated0int
canMapHostMemory1int
computeMode0int
maxTexture1D65536int
maxTexture1DMipmap16384int
maxTexture1DLinear134217728int
maxTexture2D0x0020f914 {65536, 65536}int[2]
maxTexture2DMipmap0x0020f91c {16384, 16384}int[2]
maxTexture2DLinear0x0020f924 {65000, 65000, 1048544}int[3]
maxTexture2DGather0x0020f930 {16384, 16384}int[2]
maxTexture3D0x0020f938 {4096, 4096, 4096}int[3]
maxTexture3DAlt0x0020f944 {2048, 2048, 16384}int[3]
maxTextureCubemap16384int
maxTexture1DLayered0x0020f954 {16384, 2048}int[2]
maxTexture2DLayered0x0020f95c {16384, 16384, 2048}int[3]
maxTextureCubemapLayered0x0020f968 {16384, 2046}int[2]
maxSurface1D65536int
maxSurface2D0x0020f974 {65536, 32768}int[2]
maxSurface3D0x0020f97c {65536, 32768, 2048}int[3]
maxSurface1DLayered0x0020f988 {65536, 2048}int[2]
maxSurface2DLayered0x0020f990 {65536, 32768, 2048}int[3]
maxSurfaceCubemap32768int
maxSurfaceCubemapLayered0x0020f9a0 {32768, 2046}int[2]
surfaceAlignment512unsigned int
concurrentKernels1int
ECCEnabled0int
pciBusID1int
pciDeviceID0int
pciDomainID0int
tccDriver0int
asyncEngineCount1int
unifiedAddressing0int
memoryClockRate900000int
memoryBusWidth128int
l2CacheSize262144int
maxThreadsPerMultiProcessor2048int
streamPrioritiesSupported0int
globalL1CacheSupported0int
localL1CacheSupported1int
sharedMemPerMultiprocessor49152unsigned int
regsPerMultiprocessor65536int
managedMemory0int
isMultiGpuBoard0int
multiGpuBoardGroupID0int
hostNativeAtomicSupported0int
singleToDoublePrecisionPerfRatio24int
pageableMemoryAccess0int
concurrentManagedAccess0int

——————

资料来源:
[1]. pascal-architecture-whitepaper
[2]. 《大规模并行处理器编程实战》


知识共享许可协议
本作品采用知识共享署名-非商业性使用-相同方式共享 3.0 中国大陆许可协议进行许可。

  • 0
    点赞
  • 2
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值