Làm chủ CUDA và Máy tính hiệu năng cao, Phần VII

Nơi Phần VI để lại cho chúng ta

Phần VI kết thúc bằng một câu đáng được bóc mẽ:

Hướng dẫn cp.async không đặt bảng điểm dài.

Tệp thanh ghi không liên quan nên không có bit thanh ghi nào được đánh dấu đang chờ xử lý.

SMSP phát hành cp.async, công cụ sao chép nhận nó và SMSP ngay lập tức có quyền đưa ra hướng dẫn tiếp theo cho sợi dọc đó.

Đây không phải là một lưu ý tối ưu hóa nhỏ.

Đó là mô tả về một mô hình thực thi khác về cơ bản: mô hình yêu cầu bạn từ bỏ mô hình tinh thần về “các vấn đề về hướng dẫn, kết quả đến, hướng dẫn tiếp theo được tiến hành” và thay thế nó bằng một thứ gì đó giống như một quy trình sản xuất trong nhà máy:

các giai đoạn chồng chéo , bộ đệm tồn tại giữa chúng và thông lượng được xác định theo giai đoạn chậm nhất chứ không phải tổng của tất cả độ trễ của giai đoạn.

Trước khi có thể làm cho cp.async hoạt động hữu ích, chúng tôi cần một mô hình chính xác về những gì nó đang ẩn: hệ thống phân cấp bộ nhớ.

Phân cấp bộ nhớ của A100

A100 SXM4 có sáu cấp độ bộ nhớ quan trọng đối với các lập trình viên kernel. Chúng không được ghi chép đầy đủ và các con số trong tài liệu tiếp thị thường không phải là con số trong mã sản xuất.

Đăng ký

Mỗi SM trên Ampere có tệp đăng ký 256 KB, được chia sẻ trên bốn SMSP: 64 KB mỗi SMSP, với cổng đọc 256 bit trên mỗi chu kỳ.

Độ trễ truy cập tệp đăng ký thực tế là 0 chu kỳ trong trường hợp bỏ qua; đối với các lần đọc không được bỏ qua, chi phí sẽ được tính vào đường dẫn FMA 4 chu kỳ. Đăng ký không phải là nguồn trễ. Chúng là nguồn dung lượng và băng thông.

Giới hạn dung lượng là vấn đề quan trọng: mỗi luồng có thể sử dụng tối đa 255 thanh ghi.

Áp suất trên mức này khiến trình biên dịch tràn các giá trị vào bộ nhớ cục bộ; vùng riêng tư trên mỗi luồng được ánh xạ tới L1/L2/DRAM.

Không thể phân biệt được sự cố tràn với bất kỳ quyền truy cập bộ nhớ chung nào khác ở cấp độ phần cứng: chúng đi qua đơn vị MIO, đặt bảng điểm dài và chờ hơn 400 chu kỳ cho DRAM. Mỗi thanh ghi bị tràn tốn hai thao tác MIO.

Bộ nhớ dùng chung / Bộ nhớ đệm L1

Ampere trên mỗi SM L1 là một nhóm 192 KB được phân vùng giữa bộ nhớ dùng chung và bộ đệm dữ liệu L1 phần cứng.

Phần phân chia có thể định cấu hình (0/192, 32/160, 64/128, 100/92, 132/60, 160/32 (shared/cache, tính bằng KB)) qua cudaFuncSetAttribution với cudaFuncAttributionPreferredSharedMemoryCarveout.

Bộ nhớ dùng chung có 32 dãy, mỗi dãy rộng 4 byte.

Chỉ mục ngân hàng cho địa chỉ byte:

ngân hàng = (địa chỉ >> 2) & 31

Mẫu truy cập trong đó nhiều luồng trong một sợi dọc truy cập vào các địa chỉ khác nhau trong cùng một ngân hàng được tuần tự hóa.

