[CUDA] cuda编程实践

cuda编程实战

1. cuda 理论介绍

1.1 cuda中的一些概念

• SP（streaming Processor）：是Nividia GPU硬件概念，可能对应的是线程束的执行单元
• SM（Streaming Multiprocessor）： 包含多个SP， 一个SM 会根据其内部SP数目分配warp，但是SM 不见得会一次就把这个warp 的所有指令都执行完；当遇到正在执行的warp 需要等待的时候（例如存取global memory 就会要等好一段时间），就切换到别的warp来继续做运算，借此避免为了等待而浪费时间。所以理论上效率最好的状况，就是在SM 中有够多的warp 可以切换，让在执行的时候，不会有「所有warp 都要等待」的情形发生；因为当所有的warp 都要等待时，就会变成SM 无事可做的状况了。
  • 一个SM中有warp scheduler且每个SM中可以同时执行多个warp线程束，一个warp有32个thread
  • 而在block 的方面，一个SM 可以处理多个线程块block，当其中有block 的所有thread 都处理完后，他就会再去找其他还没处理的block 来处理
• 但是一个block应该不会同时在多个SM上执行，因为要共享shared memory，且因为grid 包含多个block，如果grid过小，可能存在空闲SM的情况，导致硬件使用率过低
• 一个kenerl函数是怎么执行的：一个kernel程式会有一个grid，grid底下又有数个block，每个block是一个thread群组。在同一个block中thread可以通过共享内存（shared memory）来通信，同步。而不同block之间的thread是无法通信的
• CUDA的设备在实际执行过程中，会以block为单位。把一个个block分配给SM进行运算；而block中的thread又会以warp（线程束）为单位，对thread进行分组计算。目前CUDA的warp大小都是32，也就是说32个thread会被组成一个warp来一起执行。同一个warp中的thread执行的指令是相同的，只是处理的数据不同。
• 为了减少硬件资源浪费，最好设置block size是32的倍数

1.2 一些参考文档和关键知识点

• 【CUDA编程系列】CUDA编程基本入门学习笔记

• CUDA中grid、block、thread、warp与SM、SP的关系: 需要注意cuda编程中的一些概念的含义，如 block, grid, threads和硬件中SM，core的区别。

• 深入GPU硬件架构及运行机制: 清楚显存中各个cache的情况

• Gemm的优化-第五章: 清楚bank conflict的情况: 一个wrap中32个线程需要访问不同的bank才行，否则访问统一bank会存在多一次访存的操作，因为有多个thread访问一个bank的冲突。

• [CUDA] kernel归约注意事项

• [CUDA] stream使用笔记

• [CUDA] cuda kernel开发记录: kernel的一些属性说明，debug技巧和打印技巧，以及写法优化和注意事项。

• 判断一个指针是GPU还是CPU

• cudaHostRegitster的使用构建pinned memory

• 制定cuda sync的等待逻辑: 包括spin, yield， blocking 三种模式。

• [CUDA] gemm优化之wmma: wmma的使用

• [CUDA] ptx使用笔记： mma的使用，ldmatrix，reg, setp等一些指令的使用，如果需要可以查看官方文档去学习相关指令。

• [CUDA] cuda开发值memory使用: cuda中如何申请gpu空间时首地址对齐类型长度。

• [CUDA] 编译调试工具objdump使用查看某些cuda函数汇编代码: 如何查看cuda库的汇编语言逻辑，来推测耗时异常问题

reference

• [1] CUDA 编程 __launch_bounds__的应用方法
• [2]stream flags的设置 对比