2025年2023年的深度学习入门指南(27) - CUDA的汇编语言PTX与SASS

科技前沿 • 2025-01-27 17:09 • 阅读 70

大家好，我是讯享网，很高兴认识大家。

通过前面的学习，我们了解了在深度学习和大模型中，GPU的广泛应用。可以说，不用说没有GPU，就算是没有大显存和足够先进的架构，也没法开发大模型。

有的同学表示GPU很神秘，不知道它是怎么工作的。其实，GPU的工作原理和CPU是一样的，都是通过指令来控制硬件的。只不过，GPU的指令集和CPU不一样。下面我们就走进GPU的内部，看看如何用汇编来写GPU的程序。

初识PTX与SASS

从上图我们可以看到，CPU的架构是复杂的几个核组合在一起。而GPU的架构是大量的简单的核组合在一起。因为GPU的每个单元架构都很简单，所以我们需要用CPU去控制GPU的每个单元，让它们协同工作。CPU上的控制代码，我们称为host代码，而GPU每个单元上运行的代码，我们称为device代码。

CUDA的汇编语言分为两种，一种叫做Parallel Thread Execution，简称PTX，另一种叫做Streaming Assembly，简称SASS。PTX是一种中间语言，可以在不同的GPU上运行，而SASS是一种特定的汇编语言，只能在特定的GPU上运行。

下面我们看几个简单的例子来找找体感。

__global__ void test(int& c){ 
    c= blockIdx.x; }

讯享网

编译成PTX代码：

讯享网.visible .entry test(int&)( .param .u64 test(int&)_param_0 ) { ld.param.u64 %rd1, [test(int&)_param_0]; cvta.to.global.u64 %rd2, %rd1; mov.u32 %r1, %ctaid.x; st.global.u32 [%rd2], %r1; ret; }

PTX中间代码使用ld指令从内存中加载数据，用st指令将数据写入内存。mov用于在寄存器之间传递数据。cvta用于作地址转换。

因为要编译成真正的汇编代码，所以生成代码就要跟硬件架构相关了。我们来看一下sm值和架构的关系：

sm50: Maxswell 麦克斯韦架构。比如sm52对应GTX 980.
sm60: Pascal 帕斯卡架构。比如sm61对应GTX 1080.
sm70: Volta 伏特架构。比如sm70对应V100.
sm75: Turing 图灵架构。比如sm75对应RTX 2080, T4
sm80: Ampere 安培架构。比如A100, RTX3080
sm90: Hopper 哈珀架构。比如H100, RTX4080

下面我们将其编译成sm50架构的SASS代码：

test(int&): MOV R1, c[0x0][0x20] MOV R2, c[0x0][0x140] S2R R0, SR_CTAID.X MOV R3, c[0x0][0x144] STG.E [R2], R0 NOP NOP EXIT

与PTX不同，麦克斯韦架构下读取内存没有用ld指令，而仍然是MOV指令。而读取特殊寄存器SR_CTAID有专门指令S2R。写全局内存有指令STG.

我们再看sm60架构汇编：

讯享网test(int&): MOV R1, c[0x0][0x20] MOV R2, c[0x0][0x140] S2R R0, SR_CTAID.X MOV R3, c[0x0][0x144] STG.E [R2], R0 NOP NOP EXIT

跟sm50的没有什么区别。

再看sm70架构汇编：

test(int&): MOV R1, c[0x0][0x28] @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ S2R R5, SR_CTAID.X MOV R2, c[0x0][0x160] MOV R3, c[0x0][0x164] STG.E.SYS [R2], R5 EXIT

继续看图灵架构的:

讯享网test(int&): MOV R1, c[0x0][0x28] S2R R0, SR_CTAID.X ULDC.64 UR4, c[0x0][0x160] STG.E.SYS [UR4], R0 EXIT

图灵架构增加了ULDC指令，它用来从常量内存中读取到通用寄存器中。

sm80架构sass:

test(int&): MOV R1, c[0x0][0x28] S2R R5, SR_CTAID.X MOV R2, c[0x0][0x160] ULDC.64 UR4, c[0x0][0x118] MOV R3, c[0x0][0x164] STG.E [R2.64], R5 EXIT

sm90架构sass:

讯享网test(int&): LDC R1, c[0x0][0x28] S2R R5, SR_CTAID.X LDC.64 R2, c[0x0][0x210] ULDC.64 UR4, c[0x0][0x208] STG.E desc[UR4][R2.64], R5 EXIT

sm80和90没有实质上的变化。

编译和反汇编工具

有了感性认识之后，我们就来让代码运行起来。然后再介绍如何用工具来查看PTX代码和进行sass反汇编。

我们先写一个可以运行起来的CUDA代码，流程如下：

2025年2023年的深度学习入门指南(27) - CUDA的汇编语言PTX与SASS

初识PTX与SASS

编译和反汇编工具

相关推荐