Một xung đột giữa 4 ngân hàng gây ra 4× độ trễ của trường hợp không có xung đột. Độ trễ không xung đột là khoảng 23 chu kỳ; xung đột 4 ngân hàng kéo dài điều này lên ~ 35 chu kỳ; xung đột 8 ngân hàng đến ~ 51 chu kỳ. Hình phạt có tỷ lệ tuyến tính.

ngoại lệ phát sóng: nếu tất cả các luồng trong một sợi dọc truy cập vào địa chỉ chính xác trong một ngân hàng thì phần cứng sẽ phục vụ điều này dưới dạng một lần đọc và phát kết quả.

Ba mươi hai luồng truy cập vào ba mươi hai địa chỉ khác nhau mà tất cả đều ánh xạ tới cùng một ngân hàng không phải là một quảng bá. Đây là sự tuần tự hóa 32 chiều.

Bộ đệm L2

A100 có bộ đệm L2 40 MB, được chia thành hai phần 20 MB. Độ trễ lần truy cập L2: khoảng 180–200 chu kỳ, cao hơn hầu hết các tài liệu ngụ ý.

Truy cập vào lát cục bộ là ~160–180 chu kỳ; truy cập vào lát cắt từ xa (yêu cầu truyền tải thanh ngang) là ~ 200–230 chu kỳ.

Băng thông L2 tổng cộng khoảng 4 TB/s. Tỷ lệ băng thông L2 so với băng thông HBM là khoảng 15:1. Việc lắp một bộ làm việc vào L2 về mặt chất lượng khác với việc đổ nó vào HBM.

HBM2e

A100 SXM4 có sáu ngăn xếp HBM2e cung cấp băng thông lý thuyết cao nhất là 2 TB/s. Trong thực tế: một hạt nhân có mẫu truy cập đều đặn đủ để bão hòa tất cả các kênh đạt được 1,6–1,9 TB/s.

Các kiểu truy cập không thường xuyên có xung đột vùng đệm hàng: 800 GB/s–1,2 TB/s. Độ chi tiết byte ngẫu nhiên đọc: hàng chục GB/s, do lãng phí dòng bộ đệm.

Độ trễ HBM2e, được đo bằng bỏ qua L1 và L2: khoảng 450–600 chu kỳ ở 1410 MHz. Bộ đệm hàng chạm đất trong khoảng 300–350 chu kỳ; trượt khoảng 550–650 chu kỳ.

Hậu quả ở 1410 MHz: 500 chu kỳ × 0,71 ns/chu kỳ ≈ 355 nano giây dừng trên mỗi sợi dọc. Trong cửa sổ đó, 500 khe phát lệnh trên SM sẽ tối đi.

Nếu mọi sợi dọc thường trú đã cấp tải HBM và đang chờ, thì bạn có một gian hàng 500 chu kỳ không có sợi dọc đủ điều kiện để giải cứu bạn.

Đây là bức tường ký ức ở dạng bê tông. Giải pháp không phải là bộ nhớ nhanh hơn: mà là cơ cấu lại hoạt động di chuyển dữ liệu để độ trễ HBM chồng chéo với quá trình tính toán.

Hướng dẫn cp.async

cp.async đã được giới thiệu trong Ampere (sm_80). Nó thực hiện chuyển trực tiếp giống như DMA từ bộ nhớ chung sang bộ nhớ dùng chung, bỏ qua hoàn toàn tệp đăng ký:

cp.async.ca.shared.global [dst], [src], kích thước;
cp.async.cg.shared.global [dst], [src], kích thước;   // bỏ qua L1

Tham số size là 4, 8 hoặc 16 byte. Biến thể 16 byte là quan trọng nhất: nó phát ra LDG.128 được vector hóa, đạt được mức sử dụng giao diện bộ nhớ tối đa.

“Bỏ qua tệp đăng ký” thực sự có nghĩa là gì

Đường dẫn tải thông thường:

 LDG.128 R4, [R2] ; → bộ bảng điểm dài cho R4,R5,R6,R7
                         ; → treo dọc trên bất kỳ lần đọc nào của R4-R7
                         ; → 450-600 chu kỳ sau, HBM trả về dữ liệu
