Khi GPU đọc bộ nhớ: Hành trình từ lệnh SASS đến DRAM trên RTX 4090

Giới thiệu

Trong bài viết trước, chúng ta đã theo dõi một kernel cộng vector c[i] = a[i] + b[i] (mỗi thread xử lý một số float) từ mã nguồn CUDA qua biên dịch bởi nvcc cho đến các warp. Chúng ta đã đi sâu vào chi tiết cách kernel được khởi chạy, nhưng cũng bỏ qua rất nhiều khía cạnh khác.

Lần này, chúng ta sẽ giải quyết những thiếu sót đó, theo dõi con đường của một lệnh SASS quan trọng — một lệnh đọc bộ nhớ toàn cục (global load) — đi qua phần cứng. Trong trường hợp này, đó là một chiếc RTX 4090 đang nằm dưới bàn làm việc của tôi. Chúng tôi thực hiện kiểu phân tích ngược này vì lý do hiệu suất, ít nhất là về nguyên tắc (một lý do tuyệt vời được nêu trong bài báo vi mô chuẩn Citadel: "Why these details matter").

Hầu như không có chi tiết nào về con đường này được NVIDIA công bố, ít nhất là không đến mức chúng tôi mong muốn. Vì vậy, chúng tôi sẽ xác định nó bằng cách chạy các thí nghiệm định thời gian trên chính phần cứng.

Kernel CUDA mà chúng tôi khảo sát có hai dòng trong thân hàm:

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

Hành trình từ lệnh load đến cache L1

Trước tiên, lệnh đọc bộ nhớ toàn cục (global load) sẽ đi qua cache L1 của SM (Streaming Multiprocessor). Cache L1 trên RTX 4090 có kiến trúc 8-way set-associative, sử dụng địa chỉ ảo cho cả chỉ mục (index) và nhãn (tag). Điều này có thể được chứng minh bằng cách ánh xạ một vùng nhớ vật lý qua hai địa chỉ ảo khác nhau.

Cấu trúc chỉ mục (set index) của L1 là một hàm parity phức tạp, trong đó mỗi bit của chỉ mục là kết quả XOR của một tập hợp cố định các bit địa chỉ. Điều này không đơn giản như việc lấy một vài bit giữa của địa chỉ. Bằng cách dùng các "minimal conflict sets", chúng ta có thể tìm ra chính xác các bit nào tham gia vào hàm này.

Một slot trong cache có thể chứa một dòng (line) chỉ với một số sector (phân đoạn 32 byte) hiện diện. Khi một lệnh load bị miss, yêu cầu gửi xuống L2 chỉ nêu tên các sector mà warp cần.

Kiến trúc L2 và khái niệm "slice"

L2 cache là một thành phần phức tạp hơn nhiều. Không giống như L1, L2 được đánh địa chỉ và gắn nhãn vật lý (physically tagged and indexed). Điều này có nghĩa là quá trình dịch địa chỉ ảo sang vật lý phải xảy ra trước khi truy cập L2.

Một trong những chi tiết hấp dẫn nhất là chia sẻ L2 thành nhiều slice khác nhau. Trên AD102 (chip của RTX 4090), L2 được chia thành 12 controller bộ nhớ, mỗi controller phụ trách 3 slice (trên RTX 4090), trong khi trên L40S, mỗi controller có 4 slice. Mỗi địa chỉ vật lý được ánh xạ đến một slice cụ thể thông qua một hàm băm phức tạp.

Khám phá hàm băm slice

Hàm ánh xạ từ địa chỉ vật lý đến slice là một bài toán giải mã thú vị. Chúng tôi đã xác định được hàm này bằng thực nghiệm. Nó sử dụng các phép toán XOR trên các bit địa chỉ, kết hợp với một hàm parity và một phép toán modulo 3. Kết quả là một giá trị cho biết slice cụ thể.

Dưới đây là một phần của hàm (ở dạng Python) cho thấy độ phức tạp của nó:

