cuda学习笔记

By findhao at 2017-10-09 • 0人收藏 • 231人看过

整理一些内容,比较多以后再发到博客。


一些cuda运行时API在主机和设备之间是隐式同步的。比如cudaMemcpy,主机端程序必须等待数据拷贝完成后才能继续执行。


25 个回复 | 最后更新于 2017-10-20
2017-10-09   #1

所有cuda核函数的启动都是异步的。

2017-10-09   #2

CUDA C编程权威指南 P32 处理错误。可以定义错误处理宏封装所有CUDA API调用。

2017-10-09   #3

把线程和块索引映射到矩阵坐标上

ix = threadIdx.x + blockIdx.x * blockDim.x

iy = threadIdx.y + blockIdx.y * blockDim.y


idx = iy * nx + ix

2017-10-09   #4

线程块一旦被调度到一个SM上,其中的线程只会在那个指定的SM上并发执行。

2017-10-09   #5

SIMT(单指令多线程)架构和SIMD相似,都是将相同的指令广播给多个执行单元来实现并行,一个关键的区别是SIMD要求同一个向量中的所有元素要在一个统一的同步组中一起执行,而SIMT允许同一线程束的多个线程独立执行。SIMT确保可以编写独立的线程级并行代码、标量线程以及用于协调线程的数据并行代码。

SIMT模型包含3个SIMD所不具备的关键特征:

  1. 每个线程都有自己的指令地址计数器

  2. 每个线程都有自己的寄存器状态

  3. 每个线程可以有一个独立的执行路径

2017-10-09   #6

CUDA C编程权威指南 P65页 表3-1 计算能力概览

2017-10-09   #7

在设备上第一次运行可能会增加间接开销。

2017-10-09   #8

3.2.4是延迟隐藏,需要重新学习。

2017-10-10   #9

线程束分化问题:在kernel函数中,如果有类似if tid % 2 == 0 do something的语句,那么将有一半线程执行else,如果没有else,则其将空闲。 3.4.3 

nvprof --metrics inst_per_warp ./XXXX

每个线程束里执行的指令数,优化后越多,则每个线程束做的任务越多,一般情况下,表示代码写的越好。

2017-10-10   #10

归约求和问题大概讲了以下优化方法:

  1. 相邻两个合并

  2. 在1的基础上,使得每次工作线程id连续,避免线程束分化

  3. 交错配对,初始跨度是线程块大小一半,然后折半。主要是利用访存的请求合并进行优化

  4. 循环展开。主要是编译器执行循环展开时低级指令的改进和优化

  5. 预处理多一个数据

  6. 最后剩余32个或者更少线程时的迭代可以直接完全展开。此时注意使用volatile修饰变量。


2017-10-10   #11

检测内存加载和存储效率指标:

nvprof --metrics gld_efficiency,gst_efficicency ./XXXX


2017-10-10   #12

   

cuda内存模型:

一个核函数中的线程都有自己的私有本地内存,一个线程块有自己的共享内存,对同一个线程块中的所有线程可见,其内容持续线程块的整个生命周期。所有线程都能访问全局内存。

2017-10-10   #13

核函数中声明的一个没有其他修饰符的自变量,通常存储在寄存器中。

在核函数声明的数组中,若索引是常量且编译时确定,则该数组也在寄存器中。寄存器变量对于每个线程来说都是私有的。

不同架构GPU,每个线程可以拥有的寄存器数量不同。


寄存器溢出。https://www.findhao.net/easycoding/1682 

2017-10-10   #14

共享内存是片上内存,类似CPU的一级缓存,但是可编程。__shared__

本地内存本质上和全局内存是同一块存储区域。高延迟,低带宽。

SM中的一级缓存和共享内存都使用64KB的片上内存,静态划分或者运行时动态配置。

常量内存:静态声明,对同一编译单元中的所有核函数可见。必须在主机端初始化。如果线程束里的每个线程都从不同地址空间读取数据,并且只读一次,那么常量内存就不是最佳选择,因为每从一个常量内存读取一次数据,都会广播给线程束的所有线程。

全局内存:__device__ 内存事务必须自然对其,即首地址必须是32,64,128字节的倍数。

各类存储器的重要特征  表4-2 P122

2017-10-10   #15

零拷贝内存 https://www.findhao.net/easycoding/1448 

  1. 当设备内存不足时,可利用主机内存。

  2. 避免主机和设备间的显示数据传输

  3. 提高PCIe传输率

在进行频繁的读写操作时,使用零拷贝内存将显著降低性能。因为每一次映射到内存的传输都必须经过PCIe总线。

如果CPU和GPU共享内存,零拷贝在性能和可编程性方面可能更佳。

2017-10-10   #16

对齐内存访问、合并内存访问。

当设备内存事物的第一个地址适用于事务服务的缓存粒度的偶数倍时(32字节的二级缓存或者128字节的一级缓存),就会出现对齐内存访问。运行非对齐的加载会造成带宽浪费。

当一个线程束中全部32个线程访问一个连续的内存块时,就会出现合并内存访问。

核函数的内存请求通常是在DRAM设备和片上内存间以128字节或者32字节内存事务来实现的。

如果两级缓存都启用,那么内存访问是由一个128字节的内存事务实现的,如果只使用了二级缓存,那么这个内存访问室友一个32字节的内存事务实现的。

P137页的图4-7 4-8例子非常棒

2017-10-10   #17

CPU一级缓存和GPU一级缓存之间的差异

CPU一级缓存优化了时间和空间局部性,GPU的则专为空间局部性,而不为时间局部性设计。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。

2017-10-10   #18

没有缓存的加载,在内存段的粒度(32字节)而非缓存池的粒度(128字节)执行,这是更细粒度的加载,可以为非对齐或非合并的内存访问带来更好的总线利用率。


是否使用缓存在P138-P140页的几个分类和示例中,非常明确地进行了说明!

2017-10-10   #19

数组结构体AoS

struct innerStruct{
    float x;
    float y;
}

struct innerStruct myAoS[N];

结构体数组SoA

struct innerArray{
    float x[N];
    float y[N];
}

SoA在SIMD模型上更好,因为x是连续的。

2017-10-11   #20

内部函数更快速,但是损失了部分精度。

2017-10-12   #21

CUDA上FP16计算方式的发展变化:
 


2017-10-13   #22

In the Pascal architecture, the FPU is capable of execut-

ing 2-way SIMD instruction of the half2 (FP16x2) data type. Figure 2 shows the example for adding two values in half2 using API from cuda_fp16.h provided starting from CUDA 7.5

1.png

2017-10-19   #23

Fermi hardware limit is 63 registers per thread. kepler is 256.

https://stackoverflow.com/questions/16975727/confusion-with-cuda-ptx-code-and-register-memory

2017-10-19   #24

--generate-line-info -Xptxas="-v" 几个重要的运行参数

登录后方可回帖

Loading...