CUDA basic knowledge

物理概念、CUDA基本概念、CUDA语法、调试器、错误处理

SSE增加数据吞吐(128位)

OpenMP增加线程数,所以两个一起使用能使CPU程序最大化

物理概念

本人设备:

RTX 2060 、GTX1080TI

设备 RTX2060 GTX1080TI
架构 图灵 pascal
核心数(cuda core)/SP 1920 3584
核心频率 1365MHz 1480MHz
纹理单元 120 224
Tensor核心 240
RT核心 30
ROPs单元 48 88
显存大小 6GB 11GB
显存位宽为 192bit 352bit
显存带宽 336GB/s 484GB/s

484GB/s = 495616MB/s=495.616MB/s。理论上能在1ms内把500MB的数送进流处理器。

GPU实际上是一个SM阵列(GTX1080TI有28个SM)(RTX2060有30个SM),SM每个SM中包含N个核(cuda core)也叫SP(1080TI每个SM中有128个sp))(RTX2060每个SM有64个sp)。1080TI有3584个流处理器(SP),RTX2060有1920个流处理器(SP)。

SM

每个SM都有一个寄存器文件(Register FIle),这是一组能够以与SP相同速度工作的存储单元,所以访问这组单元不需要任何等待时间。它用来存储SP上运行线程内部活跃的寄存器。

每个SM有一个内部访问的共享内存(shared memory),可以用作程序可控的高速缓存。

SM每次可同时计算32个数而不是像CPU那样只计算一个数。

内存

CPU和GPU有各自独立的内存空间,因此在GPU代码中,不可以直接访问CPU端的代码。反过来同理。

对于纹理内存、常量内存,全局内存,每一个SM都分别设置有独立访问它们的总线。

纹理内存内存是针对全局内存的一个特殊视图,用来存储插值(interpolation)计算所需的数据。(基于硬件进行插值)

常量内存用于存储那些只读的数据。常量内存也是全局内存建立的一个视图


原子操作是指那些必须一次性完成、不会被其他线程中断的操作

CUDA概念

设备CUDA参数使用自带例程deviceQuery查询

页锁定主机内存

测试demo

CUDA页锁定主机内存(主机的固定内存),好处固定内存复制到显存,或显存复制到固定内存快。

使用页锁定内存也需要申请GPU内存和主机内存啊,只不过使用cudaHostAlloc()申请主机内存的时候,申请的主机内存是不可分页的(即申请的是固定内存,使用物理地址访问),还是需要使用cudaMemcpy()来将主机内存中的数据拷贝到GPU内存中。好处就是主机和GPU之间复制数据速率更快.主机内存到显存,显存到主机内存都有加快,大概两倍。缺点就是申请的是物理内存,不是虚拟内存,内存很快就会耗尽,所以一般都会用完就释放

线程

GPU的每个线程组被送到SM中,然后N个SP开始执行代码。事实上,线程都是以每32个为一组。这个线程组叫线程束。以线程束为单位进行操作。

如果一个block中有128个线程,每个线程传递3个参数,那么就需要3X128=384个寄存器。这听起来很多,但是其实每个SM(流处理器)中至少有8192个寄存器。就是说每个线程可以使用64个寄存器,只有当运算前度很强的时候才使用64个寄存器。

线程块

线程网格

一个线程网格是由若干个线程块组成的,每个线程块是二维的,拥有X轴与Y轴。

由于数组的存在,数据往往不可能是一维的。这时我们可以使用二维的线程块。

线程束(大小为32)

线程束是GPU的基本执行单元。GPU是一组SIMD向量处理器的集合。每一组线程或每个线程束中的线程同时执行。在理想状态下,获取当前指令只需要一次访存,然后将这个指令广播到这个线程束所占用的所有SP中。

GPU利用率

注意事项:

  1. 每个线程块开启的线程数越多,就潜在地增加了等待执行较慢线程束的可能性。因为当所有线程没有到达同步点时GPU是无法继续向下执行的。因此,有时候我们会选择在每个线程块上开启较少的线程,例如,128个线程,以此来减少之前那种等待的可能性。但是这样做会严重影响性能。

CUDA语法

1
2

CUDA内核

kernel_function<<<num_blocks,num_threads>>>(param1,param2,…)




# Nsight调试器



# 错误处理
# CUDA

评论

Your browser is out-of-date!

Update your browser to view this website correctly. Update my browser now

×