STS.128 [smem_ptr], R4 ; lưu trữ các thanh ghi → bộ nhớ dùng chung

Điều này yêu cầu 4 đăng ký khi chuyển tiếp. Tải đặt bốn bit bảng điểm dài. Warp không đủ điều kiện để đọc bất kỳ lệnh nào R4–R7 cho đến khi giao dịch HBM hoàn tất.

Đường dẫn cp.async:

CP.ASYNC.CA.SHARED.GLOBAL [smem_dst], [R2], 0x10
; → không có bit bảng điểm nào được đặt (không có thanh ghi đích)
; → warp ngay lập tức đủ điều kiện để đưa ra hướng dẫn tiếp theo
; → dữ liệu đến bộ nhớ dùng chung một cách không đồng bộ

Một chuyên dụng Công cụ sao chép không đồng bộ Ampe nhận yêu cầu thông qua thiết bị MIO, sở hữu giao dịch và thực hiện tải HBM và ghi bộ nhớ dùng chung độc lập với SMSP. Thiết bị MIO được giải phóng ngay sau khi chuyển giao.

Cơ chế cam kết/chờ

Cam kết (CP.ASYNC.COMMIT_GROUP): đánh dấu tất cả các hướng dẫn cp.async trước đó là nhóm cam kết. Chỉ sổ sách kế toán, không chờ đợi gì cả.

Chờ (CP.ASYNC.WAIT_GROUP N): tạm dừng cho đến khi có tối đa N nhóm cam kết vẫn đang chờ xử lý. N=0 là đồng bộ hóa hoàn toàn.

N=1 cho phép một nhóm trên chuyến bay vẫn nổi bật trong khi bạn tính toán nhóm trước đó.

 ống tự động = cuda::make_pipeline();

cho (int i = 0; i < BATCH_SIZE; i++)
    cuda::memcpy_async(smem[0][i], &gmem[base + i], sizeof(float4), pipe);
pipe.producer_commit();

cho (int i = 0; i < BATCH_SIZE; i++)
    cuda::memcpy_async(smem[1][i], &gmem[base + BATCH_SIZE + i], sizeof(float4), pipe);
pipe.producer_commit();

pipe.consumer_wait();   // CP.ASYNC.WAIT_GROUP 1
__syncthreads();        // bắt buộc: truyền bá khả năng hiển thị tới tất cả chủ đề

calc(smem[0]);

__syncthreads() sau consumer_wait là bắt buộc. consumer_wait đảm bảo dữ liệu nằm trong bộ nhớ dùng chung theo quan điểm của sự cố này.

Các lỗi khác trong khối luồng có thể không nhìn thấy nội dung ghi cho đến khi __syncthreads() truyền chúng qua Miền kết hợp của SM.

Bỏ qua nó là điều kiện chạy đua: điều kiện tạo ra kết quả chính xác trong hầu hết thời gian và kết quả không chính xác không thể đoán trước dưới áp lực bộ nhớ nặng nề.

Mẫu bộ đệm đôi

Vòng lặp GEMM xếp kề tiêu chuẩn Vòng lặp GEMM hoàn toàn tuần tự: tải ô, đồng bộ hóa, tính toán, đồng bộ hóa, lặp lại. Dòng thời gian là một chuỗi phụ thuộc phẳng. Đối với các vấn đề nhỏ hơn hoặc các ô mỏng hơn có T_load / T_compute > 1, hạt nhân bị giới hạn bộ nhớ.

Mẫu bộ đệm đôi phá vỡ chuỗi đó:

 Iter k: |-- cp.async A[k] --|-- cp.async B[k] --|-- commit --|
                                                                |-- chờ(k-1) --|-- tính toán(k-1) --|
Iter k+1: |-- cp.async A[k+1] --|-- cp.async B[k+1] --|-- commit --|
                                                                    |-- chờ(k) --|-- tính toán(k) --|

