CUDA 学习笔记

The_beard

已于 2024-12-04 15:03:28 修改

阅读量2.2k

点赞数 63

文章标签： c语言 c++ 深度学习

于 2024-11-29 11:12:42 首次发布

本文链接：https://blog.youkuaiyun.com/the_beard/article/details/144115718

版权

前言

本人目前研三在读，恰好也没什么事干，趁此空档期学习一下cuda

参考书目：《CUDA C编程权威指南》（Professinal CUDA C Programming）

参考blog：人工智能编程 | 谭升的博客人工智能编程 | 谭升的博客人工智能编程 | 谭升的博客

由于该大佬只写了前六章，且因cuda版本距原书出版时间已迭代了多个版本，像是nvprof等工具已基本被nsys、ncu所替代，因而想着可以边学边记点东西，本人记忆力堪忧，未来忘了还可以回过头来喵喵（可能就是看完的第二天...）

前六章内容可参考上面blog链接，本人仅记后面部分，默认读者的先验知识已囊括谭升大佬blog中所涉及的知识。以下所有内容仅代表个人观点，认为博主翻译错误或表述错误的可参考原书，并且理解错误的地方还望各cuda大佬指正，不胜感激！

第七章 Tuning Instruction-Level Primitives

一个程序的时间瓶颈通常有两方面来决定：I/O和计算。本章主要focus优化计算部分。

由于gpu采用SIMT（Single Instruction, Multiple Threads）指令，其吞吐量通常会高于其他处理器。但光是处理器速度快并不意味着你的程序是高效、高速的，如果没有达到预期的结果或者算法收敛，那一切的计算都是无用功。因而了解不同低级原语在性能、数值精度和线程安全性方面的优缺点非常重要。知道您的内核代码何时被编译成某种原语，允许您根据需求调整编译器的代码生成。

个人理解：有些代码为加快速度，底层使用的汇编或者机器指令会不尽相同，例如用一条指令实现相乘和相加两个操作，理想情况下这样做会比用两条指令分别做乘和加所花费的时间减少一半，但代价是计算结果的精度变差。因而理解这块内容，能够更好理解为什么有些情况下理论上某个算法应该收敛，但实际计算结果却一直没有收敛。参考以下这个例子。

double value = in1 * in2 + in3;

这行代码代表两个计算操作，in1与in2相乘，后加in3。这种arithmetic pattern称之为multiply-add（MAD）。放在古董级cpu架构中，可能会分两条机器指令进行执行，但由于该操作非常常见，因而在NVIDIA GPUs框架下，该操作可以通过一条机器指令执行完毕，所需要的时钟周期减少，但相应的结果的精度会变差。原文：

As a result, the number of cycles to execute the MAD operation is halved. This performance does not come for free. The results of a single MAD instruction are often less numerically accurate than with separate multiply and add instruc- tions.

7.1 INTRODUCING CUDA INSTRUCTIONS

该小节主要囊括三个对 CUDA 内核生成的指令有重大影响的主题：浮点运算（floating-point operations），内置和标准函数（intrinsic and standard functions），原子操作（atomic operations）。

浮点运算

浮点的概念这里不作解释，可自行查询IEEE-754标准（不了解该概念的童鞋可自行回顾or学习计算机组成原理课程）。

在了解IEEE-754标准的基础上，需要了解到两个点

1. 大家应该都知道浮点数能表达的其实是离散的、有限的，而非真正意义上连续、无限的。（即有些数字你用32bit or 64bit的浮点数是表示不了的）

举个例子：

float a = 3.1415927f;
float b = 3.1415928f; 
if (a == b) 
{
    printf("a is equal to b\n"); 
} else 
{
    printf("a does not equal b\n"); 
}

显然a和b是不相等的，但如果你运行的机子是服从IEEE-754标准的框架的话，结果会输出：

a is equal to b

究其原因就是这两个数都不能用单精度浮点数准确表达，在采用round-to-nearest方法近似前提下，这两个数字会近似到同一个数字，因而出现该结果也不足为奇。除了round-to-nearest，还有round-up， round-down，round-to-zero。（again计组内容）在这个例子中，如果你使用double类型，结果是不同的，可以自己验证下。

2. 浮点数所表达的数字的颗粒度不同。通俗的讲，就是两个相邻数字的一范数距离是不同的。

这里你可以举几个十进制下的科学计数法表达的数字理解下。或者可以用C语言自带的nextafterf测试下。

从上述这个例子中你应该意识到一个关键的问题，随着x增大，其损失的精度也越来越严重，与之带来的一个问题就是，你选择哪种round方法(e.g. round-to-nearest)会严重影响到输出的结果。

内置和标准函数

标准函数（standard functions）是指在host和device都能执行的函数，例如C语言标准数学库中的sqrt、exp等函数，单指令操作例如乘、加也同样属于标准函数。

