# Ch2 Programming Model
> Lightweight thread creation, zero-overhead thread scheduling, and fast barrier synchronization efficiently support very fine-grained parallelism.
- The concurrent threads of a thread block express fine-grained data parallelism and **thread parallelism**.
- The independent thread blocks of a grid express coarse-grained **data parallelism**.
- Independent grids express coarse-grained **task parallelism**

## SIMD

What?

- SIMD 是一種 parallel 的架構，允許單一指令同時處理不同 data-point，非常適合於重複、數據密集型任務。
- 目前的 GPU 主要使用 SIMD (Single Instruction, Multiple Data) 的架構來實現 GPU 應用程式中內含的 data-level parallelism

不過在寫程式時，CUDA 和 OpenCL 等 GPU 運算 API 不是直接將 SIMD 硬體暴露給程式設計師，而是具有類似 MIMD 的程式設計模型，允許程式設計師在 GPU 上啟動大量 scalar thread。其中的每一個 scalar thread 都可以遵循其獨特的執行路徑，並且可以存取任意記憶體位置
::: info
**SIMD 架構和 SIMT 模型**
在運行時，GPU 硬體在 SIMD 架構上以 warp（或 AMD 中的 wavefront）為單位來執行運算。這種執行模型稱為單指令、多執行緒 (SIMT) [Lindholm et al., 2008a, Nickolls and Reusch, 1993]。
:::

### 補充：MIMD

What?

- GPU 計算用的 API（如 CUDA 和 OpenCL）提供了一種類似 MIMD（多指令多數據）的編程模型。MIMD 是另一種形式的平行計算，在這種計算方式中，**多個處理器可以對不同的數據執行不同的指令。**
- 這種方法利用了硬體執行平行操作的能力，同時提供了更簡單的編程介面。

How?

- 在 API 提供的類似 MIMD 模型中，工程師可以在 GPU 上啟動大量 scalar threads，抽象化了底層的 SIMD 硬體
- 每個 scalar threads 都能夠進行獨立操作，執行不同的指令並獨立地訪問任意 memory location。
- 在 runtime 時，GPU 硬體執行 groups of scalar threads（稱為 warp 或 wavefront），並在 SIMD 硬體上同步執行。

整理:
| Aspect | SIMD(Single Instruction Multiple Data): | MIMD(Multiple Instruction Multiple Data): |
|------------------------------|-----------------------------------------------------------------|------------------------------------------------------------------|
| **Instruction Complexity** | Low; single instruction stream for multiple data. | High; multiple instruction streams for multiple data. |
| **Scalability** | High for uniform tasks; limited for diverse operations. | Hard due to complex architecture. |
| **Core Speed** | Slower in GPU (in-order execution) | Faster than SIMD ( out-of-order execution, branch prediction) |
| **Flexibility** | Low; best for tasks that require the same operation on all data points. | High; capable of handling diverse tasks simultaneously. |
| **Task Suitability** | Ideal for data-parallel tasks such as graphics and simulations. | Suited for a variety of independent, parallel tasks. |
| **Design Optimization** | Can be highly optimized for parallel execution of the same operation. | Requires complex coordination and management logic. |
| **Application** | GPU | Multicore CPU |

# 2.1 EXECUTION MODEL

## CPU 在 GPU 計算中的角色

1. **discrete GPU**
- **記憶體分配**：CPU 會為「將在 GPU 上進行的計算」分配 memory。而 GPU 也要在其 memory 中預留空間來存儲 data 和結果。
- **數據傳輸**：將計算所需的輸入 data 從 CPU memory 轉移到 GPU memory。此步驟很重要，因為 GPU 需要 access 此 data 以執行計算。
- **啟動 Computational Kernel**：CPU 會啟動要在 GPU 上執行 Computational Kernel(`lauch kernel`)。
- Kernel 是在 GPU 上執行的一個函數或一組指令，由 GPU 執行。Kernel 通常被設計成可以平行運行於多個 GPU thread。
2. **Integrated GPU**
- 對於與 CPU 合併的 GPU（即與 CPU 在同一芯片上），步驟稍有不同。
- 由於 CPU 和 GPU 共享相同的 memory，因此無需進行獨立的 memory 分配和數據傳輸步驟。
- CPU 只需要直接在 GPU 上啟動 Computational Kernel。

**Computational Kernel and Threads**