def _state(a, N=36):
    b35 = (a >> 35) & 1
    P1 = parity(a & 0x0BD654C80)
    P2 = parity(a & 0x2882B0800)
    A = ((a >> 15) + 2*parity(a & 0x3C9041000) + parity(a & (0x2882B0800 ^ b35)) + 2) % 3
    g = ((a + (1 if b35 else 0)) >> 17) % 9
    q0 = parity(a & 0x8000)
    q1 = parity(a & 0x5985E0500)
    q2 = parity(a & (0x2354E4400 ^ b35))
    q3 = parity(a & 0x3C9041000)
    carry = 1 if q0 + q1 + q2 >= 2 else 0
    start = (5 + 7*q0 + 5*q1 + 2*q2 + q3 - carry) % 9
    o = (g - SHIFT[A] - start) % 9
    Lf = 2 if (q0 ^ q1 ^ q2) == 0 else 1
    return P1c, P1, P2, A, q2, o // 3, (1 if (o % 3) >= Lf else 0)

Ý tưởng cốt lõi là nếu bạn có một hàm hoạt động chính xác, việc kiểm tra rất dễ dàng: bằng cách tải dữ liệu từ chính xác k địa chỉ mà hàm dự đoán nằm trong cùng một slice, bạn sẽ thấy hiệu suất băng thông giảm đi đáng kể so với việc rải đều khắp 36 slice.

Kết quả từ bảng dưới đây cho thấy rõ ràng hiệu ứng của việc tập trung dữ liệu vào các slice dự đoán:

Số dòng được vẽ (L2-resident)	Tốc độ tải (Mload/s)	Tốc độ chuẩn hóa (vs k=1 predicted slice)
1,957	1.00x	1.00x
3,917	2.00x	2.00x
7,826	4.00x	4.00x
17,582	8.98x	8.98x
34,446	17.60x	17.60x
Tất cả 36	68,085	34.78x

Cấu trúc set index bên trong một slice

Sau khi xác định được hàm slice, bạn có thể tiếp tục phân tích cấu trúc cache bên trong một slice cụ thể. Kết quả cho thấy mỗi slice là 16-way set-associative. Chỉ mục set bên trong một slice cũng là một hàm parity phức tạp, tương tự như L1, nhưng khác nhau giữa các slice:

def set_index(a):
    q = a // 1152
    b0 = parity(a & 0x0bd654c80) ^ parity(q & 0x00e500)
    b1 = parity(a & 0x0bd654c80) ^ parity(q & 0x010000)
    # ... (các bit khác)
    b9 = parity(a & 0x13bc21180) ^ parity(q & 0x006400) ^ int((a >> 15) % 9 in (2, 6))
    return sum(b << i for i, b in enumerate([b0, b1, ..., b9]))

Truy cập DRAM: Hàng và cột

Cuối cùng, khi cache L2 bị miss, yêu cầu được chuyển đến bộ nhớ DRAM (GDDR6X trên RTX 4090). Bộ nhớ DRAM được tổ chức thành các hàng (row) và cột (column). Việc đóng/mở một hàng mất nhiều thời gian hơn là truy cập tuần tự trong cùng một hàng.

Bằng các thí nghiệm định thời gian, chúng tôi xác định rằng kích thước một hàng là 1 KiB (32 cột, mỗi cột 32 byte). Truy cập vào cùng một hàng chỉ tốn thêm khoảng 3.4 ns cho mỗi lần đọc, trong khi đọc một hàng mới tốn thời gian gấp khoảng 15 lần.

Kết luận

Hành trình mà một GPU đọc bộ nhớ là một quá trình nhiều tầng phức tạp, từ việc biên dịch SASS, đi qua các tầng cache với các hàm băm phức tạp, đến việc quản lý các hàng/cột trong DRAM. Những chi tiết như cách chia L2 thành các slice và hàm ánh xạ địa chỉ cụ thể không được NVIDIA công bố, nhưng có thể được phát hiện thông qua các kỹ thuật đo đạc timing chính xác và phân tích ngược.

Kiến thức này không chỉ thú vị về mặt kỹ thuật mà còn có ý nghĩa thực tiễn lớn trong việc tối ưu hiệu suất cho các ứng dụng GPU. Ví dụ, hiểu được cách dữ liệu được phân chia giữa các slice có thể giúp thiết kế các mẫu truy cập để giảm xung đột và tăng băng thông hiệu quả. Việc tránh các mẫu truy cập gây ra xung đột trên cùng một slice có thể cải thiện đáng kể thời gian thực thi các kernel nhạy cảm với băng thông.

Tham khảo

• Citadel Microbenchmarking Paper
• NVIDIA CUDA C++ Programming Guide
• Blog Doubleword