---
title: 编译入门那些事儿（9）：SIMT 编程及编译-官方技术文章-鲲鹏社区
description: 本文主要向大家介绍SIMT 的编程界面，通过对比普通的C语言标量编程及 SIMD（Single Instruction Multiple Data）向量编程，展示其易用性。
keywords: SIMT,编译入门那些,事儿,编程及编译,官方技术文章,鲲鹏社区,引言,编程及其易用性分析
url: https://www.hikunpeng.com/developer/techArticles/20231127-10?envFlag=1
section: (其他)
---

# 编译入门那些事儿（9）：SIMT 编程及编译-官方技术文章-鲲鹏社区

URL: https://www.hikunpeng.com/developer/techArticles/20231127-10?envFlag=1
描述: 本文主要向大家介绍SIMT 的编程界面，通过对比普通的C语言标量编程及 SIMD（Single Instruction Multiple Data）向量编程，展示其易用性。
关键词: SIMT,编译入门那些,事儿,编程及编译,官方技术文章,鲲鹏社区,引言,编程及其易用性分析

官方技术文章 [了解详情](https://www.hikunpeng.com/zh/developer/techArticles)

编译入门那些事儿（9）：SIMT 编程及编译

编译入门那些事儿（9）：SIMT 编程及编译

编译

发表于 2023/11/27

0

## 引言

随着人工智能领域的发展，GPGPU（General-Purpose Graphics Processing Unit）领域获得了前所未有的关注，凭借其高性能计算的特点，以及 SIMT（Single Instruction Multiple Threads）编程的易用性，促进了整个生态的蓬勃发展。

这里首先介绍 SIMT 的编程界面，通过对比普通的C语言标量编程及 SIMD（Single Instruction Multiple Data）向量编程，展示其易用性。再结合 GPGPU 的架构，分析其代码执行过程。最后介绍 SIMT 代码编译流程。

## SIMT 编程及其易用性分析

高性能计算，最常见的操作便是 Element-wise 操作，这里以N个整型数据进行 Element-wise 的加法计算为例。

普通的C语言编程，常见的写法如下：

```
void foo(int *x, int *y, int N) {
```

上面的写法执行效率低，需要循环执行N次。为了提高性能，这里展示常见的 SIMD 编程方案（以 Arm 的 Neon 为例）：

```
#include <arm_neon.h>
```

以及常见的 SIMT 编程方案（以 CUDA 为例）：

```
__global__ void foo(int *x, int *y) {
```

上述 SIMD 和 SIMT 的编程方案均能达到 Element-wise 的加法效果，具体过程见代码标注。从编程界面来看，SIMT 编程方案最符合一般编程人员的编程习惯，使用常规的C语言运算符即可满足，而 SIMD 编程需要使用很多内置的数据类型（int32x4_t）和函数（例如vld1q_s32、vaddq_s32 和 vst1q_s32），所以 SIMT 的编程易用性更好。

介绍完 SIMT 的编程，下面介绍下 SIMT 代码如何在 GPGPU 硬件执行。

## SIMT 代码执行及性能分析

为了更加清晰的介绍 SIMT 代码的编译和执行，这里以 NVIDIA 的 Maxwell_GeForce-GTX-750 芯片为例，如下图所示。

图1 Maxwell_GeForce-GTX-750

从图中，可以发现，每个 SM 核中有 Warp Scheduler、Instruction Fetch/Dispatch、RegisterFile、Core、Shared Memory等。还是以上个章节介绍的 SIMT 代码为例，最终编译的硬件指令序列如下：

```
MOV R1, c[0x0][0x20] ;
```

上述 SIMT 代码，总的 thread 数 N 为 256，对应 Warp 个数为 8（Warp size=32），总的指令数为 13，那么代码执行过程如下图所示：

图2 SIMT代码执行过程

所有 Warp 均由 Warp Scheduler 统一下发，每个 Warp 的指令顺序执行，可以发现，执行过程中，第 n 个 Warp 的指令 a 在 suspended 状态下，可以切换到第 m 个 Warp 下发 b 指令，等到第 n 个 Warp 的 a 指令执行完成后，又可以下发 a+1 指令。由于同时存在多个 Warp 下发执行，这样就可以进行指令 latency 的隐藏，提高计算并行度。

## SIMT 代码编译

介绍完 SIMT 代码的硬件执行逻辑，这里以 CUDA 为例介绍 SIMT 代码的编译流程。CUDA 的官方编译流程，主要包括 PTX 和 SASS 两个阶段，如下图所示：

图3 CUDA编译流程

其中 cicc 编译阶段主要实现 PTX 的编译，包括 PTX 指令选择、Generic 地址空间推导、指令调度等。ptxas 编译阶段主要实现硬件指令 SASS 的编译，包括寄存器分配、指令调度、divergence 处理、指令依赖关系处理（fix latency 和 variable latency 指令分开处理）等。

由于 PTX 指令集开源，LLVM 社区维护了 NVPTX 后端实现，可以完成 CUDA 到 PTX 的 CICC 编译。这边整理了下 NVPTX 的编译流程中的关键 pass，如下图所示。可以看到，该编译流程核心实现了 PTX 的指令选择，并没有实现寄存器分配、divergence 等复杂逻辑。

图4 LLVM NVPTX后端编译流程

## 结语

在高性能计算领域，相比 SIMD 的编程，SIMT 由于其编程的易用性，生态蓬勃发展。

本文以 CUDA 为例，简单介绍了下 SIMT 的硬件指令在硬件上的执行逻辑，通过 Warp 间的隐藏实现代码的并行执行。

由于 SIMT 是标量编程，所以编译流程遵循传统的标量编译流程方案。不过由于硬件本身是多线程的并行执行，特别需要处理的是 divergence 及指令依赖关系的逻辑。

## 参考

https://www.techpowerup.com/gpu-specs/docs/nvidia-gtx-750-ti.pdf [了解详情](https://www.techpowerup.com/gpu-specs/docs/nvidia-gtx-750-ti.pdf)

https://llvm.org/docs/NVPTXUsage.html [了解详情](https://llvm.org/docs/NVPTXUsage.html)

https://docs.nvidia.com/cuda/index.html [了解详情](https://docs.nvidia.com/cuda/index.html)

本页内容