Tải cho lần lặp k+1 trùng lặp với việc tính toán lần lặp k. Độ trễ bộ nhớ được ẩn miễn là T_load(k+1) < T_compute(k). Sau đó, quy trình sẽ chạy ở tốc độ tính toán mà không bị dừng bộ nhớ.

Điều này yêu cầu hai bộ đệm bóng bàn trong bộ nhớ dùng chung, tăng gấp đôi yêu cầu về bộ nhớ dùng chung.

Nhân đôi bộ nhớ dùng chung trên mỗi khối luồng sẽ giảm một nửa số khối luồng thường trú tối đa trên mỗi SM, giảm tỷ lệ chiếm chỗ. Sự đánh đổi là rõ ràng và có thể tính toán được.

Tín hiệu chẩn đoán: nếu smsp__warp_issue_stalled_long_scoreboard.avg.pct_of_peak_sustained_active vượt quá 20%, độ trễ bộ nhớ sẽ không bị ẩn. Sự can thiệp đầu tiên là công suất phòng cao hơn.

Điều thứ hai, khi công suất sử dụng đã gần đạt mức tối đa, là cp.async đường dẫn, loại bỏ hoàn toàn bảng điểm dài khỏi phương trình.

Mẫu hạt nhân đầy đủ

 constexpr int TILE_M = 128, TILE_N = 128, TILE_K = 32;
constexpr int NUM_STAGES = 2;

