Chương 7: Thiết kế Bộ xử lý đa luồng Streaming trong thiết kế GPU

Bộ xử lý đa luồng Streaming (SM) là khối xây dựng cơ bản của kiến trúc GPU của NVIDIA. Mỗi SM chứa một tập các lõi CUDA thực hiện các chỉ thị theo kiểu SIMT (Single Instruction, Multiple Thread). SM chịu trách nhiệm quản lý và lập lịch các warp, xử lý sự phân kỳ của nhánh, và cung cấp truy cập nhanh chóng vào bộ nhớ chia sẻ và bộ đệm. Trong chương này, chúng ta sẽ khám phá vi kiến trúc của SM, bao gồm các ống dẫn, cơ chế lập lịch warp, thiết kế tệp thanh ghi và tổ chức bộ nhớ chia sẻ và bộ đệm L1.

Vi kiến trúc và Ống dẫn của SM

SM là một bộ xử lý được thiết kế rất song song và có đường ống nhằm thực thi hiệu quả hàng trăm luồng đồng thời. Hình 7.1 hiển thị một sơ đồ khối đơn giản của một SM trong kiến trúc Volta của NVIDIA.

                                 Bộ nhớ Cache Chỉ thị
                                         |
                                         v
                                    Lập lịch Warp
                                         |
                                         v
                               Đơn vị Phân phối (4 warp)
                                 |   |   |   |
                                 v   v   v   v
                               Lõi CUDA (FP64/FP32/INT)
                               Lõi CUDA (FP64/FP32/INT)
                               Lõi CUDA (FP64/FP32/INT)
                               ...
                               Tensor Core
                               Tensor Core
                               ...
                               Đơn vị Tải/Lưu
                               Đơn vị Tải/Lưu
                               ...
                               Đơn vị Hàm đặc biệt
                                         ^
                                         |
                                Tệp Thanh ghi (64 KB)
                                         ^