内置函数（intrinsic functions）是指只能在device上执行的代码。这里为避免博主错误的理解而误导读者，便不作翻译，贴上原文：

In programming, a function being intrinsic, or built-in, implies that the compiler has special knowledge about its behavior, which enables more aggressive optimization and specialized instruction generation. This is true for CUDA intrinsic functions. In fact, many trigonometric functions are directly implemented in hardware on GPUs because they are used heavily in graphics applications (to perform translations, rotations, and so on for 3D visual applications).

个人理解：gpu为加速计算，底层硬件有其固有的指令集（参考RISC与CISC的区别），因而有些代码可以直接编译成gpu独有的机器指令。

在CUDA中，有许多内置函数的作用同标准函数相同，例如标准函数sqrt对应于内置函数__dsqrt_rn，以及单精度的__fdividef。内置函数比其等效的标准函数decompose成更少的指令，好处与坏处不再赘述。

之后的章节（Standard vs. Intrinsic Functions）中会更具体地介绍如何利用这两种函数来平衡求解速度和求解精度。

原子操作

原子性的概念可以参考操作系统课程中的定义，本书中给出的解释为：

An atomic instruction performs a mathematical operation, but does so in a single uninterruptable operation with no interference from other threads.

通俗地讲，在进行原子操作时，该操作并不会因为其他线程的干扰而中断。好吧，只是把上面后半句话翻译了一下，具体还是用个例子来解释。例如你想要对一个变量先进行加1的操作，从传统cpu架构来考虑，使用汇编需要三个步骤：1.从内存读取变量至寄存器 2.从寄存器中读取，进行加1操作，再存回寄存器 3.从寄存器中保存至内存。那么在执行过程中，类似于流水线技术，可能会产生读后写或者写后读的冲突，在你从寄存器读取变量之前，寄存器内的内容已经被另一线程修改了，这样就会导致结果的不确定。而操作的原子性保证这三个步骤是一个整体。（这里的解释仅为方便理解原子性，实际情况可能并非如此，要理解原子性，也可以考虑操作系统中的各种经典同步问题，例如读者写者问题。仅代表个人观点，如有错误，请轻喷...）

有了基本的概念后，我们再来看书中给出的代码例子：

__global__ void incr(int *ptr) 
{ 
    int temp = *ptr;
    temp = temp + 1;
    *ptr = temp;
}

假如我启动一线程束（32线程），写这个核函数目的是使ptr指向的变量增加32，但如果kernel写成这样的话，大概率结果并非如此，即如果ptr指向的变量初始值为0，结果大概率不是32，原因不再赘述。CUDA提供了加法的原子函数，将以上代码修改为如下即可得到预期结果：

__global__ void incr(__global__ int *ptr) 
{ 
    int temp = atomicAdd(ptr, 1);
}

基本的操作加减乘除（arithmetic functions）、与或异或（bitwise functions）、交换（swap functions）都有对应的原子函数，具体有哪些及各个参数含义都可以查官方文档的API。

再看一个书中的例子：

__global__ void check_threshold(int *arr, int threshold, int *flag) 
{ 
    if (arr[blockIdx.x * blockDim.x + threadIdx.x] > threshold) 
    {
        *flag = 1; 
    }
}

这个核函数的目的是，线程数超过threshold，那么就将flag全局变量修改为1。可以考虑到写操作其实是不安全的（if multiple values are above the threshold then the assignment to flag is unsafe）。那么使用原子指令就能保证写的安全性：


__global__ void check_threshold(int *arr, int threshold, int *flag) 
{ 
    if (arr[blockIdx.x * blockDim.x + threadIdx.x] > threshold) 
    {
        atomicExch(flag, 1); 
    }
}

当然在这个例子中，你使不使用原子指令，最后的结果都会是一致的。举这个例子主要是为了说明原子操作的代价是gpu计算性能的降低（using atomicExch and other atomic operations may significantly degrade performance）。同时值得注意的是，原子操作并不会像内置函数一样会降低结果的精度。

有关原子操作的内容会在之后的章节Understanding Atomic Instructions中具体再介绍。

小结

该小节概念介绍偏多，没有太多需要自己code的地方，如果对操作系统、计组比较熟悉的童鞋应该是比较好理解的，当然细节上或许有些不同，但博主理解大体的思路是差不多的（again个人观点...）第一次写blog可能存在诸多不足，还请各位看官多多体谅以及给予一些建议。因为原书是全英文的，所以有些翻译也可能存在偏差，还望海涵（blame on GPT）。

参考自：

人工智能编程 | 谭升的博客

Cheng, J., Grossman, M. and McKercher, T., 2014. Professional CUDA c programming. John Wiley & Sons.