__global__ void gemm_async_kernel(
    const __nv_bfloat16* __restrict__ A,
    const __nv_bfloat16* __restrict__ B,
    nổi* __hạn chế__ C,
    int M, int N, int K
) {
    __shared__ __nv_bfloat16 smem_A[NUM_STAGES][TILE_M][TILE_K];
    __shared__ __nv_bfloat16 smem_B[NUM_STAGES][TILE_K][TILE_N];
    float acc[4][4] = {};

    ống tự động = cuda::make_pipeline();
    const int k_tiles = K/TILE_K;

    // MỞ ĐẦU: tạo ô 0 trước vòng lặp chính
    nếu (k_tiles > 0) {
        int row_a = threadIdx.x / TILE_K, col_a = threadIdx.x % TILE_K;
        nếu (hàng_a < TILE_M)
            cuda::memcpy_async(&smem_A[0][row_a][col_a],
                               &A[(blockIdx.y * TILE_M + row_a) * K + col_a],
                               sizeof(__nv_bfloat16), ống); pipe.producer_commit();
    }

    // VÒNG LẶP CHÍNH
    for (int k = 1; k < k_tiles; k++) {
        const int sw = k % 2, sr = (k - 1) % 2;

        int row_a = threadIdx.x / TILE_K, col_a = threadIdx.x % TILE_K;
        nếu (hàng_a < TILE_M)
            cuda::memcpy_async(&smem_A[sw][row_a][col_a],
                               &A[(blockIdx.y * TILE_M + row_a) * K + (k * TILE_K + col_a)], sizeof(__nv_bfloat16), ống);
        pipe.producer_commit();

        pipe.consumer_wait();   // CP.ASYNC.WAIT_GROUP 1
        __syncthreads();

        cho (int ki = 0; ki < TILE_K; ki++)
            vì (int i = 0; i < 4; i++)
                vì (int j = 0; j < 4; j++)
                    acc[i][j] += __bfloat162float(smem_A[sr][threadIdx.y*4+i][ki])
                               * __bfloat162float(smem_B[sr][ki][threadIdx.x*4+j]); __syncthreads();
    }

    // PHẦN MỞ
    pipe.consumer_wait();   // CP.ASYNC.WAIT_GROUP 0
    __syncthreads();

Ba điều cần hiểu về cấu trúc này:

Phần mở đầu không phải là tùy chọn. Nếu không cấp ô 0 trước vòng lặp, consumer_wait đầu tiên sẽ chặn một nhóm cam kết không tồn tại. Hành vi không xác định. Đoạn mở đầu thiết lập bất biến “trước một giai đoạn” mà vòng lặp phụ thuộc vào.

Cả hai nguyên tắc đồng bộ hóa đều được yêu cầu. consumer_wait đảm bảo công cụ DMA đã ghi dữ liệu vào bộ nhớ dùng chung cho sợi dọc này. __syncthreads() đảm bảo tất cả các luồng trong khối đã đạt đến điểm này trước khi bất kỳ luồng nào được đọc.

Họ giải quyết các vấn đề khác nhau. Không thay thế cho cái kia.

Đọc giai đoạn và ghi giai đoạn không bao giờ bằng nhau. Số học mô-đun đảm bảo sw ≠ sr cho NUM_STAGES = 2. Công cụ DMA ghi vào một bộ đệm trong khi các luồng đọc từ bộ đệm khác.

Với NUM_STAGES ≥ 3, bạn đào sâu quy trình, ẩn độ trễ nhiều hơn, tiêu tốn nhiều bộ nhớ dùng chung hơn.

Tổng quát hóa giai đoạn N

Với N giai đoạn, bạn phát hành cp.async có giá trị N ô trước khi lần tính toán đầu tiên bắt đầu. Độ trễ bị ẩn khi T_compute(tile) > T_HBM_load / N.

CUTLASS triển khai tối đa quy trình 5 giai đoạn cho hạt nhân Ampere GEMM, với số lượng giai đoạn dưới dạng tham số mẫu tại thời điểm biên dịch được trình phân tích tài nguyên quét tại thời điểm điều chỉnh. Chi phí bộ nhớ chia sẻ tỷ lệ tuyến tính với số lượng giai đoạn.

Tại một số điểm giao nhau, yêu cầu bộ nhớ dùng chung buộc phải giảm tỷ lệ sử dụng vượt quá lợi ích của đường ống.

Sự kết hợp này phụ thuộc vào hạt nhân cụ thể và kích thước vấn đề, đó là lý do tại sao CUTLASS hiển thị tham số thay vì mã hóa cứng nó.

Trình hồ sơ hiển thị cho bạn những gì

Trước khi lắp đặt đường ống (tải LDG thông thường):

smsp__warp_issue_stalled_long_scoreboard — 40–70%, chiếm ưu thế
smsp__pipe_fma_cycles_active — 30–60%, thiếu tính toán

Sau khi tạo đường dẫn (cp.async, bộ đệm đôi):

smsp__warp_issue_stalled_long_scoreboard. <5%, cp.async không đặt bit bảng điểm
smsp__pipe_fma_cycles_active. 70–90% cho hạt nhân được điều chỉnh tốt
Theo dõi smsp__warp_issue_stalled_mio_throttle; nếu bạn phát hành cp.async nhanh hơn mức đơn vị MIO có thể phục vụ chúng (~1 trên 4 chu kỳ cho mỗi SMSP đối với truyền 128 bit), thì gian hàng này sẽ thay thế gian hàng của bảng điểm.

Cách khắc phục là các ô lớn hơn hoặc chấp nhận điều tiết nếu thông lượng MIO vẫn vượt quá thông lượng điện toán.

Xung đột ngân hàng

Mô hình 32 ngân hàng đã được ghi lại. Ý nghĩa thực tế của các mẫu truy cập ma trận là không.

Trong GEMM xếp chồng, ô A được tải vào bộ nhớ dùng chung theo bố cục hàng chính, sau đó đọc theo cột trong quá trình nhân.

Đối với TILE_K = 32 và phần tử BF16 (mỗi phần tử 2 byte), phần tử [j][i] nằm ở độ lệch byte j × 64 + i × 2. Chỉ số ngân hàng: (j × 16 + i/2) & 31.

Đối với cột đọc cong i (i đã sửa, j chạy 0..31), mọi cặp luồng đều ánh xạ tới cùng một ngân hàng. Đây là xung đột ngân hàng 2 chiều trên mỗi cột được đọc.

Cách khắc phục là phần đệm:

__shared__ __nv_bfloat16 smem_A[TILE_M][TILE_K + 2];   // +2 BF16 = +4 byte mỗi hàng

Với phần đệm, phần tử [j][i] nằm ở độ lệch byte j × 68 + i × 2. Chỉ số ngân hàng: (j × 17 + i/2) & 31. Vì gcd(17, 32) = 1, các chỉ số ngân hàng dưới dạng j chạy 0..31 tạo thành một hoán vị hoàn chỉnh là 0..31. Không có xung đột.

Chi phí bộ nhớ dùng chung là TILE_M × 4 byte cho mỗi bộ đệm: 512 byte cho TILE_M = 128, không đáng kể so với ô 8 KB.

Kỹ thuật Swizzle của CUTLASS đạt được kết quả tương tự thông qua hoán vị bit địa chỉ thay vì đệm tuyến tính, xử lý rõ ràng các kích thước ô không có lũy thừa hai.

Phép tính bên dưới giống hệt nhau.

Chính sách bộ đệm L1

Hành vi của bộ nhớ đệm trên Ampere có thể định cấu hình ở cấp độ hướng dẫn:

Hành vi vòng loại LDG.CA Bộ nhớ đệm trong L1 (mặc định) LDG.CG Bỏ qua L1, đi tới L2 LDG.CS Truyền phát: chèn tại vị trí LRU LDG.CV Bỏ qua tất cả bộ đệm (hầu như không bao giờ chính xác)

Trong CUDA: __ldg() dành cho bộ nhớ đệm L1, __ldcg() / __ldcs() dành cho các biến thể bỏ qua. Trình biên dịch mặc định là LDG.CA khi không chắc chắn.

Đối với các hạt nhân xử lý từng phần tử đầu vào chính xác một lần, các thao tác theo từng phần tử, rút gọn, bất kỳ thao tác nào không cần sử dụng lại, __ldcg() sẽ loại bỏ ô nhiễm L1 và duy trì dung lượng L1 cho dữ liệu được hưởng lợi từ bộ nhớ đệm.

Hiệu ứng trong trình phân tích hồ sơ: tỷ lệ trúng L1 thấp hơn, tỷ lệ trúng L2 không thay đổi. Dữ liệu bỏ qua một cấp độ bộ đệm mà không làm giảm băng thông hiệu quả ở cấp độ thực sự tồn tại việc sử dụng lại.

Mô hình mái nhà

Mô hình đường mái (Williams, Waterman, Patterson, 2009) biểu thị FLOP/s theo cường độ số học (FLOP/byte lưu lượng DRAM). Đối với A100 trong FP32:

Tính toán tối đa: ~19,5 TFLOP/s
Băng thông HBM cao nhất: ~2 TB/s
Điểm sườn: ~9,75 FLOP/byte

Bên dưới sườn núi: trí nhớ bị ràng buộc. Ở trên: giới hạn tính toán. Lỗi phổ biến là coi băng thông DRAM là dòng duy nhất quan trọng.

Đường mái dựa trên L2 có đường gờ ở mức ~4,9 FLOP/byte. Đường mái dựa trên L1 có đường gờ ở mức ~1 FLOP/byte.

Một hạt nhân có tái sử dụng L1 mạnh có thể được tính toán ở cường độ số học trông giống như bị giới hạn bộ nhớ trên đường mái DRAM.

Một hạt nhân tấn công L2 sẽ hoạt động kém hơn đường giới hạn DRAM vì băng thông hiệu dụng của nó thấp hơn mức đỉnh theo lý thuyết. Biểu đồ đường mái của NCU hiển thị đồng thời cả ba.

Chẩn đoán chính xác đầu tiên là phân tích băng thông theo cấp bậc. Không phải “nó bị ràng buộc bởi trí nhớ”; đó là một thể loại.

Chẩn đoán hữu ích là “nó bị giới hạn bộ nhớ ở cấp độ L2, đạt 60% mức đỉnh L2 vì 40% băng thông L2 bị lãng phí đối với dữ liệu không được sử dụng lại và bị loại bỏ trước khi sử dụng lần thứ hai.” Điều đó cho bạn biết cách khắc phục.

Bộ tăng tốc bộ nhớ Tensor

Ampere giới thiệu cp.async. Hopper (sm_90, H100) đã giới thiệu Bộ tăng tốc bộ nhớ Tensor (TMA), ý tưởng tương tự đã được đưa đến kết luận hợp lý.

Với cp.async, lập trình viên vẫn tính toán địa chỉ bộ nhớ chung của mọi phần tử và xây dựng luồng lệnh.

Đối với ô 128×128 BF16, đó là 512 lệnh được vector hóa 128-bit cp.async tiêu tốn băng thông lệnh SMSP, ngay cả khi quá trình truyền không đồng bộ.

TMA chấp nhận một bộ mô tả tensor (địa chỉ cơ sở, kích thước, bước tiến, loại phần tử) và đưa ra một lệnh duy nhất:

 cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
    [smem_dst], [gmem_desc, {coord_y, coord_x}], [mbar];

Một hướng dẫn. Một ô 128×128 BF16. Đơn vị TMA tạo ra tất cả địa chỉ, quản lý tất cả giao dịch và hoàn thành tín hiệu thông qua mbarrier nguyên thủy;

cơ chế đồng bộ hóa nhẹ hơn __syncthreads(), được thiết kế để phối hợp nhà sản xuất-người tiêu dùng mà không có rào cản SM đầy đủ.

Hậu quả: trên Hopper, tỷ lệ lệnh tính toán trên tải trong vòng lặp bên trong GEMM tiếp cận ∞ từ góc nhìn của SMSP.

SMSP chạy wgmma.mma_async liên tục; bộ phận TMA xử lý mọi chuyển động dữ liệu một cách độc lập. CUTLASS 3.x được thiết kế dựa trên mô hình này. Phần VIII sẽ trình bày đầy đủ.

Kết luận

Dòng từ Phần VI, “SMSP có quyền ngay lập tức đưa ra hướng dẫn tiếp theo cho sợi dọc đó”, là điểm mấu chốt mà bài viết này hướng tới.

Hệ thống phân cấp bộ nhớ áp đặt độ trễ không thể thương lượng tính bằng nano giây: 23 chu kỳ đối với bộ nhớ dùng chung, 180 đối với L2, 500 đối với HBM. Những con số này không thay đổi bằng cách phàn nàn về chúng.

Chúng thay đổi bằng cách cấu trúc mã để phát sinh độ trễ trước khi cần kết quả: đưa ra yêu cầu bộ nhớ trong khi tính toán trên dữ liệu đã tải trước đó.

cp.async là cơ chế. Đường ống phần mềm là mô hình. Bộ đệm đôi là trường hợp khả thi tối thiểu. Giao thức cam kết/chờ duy trì tính chính xác trong khi công cụ DMA và công cụ điện toán chạy đồng thời.

Phân tích xung đột ngân hàng và thảo luận Bỏ qua L1 là những phần mở rộng của cùng một ý tưởng: giảm thiểu độ trễ và tối đa hóa băng thông hiệu quả ở mọi cấp độ phân cấp, sao cho đến thời điểm dữ liệu đến quá trình tính toán, nó đã truyền qua phần cứng hiệu quả như vật lý cho phép.

Giới hạn của cách tiếp cận này trên Ampere chính là động lực thúc đẩy TMA trên Hopper: một kiến trúc trong đó khoảng cách giữa những gì lập trình viên thể hiện và những gì phần cứng thực thi thu hẹp hơn nữa, tiếp cận chế độ trong đó lập trình viên mô tả cái gì nên di chuyển và phần cứng quyết định khi nào.

Phần VIII bắt đầu từ đó.