```Bộ nhớ dùng chung / Cache L1 (96 KB)

Hình 7.1: Sơ đồ khối đơn giản của một SM trong kiến trúc NVIDIA Volta.

Các thành phần chính của SM bao gồm:

Bộ nhớ cache lệnh: Lưu trữ các lệnh được truy cập thường xuyên để giảm độ trễ và cải thiện hiệu suất.
Bộ lập lịch warp: Chọn các warp sẵn sàng để thực thi và gửi chúng đến các đơn vị thực thi sẵn có.
Đơn vị phân phối: Lấy và giải mã các lệnh cho tối đa 4 warp mỗi chu kỳ và gửi chúng đến các đơn vị thực thi thích hợp.
Lõi CUDA: Các đơn vị thực thi có thể lập trình hỗ trợ một loạt các phép toán số nguyên và số dấu phẩy động. Mỗi SM trong Volta chứa 64 lõi CUDA.
Lõi Tensor: Các đơn vị thực thi chuyên dụng được thiết kế để tăng tốc học sâu và các tải công việc AI. Mỗi SM trong Volta chứa 8 Lõi Tensor.
Đơn vị tải/lưu trữ: Xử lý các thao tác bộ nhớ, bao gồm cả tải và lưu trữ vào bộ nhớ toàn cục, bộ nhớ dùng chung và bộ nhớ cache.
Đơn vị chức năng đặc biệt: Thực hiện các phép tính super-hình học và các phép tính toán phức tạp khác.
Tệp thanh ghi: Cung cấp quyền truy cập nhanh vào các thanh ghi riêng tư của thread. Mỗi SM trong Volta có một tệp thanh ghi 64 KB.
Bộ nhớ dùng chung / Cache L1: Một không gian bộ nhớ có thể cấu hình được sử dụng làm bộ nhớ cache do phần mềm quản lý (bộ nhớ dùng chung) hoặc làm bộ nhớ cache L1 do phần cứng quản lý.

Đường ống SM được thiết kế để tối đa hóa hiệu suất bằng cách cho phép nhiều warp thực thi đồng thời và ẩn độ trễ bộ nhớ. Hình 7.2 minh họa một cái nhìn đơn giản về đường ống SM.

    Lấy lệnh
            |
            v
    Giải mã lệnh
            |
            v
    Nhận toán hạng
            |
            v
    Thực thi (Lõi CUDA, Lõi Tensor, Đơn vị tải/lưu trữ, Đơn vị chức năng đặc biệt)
            |
            v
    Ghi lại

Hình 7.2: Đường ống SM đơn giản.

Các giai đoạn của đường ống như sau:

Lấy lệnh: Bộ lập lịch warp chọn một warp sẵn sàng để thực thiDưới đây là bản dịch tiếng Việt của file Markdown này. Đối với phần code, tôi không dịch code, mà chỉ dịch các comment:
Lấy Lệnh: Lấy lệnh tiếp theo cho warp đó từ bộ nhớ cache lệnh.
Giải Mã Lệnh: Lệnh được lấy được giải mã để xác định loại hoạt động, toán hạng và các thanh ghi đích.
Thu Thập Toán Hạng: Các toán hạng cần thiết cho lệnh được thu thập từ tệp thanh ghi hoặc bộ nhớ chung.
Thực Thi: Lệnh được thực thi trên đơn vị thực thi thích hợp (Lõi CUDA, Lõi Tensor, Đơn vị Tải/Lưu hoặc Đơn vị Chức Năng Đặc Biệt).
Ghi Lại: Kết quả của quá trình thực thi được ghi lại vào tệp thanh ghi hoặc bộ nhớ chung.

Để đạt được hiệu suất cao, SM sử dụng nhiều kỹ thuật để tối đa hóa việc sử dụng tài nguyên và ẩn đi độ trễ:

Phát Hành Song Song: SM có thể phát hành hai lệnh độc lập mỗi warp trong một chu kỳ, cho phép tăng tính song song cấp lệnh.
Đơn Vị Thực Thi Được Đường Ống: Các đơn vị thực thi được đường ống, cho phép SM bắt đầu một hoạt động mới trên một đơn vị trước khi hoạt động trước đó hoàn thành.
Ẩn Độ Trễ: SM có thể chuyển đổi giữa các warp theo chu kỳ, cho phép nó ẩn độ trễ của các truy cập bộ nhớ và các hoạt động có độ trễ dài bằng cách thực thi các lệnh từ các warp khác.

Ví dụ 7.1 cho thấy một hạt nhân CUDA đơn giản thực hiện phép cộng phần tử-theo-phần tử của hai vector.

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

Ví dụ 7.1: Hạt nhân CUDA cho việc cộng vector.

Trong ví dụ này, mỗi luồng trong hạt nhân tính tổng của các phần tử tương ứng từ các vector đầu vào a và b và lưu kết quả vào vector đầu ra c. SM thực hiện hạt nhân này bằng cách gán mỗi luồng cho một Lõi CUDA và lập lịch các warp luồng để thực hiện trên các lõi có sẵn. Các đơn vị tải/lưu được sử dụng để lấy dữ liệu đầu vào từ bộ nhớ toàn cục và ghi kết quả trở lại.

Lập Lịch Warp và Xử Lý Phân Kỳ

Ví...Dưới đây là bản dịch tiếng Việt của file markdown, với phần code giữ nguyên bản tiếng Anh và chỉ dịch các phần bình luận:

Lập lịch warp hiệu quả là rất quan trọng để tối đa hóa hiệu suất của SM. Bộ lập lịch warp chịu trách nhiệm chọn các warp sẵn sàng để thực thi và phân phối chúng đến các đơn vị thực thi có sẵn. Mục tiêu chính của bộ lập lịch warp là giữ cho các đơn vị thực thi luôn bận rộn bằng cách đảm bảo rằng luôn có các warp sẵn sàng để thực thi.

SM sử dụng cơ chế lập lịch warp hai cấp:

Lập lịch warp: Bộ lập lịch warp chọn các warp sẵn sàng để thực thi dựa trên một chính sách lập lịch, chẳng hạn như vòng tròn hoặc cũ nhất trước. Các warp được chọn sau đó được phân phối đến các đơn vị thực thi có sẵn.
Lập lịch chỉ thị: Trong mỗi warp, SM lập lịch các chỉ thị dựa trên sự phụ thuộc của chúng và khả năng của các đơn vị thực thi. SM có thể cấp phát nhiều chỉ thị độc lập từ cùng một warp trong một chu kỳ để tối đa hóa song song cấp độ chỉ thị.

Hình 7.3 minh họa cơ chế lập lịch warp hai cấp.

    Warp Pool
    Warp 1 (Sẵn sàng)
    Warp 2 (Đang chờ)
    Warp 3 (Sẵn sàng)
    ...
    Warp N (Sẵn sàng)
        |
        v
    Bộ lập lịch warp
        |
        v
    Đơn vị phân phối
        |
        v
    Các đơn vị thực thi

Hình 7.3: Cơ chế lập lịch warp hai cấp.

Một trong những thách thức chính trong lập lịch warp là xử lý sự phân kỳ nhánh. Trong mô hình thực thi SIMT, tất cả các luồng trong một warp thực hiện cùng một chỉ thị theo thời gian thực. Tuy nhiên, khi một warp gặp một chỉ thị phân nhánh (ví dụ, một câu lệnh if-else), một số luồng có thể đi theo nhánh if trong khi những luồng khác đi theo nhánh else. Tình huống này được gọi là sự phân kỳ nhánh.

Để xử lý sự phân kỳ nhánh, SM sử dụng một kỹ thuật được gọi là gán điều kiện (predication). Khi một warp gặp một nhánh phân kỳ, SM thực hiện cả hai đường dẫn của nhánh tuần tự, che đi các luồng không đi theo mỗi đường dẫn. Kết quả sau đó được kết hợp bằng các thanh ghi điều kiện để đảm bảo rằng mỗi luồng nhận được kết quả chính xác.

Ví dụ 7.2 cho thấy một kernel CUDA với sự phân kỳ nhánhĐây là bản dịch Tiếng Việt của tệp Markdown này. Đối với phần mã, tôi không dịch mã, chỉ dịch các nhận xét.

__global__ void divergentKernel(int *data, int *result) {
    // Lấy ID của thread
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    // Nếu data[tid] > 0
    if (data[tid] > 0) {
        // Gán result[tid] = data[tid] * 2
        result[tid] = data[tid] * 2;
    } else {
        // Gán result[tid] = data[tid] * 3
        result[tid] = data[tid] * 3;
    }
}

Ví dụ 7.2: Hạt nhân CUDA với một nhánh phân kỳ.

Trong ví dụ này, điều kiện nhánh data[tid] > 0 có thể khiến một số thread trong một warp thực hiện đường dẫn if, trong khi những thread khác thực hiện đường dẫn else. SM xử lý sự phân kỳ này bằng cách thực hiện cả hai đường dẫn tuần tự và gán mặt nạ cho các thread không hoạt động trong mỗi đường dẫn.

Hình 7.4 minh họa quá trình gán mặt nạ cho một warp với các thread phân kỳ.

    Warp (32 thread)
    Thread 1: data[1] = 5, result[1] = 10
    Thread 2: data[2] = -3, result[2] = -9
    ...
    Thread 32: data[32] = 7, result[32] = 14

    Nhánh phân kỳ:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    Gán mặt nạ:
    Bước 1: Thực hiện đường dẫn if với mặt nạ
        Thread 1: result[1] = 10
        Thread 2: (bị gán mặt nạ)
        ...
        Thread 32: result[32] = 14

    Bước 2: Thực hiện đường dẫn else với mặt nạ
        Thread 1: (bị gán mặt nạ)
        Thread 2: result[2] = -9
        ...
        Thread 32: (bị gán mặt nạ)

    Kết quả cuối cùng:
    Thread 1: result[1] = 10
    Thread 2: result[2] = -9
    ...
    Thread 32: result[32] = 14

Hình 7.4: Quá trình gán mặt nạ cho một warp với các thread phân kỳ.

Bằng cách sử dụng gán mặt nạ, SM có thể xử lý sự phân kỳ nhánh mà không cần các hướng dẫn nhánh rõ ràng hoặc sự phân kỳ luồng điều khiển. Tuy nhiên, các nhánh phân kỳ vẫn có thể ảnh hưởng đến hiệu suất, vì SM phải thực hiện cả hai đường dẫn tuần tự, làm giảm hiệu suất song song.

Tệp thanh ghi và Trình thu thập toán hạng

Tệp thanh ghi là một thành phần quan trọng của SM, cung cấp truy cập nhanh chóng đến các thanh ghi riêng của thread. Mỗi SM có một tệp thanh ghi lớn để hỗ trợ nhiều thread hoạt động và cho phép chuyển đổi ngữ cảnh hiệu quả giữa các warp.Dưới đây là bản dịch tiếng Việt của tệp Markdown trên:

Trong kiến trúc NVIDIA Volta, mỗi SM có một tệp thanh ghi 64 KB, được tổ chức thành 32 ngân hàng, mỗi ngân hàng 2 KB. Tệp thanh ghi được thiết kế để cung cấp băng thông cao và truy cập có độ trễ thấp để hỗ trợ số lượng lớn các luồng song song.

Để giảm thiểu xung đột ngân hàng và cải thiện hiệu suất, SM sử dụng một kỹ thuật được gọi là thu thập toán hạng. Bộ thu thập toán hạng là những đơn vị chuyên dụng thu thập toán hạng từ các ngân hàng tệp thanh ghi và chuyển chúng đến các đơn vị thực thi. Bằng cách sử dụng bộ thu thập toán hạng, SM có thể giảm tác động của xung đột ngân hàng và cải thiện mức độ sử dụng của các đơn vị thực thi.

Hình 7.5 hiển thị sơ đồ đơn giản của tệp thanh ghi và bộ thu thập toán hạng trong một SM.

    Tệp thanh ghi (64 KB)
    Ngân hàng 1 (2 KB)
    Ngân hàng 2 (2 KB)
    ...
    Ngân hàng 32 (2 KB)
        |
        v
    Bộ thu thập toán hạng
        |
        v
    Các đơn vị thực thi

Hình 7.5: Tệp thanh ghi và bộ thu thập toán hạng trong một SM.

Các bộ thu thập toán hạng hoạt động bằng cách thu thập toán hạng từ nhiều hướng dẫn và nhiều warp, cho phép SM phát hành hướng dẫn từ các warp khác nhau đến các đơn vị thực thi trong một chu kỳ duy nhất. Điều này giúp che giấu độ trễ của việc truy cập tệp thanh ghi và cải thiện tổng thể hiệu suất của SM.

Ví dụ 7.3 hiển thị một kernel CUDA thực hiện tích vô hướng của hai vector.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // Mỗi luồng tính toán một phần tổng của tích vô hướng
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partialSum[tid] = 0;
 
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    __syncthreads();
 
    // Tổng hợp các phần tổng riêng lẻ
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    // Luồng 0 lưu kết quả tích vô hướng
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

Trong ví dụ này, mỗi luồng tính toán một phần tổng của tích vô hướng bằng cách sử dụng chỉ số được gán cho nóDưới đây là bản dịch tiếng Việt của tệp Markdown:

Các phần tử từ các vector đầu vào. Các tổng phần được lưu trữ trong mảng bộ nhớ được chia sẻ partialSum. Sau khi tất cả các luồng đã tính toán các tổng phần của chúng, một giảm thiểu song song được thực hiện để cộng các tổng phần và thu được kết quả tích vô hướng cuối cùng.

Bộ sưu tập toán hạng đóng vai trò quan trọng trong ví dụ này bằng cách hiệu quả thu thập các toán hạng cho các truy cập bộ nhớ được chia sẻ và các phép tính số học. Nó giúp tránh xung đột ngân hàng và cải thiện việc sử dụng các đơn vị thực thi.

Kết luận

Bộ xử lý đa lõi dòng chảy là đơn vị tính toán cốt lõi trong các kiến trúc GPU hiện đại. Thiết kế của nó tập trung vào việc tối đa hóa thông lượng và che dấu độ trễ bộ nhớ thông qua sự kết hợp của đa luồng mịn, thực thi SIMT và thu thập toán hạng hiệu quả.

Các thành phần chính của SM bao gồm bộ lập lịch warp, chọn warp để thực thi; ngăn xếp SIMT, xử lý sự phân nhánh và hội tụ; bộ nhớ đăng ký và bộ sưu tập toán hạng, cung cấp truy cập nhanh chóng vào các đăng ký riêng tư của luồng; và bộ nhớ được chia sẻ và bộ nhớ đệm L1, cho phép chia sẻ và tái sử dụng dữ liệu với độ trễ thấp.

Khi các kiến trúc GPU tiếp tục phát triển, nghiên cứu trong các lĩnh vực như xử lý phân nhánh, lập lịch warp và thiết kế bộ nhớ đăng ký sẽ rất quan trọng để cải thiện hiệu suất và hiệu quả của các GPU trong tương lai. Các kỹ thuật mới như hình thành warp động, nén khối luồng và bộ nhớ đệm tái sử dụng toán hạng có tiềm năng để cải thiện đáng kể khả năng của SM và cho phép các mức độ hiệu suất mới trong các tải công việc tính toán song song.

Chapter 6 Gpu Performance Metrics and Analysis Chapter 8 Interconnect and on Chip Networks