- Computational Kernel 可能有成千個 threads 組成。每個 thread 運行相同的指令集，但可以處理不同的 data，並根據計算結果遵循不同的執行路徑。
- 這種平行執行模型是 GPU 在某些應用中高性能的關鍵。每個 thread 就像一個大團隊中的工人，每個工人都執行相同的任務，但對不同 data 進行處理。

**性能考慮**：

- 需要注意的是，特別針對一種架構（如 GPU）優化的代碼，在另一種架構（如 CPU）上可能表現不佳。這是因為 GPU 和 CPU 在其架構和處理模型上有根本區別。

## CUDA program (補充)

::: success
ch1 相關筆記內容 [連結](https://hackmd.io/N4Enq8iSTwWreOhASRld0w?view#GPU-%E5%95%9F%E5%8B%95%E9%81%8B%E7%AE%97%E7%9A%84%E7%A8%8B%E5%BA%8F)
:::
CUDA 提供三個關鍵的 abstractions
1. a hierarchy of thread groups,
2. shared memories,
3. barrier synchronization

CUDA program 存在 host code 和 device code

- 單純的 c program 就是 host code (CPU serial code)
- kernel 就是 device code
- When a kernel function is called, a large number of threads are **launched** on a device to execute the kernel.
- All the threads that are launched by a kernel call are collectively called a **grid**
- When all threads of a grid have completed their execution, the grid terminates, and the execution continues on the host until another grid is launched.
- 如圖，有兩個 grid 依序被呼叫執行
- 這些 threads 在 GPU 上，會在幾個 cycle 內被產生和排程；而 CPU 通常會需要上千個 cycle 來完成相同的事情
- In CUDA, the execution of each thread is sequential as well.

::: danger
圖中的 CPU 和 GPU 是經過簡化過的模型，兩者實際上應該是要 Overlap execution
:::


## SAXPY 算式

> single- precision scalar value A times vector value X plus vector value Y, known as SAXPY.

// 用 `__global__` 標示出 kernel，即會在 GPU 上執行的 function
__global__ void saxpy(int n, float a, float *x, float *y)
{
		// for loop 平行化的部分-----
	 int i = blockIdx.x*blockDim.x + threadIdx.x;  //Thread indexing
	  if(i<n)
		   y[i] = a*x[i] + y[i];
		// -------------------------
}
int main() {
	float *h_x, *h_y;
	int n;
	// omitted: allocate CPU memory for h_x and h_y and initialize contents
	float *d_x, *d_y;
	int nblocks = (n + 255) / 256;
	// cudaMalloc: 呼叫 GPU driver，讓他去在 GPU 上分配記憶體給 program
    // d_x 設定為指向 GPU 記憶體的一個區域
	cudaMalloc( &d_x, n * sizeof(float) );
	cudaMalloc( &d_y, n * sizeof(float) );
	//cudaMemcopy: 把 h_x 指向的CPU memory content，複製給 d_x 指向的 GPU memory
	cudaMemcpy( d_x, h_x, n * sizeof(float), cudaMemcpyHostToDevice );
	cudaMemcpy( d_y, h_y, n * sizeof(float), cudaMemcpyHostToDevice );

	//⬇️ kernel launch
    // CPU 啟動由 nblocks 個 thread blocks 組成的單一 grid
    // 其中每個 thread blocks 包含 256 個執行緒
	 saxpy<<<nblocks, 256>>>(n, 2.0, d_x, d_y);
	//⬇️ CPU 呼叫 GPU driver，把資料從 GPU memory 複製回 CPU memory
	 cudaMemcpy( h_x, d_x, n * sizeof(float), cudaMemcpyDeviceToHost );
	  // omitted: use h_y on CPU, free memory pointed to by h_x, h_y, d_x, and d_y 20

for loop 平行化的部分

//line 4:
y[i] = a* x[i] + y[i]

會轉換成

//calculate global thread ID index in 1D grid of 1D blocks
int i = blockIdx.x*blockDim.x + threadIdx.x;
if(i<n)
	y[i] = a*x[i] + y[i];

相當於，**每個 iteration 都會由一個 thread 來執行**

### 細節介紹

在範例中，考慮由程式設計師管理的獨立 GPU 和 CPU 記憶體(相對於 unified memory)

naming:
`h_` (host) 是 CPU 上的 Memory，`d_`(device) 是 GPU 上的 Memory
`cudaMalloc`: 呼叫 GPU driver，讓他去在 GPU 上分配記憶體給 program
`cudaMemcopy`: 把 h_x 指向的 CPU memory content，複製給 d_x 指向的 GPU memory

記憶體存取:

- thread block 內的 thread 透過 shared memory 進行溝通
- thread block 內的 thread 透過 barrier 進行同步。
- 每個 SM 上可以跑多個 thread block。
- 不同 thread block 間 thread 的溝通，需要經由 global memory。

## CUDA Basics (補充)

::: info
層次結構小到大為 thread -> warp -> thread block -> grid。
一個 thread block 包含多少 thread，是軟體概念，由內 kernel 配置語法決定；
而一個 warp 包含多少 thread，是硬體概念，warp 是硬體排程的最小粒度，無法改動預設。
:::


### 0. Thread Indexing

- threadIdx.[x y z]- thread 座標。 Thread(2,1) 的 threadIdx.x = 2, threadIdx.y = 1。
- blockIdx.[x y z]- block 座標。Block 的 blockIdx.x = 1, blockIdx.y = 1。
- blockDim.[x y z]- block 的維度，表示 block 包含多少個 thread。框架圖中，blockDim.x = 5, blockDim.y = 3。
- gridDim.[x y z] - grid 的維度，表示一個 grid 包含多少個 block。框架圖中，gridDim.x = 3, gridDim.y = 2。

::: warning
(x, y)是(橫坐標，縱座標)，不是(列，行)
:::

Example and diagram:

- [CUDA Thread Indexing. When I was learning CUDA programming, I… | by Anuradha Karunarathna | Medium](https://anuradha-15.medium.com/cuda-thread-indexing-fb9910cba084)

Explain detail:

- [初识 GPU 编程 | 鲁老师 (lulaoshi.info)](https://lulaoshi.info/gpu/python-cuda/cuda-intro.html)
- [Lecture 5.pdf (wfu.edu)](https://users.wfu.edu/choss/CUDA/docs/Lecture%205.pdf)
- [cuda 中 threadIdx、blockIdx、blockDim 和 gridDim 的使用 - 知乎 (zhihu.com)](https://zhuanlan.zhihu.com/p/544864997)


平行程式設計中，會為每個執行緒分配一部分資料。為了實作這種策略，GPU 上的每個執行緒都可以在 grid 的執行緒區塊中找到自己的 id。在 CUDA 中執行此操作的機制使用 `gridDim` 、`blockDim`、`blockIDx`、`threadIDx`。

block (`blockIDx`) 和 thread (`threadIDx`)都是用 `dim3` 的內建變數來進行 index

`threadIDx`

- uint3 threadIdx -- thread index within block:
- threadIdx.x, threadIdx.y, threadIdx.z

`blockIDx`

- uint3 blockIdx -- block index within grid:
- blockIdx.x, blockIdx.y, blockIdx.z

Full global thread ID in x and y dimensions can be computed by:

- x = blockIdx.x \* blockDim.x + threadIdx.x;
- y = blockIdx.y \* blockDim.y + threadIdx.y;

`blockDim` 定義，thread block 大小

- 一個 Block 中 Thread 的個數，Number of threads in a block = blockDim.x _ blockDim.y _ blockDim.z
- Thread 在 block 中擺放的形狀
- 一個 Block 中的 `threadIdx` 最大不超過 `blockDim`；
- dim3 blockDim -- Size of block dimensions x, y, z
- x：1024
- y：1024
- z：64
- **一個 thread block 最多只能有 1024 個 threads!!!**

`gridDim` 定義，grid 大小

- 一個 Grid 中 Block 的個數，Number of blocks in grid = gridDim.x \* gridDim.y
- Block 在 Grid 中擺放的形狀
- 一個 Grid 中的 `blockIdx` 最大不超過 `gridDim`。
- dim3 gridDim -- Grid dimensions, x, y, z
- x： 2^31 − 1
- y：65535
- z：65535

**CUDA 變數解析:**

**_Number of blocks in the grid:_**

gridDim.x — number of blocks in the x dimension of the grid *(eg: 3)*

gridDim.y — number of blocks in the y dimension of the grid *(eg: 2)*

gridDim.z — number of blocks in the z dimension of the grid *(eg: 1)*

**_Number of threads in a block:_**

blockDim.x — number of threads in the x dimension of the block *(eg: 4)*

blockDim.y — number of threads in the y dimension of the block *(eg: 3)*

blockDim.z — number of threads in the z dimension of the block *(eg: 3)*

**_Block Index:_**

blockIdx.x — block’s index in x dimension

blockIdx.y — block’s index in y dimension

blockIdx.z — block’s index in z dimension

_eg: block (0,1) — blockIdx.x = 0 , blockIdx.y = 1_

**_Thread Index:_**

ThreadIdx.x — thread’s index in x dimension

ThreadIdx.y — thread’s index in y dimension

ThreadIdx.z — thread’s index in z dimension

_eg: Thread(2,1) — ThreadIdx.x = 2, ThreadIdx.y = 1_

**`dim3` 作為 integer vector type 來當參數**

For example, the following host code can be used to call the `vecAddkernel()` kernel function and generate a 1D grid that consists of 32 blocks, each of which consists of 128 threads. The total number of threads in the grid is 128*325=4096:
The first execution configuration parameter specifies the dimensions of the grid in number of blocks.

The second specifies the dimensions of each block in number of threads

Each such parameter has the type `dim3`, which is an **integer vector type** of three elements x, y, and z.
- These three elements specify the sizes of the three dimensions.
- The programmer can use fewer than three dimensions by setting the size of the unused dimensions to 1.

---

### 1. Thread:

- 定位: The **smallest unit of execution in CUDA**.
- 作用: Executes a single instance of a kernel (a function that runs on the GPU).
- Threads can be identified using unique IDs, accessible through built-in variables like `threadIdx`.
- 特性: Has its own set of registers and local memory.
- CUDA 中，每個 individual thread 其實是只用了 individual value (相對於 vector)
- 例如， 一個 thread take a single 32-bit floating-point number, perform some calculations on it, and then output another 32-bit floating-point number.

而為了增加效率，讓多個 threads 可以平行地執行運算

- NVDA 使用 Warp (32 個 threads) 為每次 SIMD 的執行單位
- 多個 Warps => A Block
- AMD 使用 Wavefront (64 個 threads) 為每次 SIMD 的執行單位
- 多個 Wavefronts => A cooperative thread array (CTA)

### 2. Block:

不同 Grid 之間的 block 是獨立的，可能在相同或不同 SM 上；
但同一個 block 的 thread 一定會在同一個 SM

- 定位: A block is a group of threads that execute together on the **same** Streaming Multiprocessor (SM).
- cannot be migrated to other SMs in GPU (except during preemption, debugging, or **CUDA dynamic parallelism**).
- **One SM can run several concurrent CUDA blocks** depending on the resources needed by CUDA blocks.
- All the thread blocks operate on different parts of the vectors. They can be executed in any arbitrary order. The programmer must not make any assumptions regarding execution order.
- 特性: Threads within a block can communicate and synchronize with each other using shared memory and synchronization primitives like `__syncthreads()`.
- `__syncthreads` 讓你在同一個 block 中，進行 thread 的同步 
- **ID 識別** => A block is identified using `blockIdx` and has a maximum size (number of threads), which is determined by the architecture and the resources used by the kernel.
- **獨立執行** => Blocks are independent; once launched, they execute independently from other blocks.
- **thread 上限** => - CUDA architecture limits the numbers of threads per block (1024 threads per block limit). - 最好是 以 32 為倍數的分配 threads 給每個 block，因為實際上是以 warp 為單位執行運算

### 3. Warp:

- 定位:
- The set of parallel threads that execute the same instruction together in a SIMT architecture.

- GPU 內，SM 排程和執行指令的單位，32 個 threads 為一個 block，threads 之間 **"execute in lockstep"**

- This **lockstep** execution ensures that **at any instruction, all threads in the warp are synchronized with each other**,
- 執行方式本身，就達到 (Implicit) Inter-Warp Synchronization
- 而不需要呼叫 explicit synchronization commands like `__syncthreads()`

- Warp as a SIMD Unit
- 同一個 warp 中的 threads，會在相同時間，對不同 data 執行相同的指令(SIMD)
- Warps 之間，會共享控制單元(warp scheduler)
- The GPU **schedules and executes instructions for warps,** not individual threads.

- 特性:

- **在同一個 warp 中的 threads，才是 Physically in parallel (物理上真的是平行)**
- 所以 Warp 才不能太大，因為同一個 warp 裡面的 thread 會 synchronize，
- 在不同 warp 中的 threads，是 logically in parallel (邏輯上平行，物理上不是)
- 在很大程度上，編譯器和硬體的結合使程式設計師可以忽略執行緒執行的鎖步性質。編譯器和硬體讓 warp 中的每個執行緒看起來像是獨立執行

- 作用:
- 每個 CUDA core in an SM TAKE one thread from a warp at a time.
- 屬於「Concurrently」執行，同一時間只能做一個 thread，但可以透過 context switch 快速切換
- **Inter-Warp Synchronization**: 在相同 block 但 不同 warp 之間的同步 通常是必須的，會用到前面提過的 `__syncthreads()`. especially when dealing with shared memory or coordinating work across the block.
- 利用 Warp scheduling 達到 Latency hiding:
- While a warp is waiting for data (e.g., memory access), the **SM** **can switch to another warp** without incurring significant context-switching overhead. This is known as **latency hiding** and is key to the high performance of GPUs.

### 4. Grid:

- 定位:
- kernel launch 時，會對應到一組 grid configuration，代表此次要「分配多少 block」,「每個 block 多少 thread」
- 換句話說，一個 grid 對應到一個 kernel & kernel launch ( Each grid is dedicated to the execution of a single instance of a kernel.)
- A grid 由一組 blocks 組成， execute the same kernel.
- 同一個 Grid: 可能在不同 SM 處理，不同 Grid: 執行不同 kernel
- 特性:
- CUDA 變數：A grid is identified using `gridDim`, 代表 total number of blocks that are executing a given kernel.
- The size and dimensions of a Grid (i.e., the number and layout of Blocks and Threads) are determined at the time of the CUDA kernel call.
- 作用:
- 主要任務是，讓 blocks in a grid 可以用彈性的方式(自定義 block 和 thread 的數量)在 multiple SMs 上執行，allowing for parallel execution of blocks.
- Fig2.9: Each thread block is organized as a one-dimensional array of threads because the data are one-dimensional vectors.

### 5. Kernel: A Grid of thread blocks

- 定位:
- CUDA kernel 是在 GPU 上執行的 function
- 開頭都是用 `__global__` 來宣告
- 在支援 CUDA dynamic parallelism 的系統中, it can also be called from the device
- 特性: A kernel is executed as a grid of blocks of threads
- Each **kernel** is executed on **one device**
- CUDA supports **running multiple kernels on a device at one time**.

### 6. NVIDIA (CUDA) 與 AMD (OpenCL) 術語對照表

| NVIDIA (CUDA) | AMD (OpenCL) |
| --------------------------------------------- | ------------------------- |
| Grid | NDRange |
| Thread Block (Co-operative thread array, CTA) | Work Group |
| Warp (32 個 threads) | Wavefront (64 個 threads) |
| Thread | Work Item |

### 7. CUDA C keywords for function declaration

- `__device__`: can be called only from a kernel function or another device function
- `__host__`: By default, all functions in a CUDA program are host functions if they do not have any of the CUDA keywords in their declaration
- `__global__`: 宣告 kernel function

### 軟硬體對照

- 一個 thread 一定對應一個 CUDA Core，而 CUDA Core 可能對應多個 thread。
- 一個 Block 內的 thread 一定會在同一個 SM（Streaming Multiprocessor）內
- 一個 SM 可以運行多個 Block
- 每個 block 內的 thread 會以 warp 為單位進行運算，一個 warp 對應一條指令流，一個 warp 內的 thread 是真正同步的，同一個 warp 內的 thread 可以讀取其他 warp 的值。

## Unified Memory
> A system architecture in which the CPU and GPU share a common system memory.

傳統上，GPU 有自己的 DRAM，在用於機器學習的資料中心內的 GPU 也是如此。
而 NVIDIA 引入了統一記憶體 (Unified Memory)，它透明化地從 CPU 記憶體更新 GPU 記憶體，以及從 GPU 記憶體更新 CPU 記憶體 （Unified Memory is a single memory address space that is accessible from any processor in a system.）

- Uses only CPU system memory, omitting GPU memory from the system.
- 讓資料可以從 CPU 或 GPU 直接進行 R/W
- replace calls to `malloc()` or `new` with calls to `cudaMallocManaged()`, an allocation function that returns a pointer accessible from any processor
- `cudaError_t cudaMallocManaged(void** ptr, size_t size);`

These systems have relatively low-performance GPUs, since their achieved performance is limited by the available system memory bandwidth and increased latency of memory access, whereas dedicated GPU memory provides high bandwidth and low latency.

# 2-2 GPU 指令集架構

> NVIDIA fully documents this virtual instruction set architecture with each release of CUDA to the point that it was easy for the authors of this book to develop the GPGPU-Sim simulator to support PTX [Bakhoda et al., 2009]

本節討論 kernel 從 CUDA 和 OpenCL 等高階語言到 GPU 硬體執行的彙編層級的轉換，以及目前 GPU 指令集的形式。 GPU 架構與 CPU 架構有些不同的一個有趣方面是 GPU 生態系統已經進化以支援指令集進化的方式。例如，x86 微處理器向後相容於 1976 年發布的 Intel 8086。向後相容意味著為上一代架構編譯的程式將在下一代架構上運行而無需任何更改。因此，40 年前為 Intel 8086 編譯的軟體理論上可以在當今的任何 x86 處理器上運行。

## NVCC (NVIDIA C compiler)

- 使用 CUDA keywords to separate the host code and device code.
- Host code 就按照正常 c program 去編譯
- Device code 會經由 NVCC 來編譯成 virtual binary files called **PTX (Parallel Thread Execution)** files. These PTX files are further compiled by a runtime component of NVCC into the real object files and executed on a CUDA-capable GPU device.


## 2.2.1 NVIDIA GPU 指令集架構

### PTX

一直以來，大量供應商提供 GPU 硬體（每個都有自己的硬體設計），透過 OpenGL 著色語言(OGSL) 和微軟的高級著色語言(HLSL) ，一定程度的指令集虛擬化隨著隨著早期的 GPU 變得 programmable 而變得很常見。
當 NVIDIA 在 2007 年初推出 CUDA 時，他們決定走類似的道路，並為 GPU 運算引入了自己的 high-level 虛擬指令集架構，稱為 Parallel Thread Execution ISA，或 PTX [NVI，2017] 。 NVIDIA 在每個 CUDA 版本中都完整地記錄了這種虛擬指令集架構，以至於本書的作者很容易開發出支援 PTX 的 GPGPU-Sim 模擬器 [Bakhoda et al., 2009]。

- PTX 在許多方面類似於標準精簡指令集電腦 (RISC) 指令集架構，如 ARM、MIPS、SPARC 或 ALPHA。它也與最佳化編譯器中使用的中間表示具有相似性。一個這樣的例子是使用無限組的虛擬暫存器。

### SASS(Streaming ASSembler)

在 GPU 上執行 PTX 程式碼之前，需要將 PTX 編譯為硬體支援的實際指令集架構，NVIDIA 將此層級稱為 SASS，它是 "Streaming ASSembler" 的縮寫 [Cabral，2016]。

從 PTX 轉換為 SASS 的過程，可以透過 GPU 驅動程式或 NVIDIA 的 CUDA 工具包提供的 ptxas 來完成。

NVIDIA 沒有完整記錄 SASS。雖然這使得學術研究人員更難開發能夠捕獲所有編譯器優化效果的架構模擬器，但它使 NVIDIA 從客戶需求中解放出來，在硬體層級提供向後相容性，從而能夠從一代到下一代完全重新設計指令集架構。

因此，希望了解底層效能的開發人員開始創建自己的工具來 disassemble（反彙編） SASS。由 Wladimir Jasper van der Laan 完成並命名為 "decuda" [van der Lann] 的第一個作品於 2007 年底推出，用於 NVIDIA 的 GeForce 8 系列(G80)，當時是在第一個支持 CUDA 的硬體發布後的大約一年內。 decuda 專案對 SASS 指令集有了足夠詳細的了解，因此可以開發彙編程式。這有助於在 GPGPU-Sim 3.2.2 [Tor M. Aamodt 等人] 中開發對 SASS 的支持，直至 NVIDIA 的 GT200 架構。 NVIDIA 最終推出了一個名為 cuobjdump 的工具，並開始紀錄一部份的 SASS。
NVIDIA 的 SASS 文件 [NVIDIA Corporation, c] 目前（2018 年 4 月）僅提供了 a list of the assembly opcode names，但沒有提供有關 operand format 或 SASS 指令語義的詳細資訊。最近，隨著在機器學習中使用 GPU 的爆炸性增長以及對效能最佳化程式碼的需求，其他人已經為後續架構開發了類似於 decuda 的工具，例如 NVIDIA 的 Fermi [Yunqing] 和 NVIDIA 的 Maxwell 架構 [Gray ]。


圖 2.4 展示了我們為 NVIDIA 的 Fermi 架構 [NVI, 2009] 編譯並使用 NVIDIA 的 cuobjdump（CUDA 工具包的一部分）提取的 SAXPY 核心的 SASS 程式碼。

- 圖 2.4 中的第一列是指令的位址。第二列是 disassembly(machine code -> assembly code)，第三列是編碼指令。
- 比較圖 2.3 和圖 2.4，可以看出虛擬和硬體 ISA 等級之間的相似之處和差異。在 high-level 上存在著重要的相似之處，例如都是 RISC（都使用 **load** 和 **store** 來存取記憶體）和都使用 **predicate** [Allen et al., 1983]。
- 更細微的差異包括：
- PTX 版本具有基本上無限的可用 registers sets，因此每個定義通常使用一個新暫存器，很像靜態單一分配 [Cytron 等人，1991]，而 SASS 使用有限的一組暫存器；
- kernel 參數透過 "儲存在 SASS 中的 non-load/store 指令" 可以存取的 banked constant memory 進行傳遞，而在 PTX 中，參數會被分配到它們自己單獨的「parameter」address space。


圖 2.5 展示了 SAXPY 的 SASS 程式碼，該程式碼由相同版本的 CUDA 但針對 NVIDIA 的 Pascal 架構產生並使用 NVIDIA 的 cuobjdump 擷取。

比較圖 2.5 和圖 2.4，很明顯 NVIDIA 的 ISA 發生了顯著變化，包括指令編碼方面。

圖 2.5 包含一些沒有 disassembly 的行（例如，在第 3 行的位址 0x0000 處）。這些是在 NVIDIA Kepler 架構中引入的特殊“控制指令 (control instructions)”，以消除使用 score board 機制來進行顯式的 dependency check 的需要 [NVIDIA Corporation, b]。

Lai 和 Seznec [2013] 探討了 Kepler 架構的 control instructions encoding。正如 Lai 和 Seznec [2013] 所指出的，這些 control instructions 似乎類似於 Tera 電腦系統上的 explicit-dependence lookahead [Alverson et al., 1990]。

Gray 描述了他們能夠為 NVIDIA 的 Maxwell 架構推斷出的 control instructions encoding 的大量細節。根據 Gray 的說法：

- Maxwell 中每三個常規指令就有一個控制指令。這似乎也適用於 NVIDIA 的 Pascal 架構
- Maxwell 上的 64 位控制指令包含三組 21 位，為以下三個指令中的每一個編碼接下來的信息：
- a stall count；yield hint flag；以及 write, read, and wait dependency barriers 。
- 暫存器重用標誌(register reuse flags) 在常規指令上的使用，如圖 2.5 所示（例如，R0.reuse 用於第 8 行的整數短乘加指令(Integer Short Multiply Add instruction) XMAD 中的第一個 source operand）。

這似乎表明從 Maxwell 開始在 NVIDIA GPU 中添加了「operand reuse cache」（請參閱第 3.6.1 節中的相關研究）。這種 operand reuse cache 似乎能夠讓每個 main register file 讀取多次 register value，從而降低能耗和/或提高效能。

## 2.2.2 AMD GRAPHICS CORE NEXT ISA

略

# Reference

- Aamodt, T. M., Fung, W. W. L., & Rogers, T. G. (2018). General-Purpose Graphics Processor Architectures. Morgan & Claypool Publishers.
- Hwu, W. W., Kirk, D. B., El Hajj, I. (2023). Programming Massively Parallel Processors: A Hands-on Approach. Morgan Kaufmann.
- [CUDA Refresher: The CUDA Programming Model | NVIDIA Technical Blog](https://developer.nvidia.com/blog/cuda-refresher-cuda-programming-model/)
- [【NV-02】CUDA 并行计算--线程和索引 - 知乎](https://zhuanlan.zhihu.com/p/583258642)