Inside CUDA: Performance Engineering

Motivation & Recap#

Trong bài viết trước — Hello CUDA! — chúng ta đã cùng làm quen với CUDA và khám phá bức tranh tổng thể về kiến trúc GPU. Tiếp nối hành trình đó, trong bài này, chúng ta sẽ nói về một chủ đề cốt lõi đã làm nên sức mạnh (và cả sự phức tạp) của GPU: Performance Engineering.

GPU từ lâu đã được biết đến với khả năng tính toán song song vượt trội. Từ việc huấn luyện các mô hình deep learning hàng tỷ tham số, mô phỏng động lực học phân tử, đến phân tích rủi ro tài chính, tất cả đều cần GPU.

Trong CUDA, sự khác biệt về hiệu suất giữa code chưa tối ưu và code đã tối ưu có thể lên đến hàng chục, thậm chí hàng trăm lần. Để đạt được sự tăng tốc này không chỉ là "viết code GPU", chúng ta phải học cách tư duy song song (thinking in parallel). Điều này đòi hỏi sự hiểu biết sâu sắc về cả thuật toán song song và cách chúng tương tác trực tiếp với phần cứng. ¹

Trong bài viết này, chúng ta sẽ tìm hiểu cách giữ cho các lõi GPU luôn bận rộn và khai thác tối đa sức mạnh của phần cứng. Thay vì đi sâu vào các thuật toán song song, chúng ta sẽ tập trung vào những kỹ thuật nền tảng quyết định hiệu năng thực tế của mỗi kernel.

Hardware & Experimental Setup#

Hiệu suất luôn bị giới hạn bởi phần cứng. Một kỹ thuật tối ưu hóa có thể hoạt động hiệu quả trên kiến trúc này nhưng lại kém hiệu quả trên kiến trúc khác.

Để đo lường một cách khách quan, chúng ta sẽ sử dụng hai nền tảng phần cứng khác nhau:

• NVIDIA GeForce MX330: Một GPU laptop cũ, yếu, đại diện cho phần cứng cấp thấp (kiến trúc Pascal).
• NVIDIA RTX A4000: Một GPU hiện đại (workstation) mạnh mẽ dựa trên kiến trúc Ampere.

Việc đo lường trên cả hai sẽ cho chúng ta thấy kỹ thuật nào mang lại lợi ích tổng quát (hiệu quả trên cả hai) và kỹ thuật nào chỉ hiệu quả trên một kiến trúc cụ thể.

Dưới đây là bảng so sánh thông số kỹ thuật chính của hai GPU này:

Phân tích nhanh: Sự khác biệt là rất lớn. A4000 (Ampere) không chỉ có nhiều SM gấp 16 lần (48 vs 3) mà còn có băng thông bộ nhớ lớn hơn gần 8 lần (448 vs 56.1 GB/s). Điều này dự báo trước rằng các kernel bị giới hạn bởi bộ nhớ (memory-bound) sẽ được hưởng lợi đáng kể trên A4000, trong khi các kernel giới hạn bởi tính toán (compute-bound) sẽ được hưởng lợi từ số lượng SM/CUDA core khổng lồ.

Performance Optimization Techniques#

Chúng ta sẽ khám phá một số kỹ thuật tối ưu hóa quan trọng nhất trong CUDA, bắt đầu từ công cụ mạnh mẽ nhất.

1. Shared Memory (SMEM)#

Truy cập Global Memory (VRAM của GPU) là một trong các thao tác tốn kém nhất trong một kernel CUDA. Nó có latency (độ trễ) rất cao, có thể lên tới hàng trăm chu kỳ xung nhịp. Nếu kernel của bạn liên tục đọc và ghi từ global memory, hiệu năng sẽ bị "thắt cổ chai" nghiêm trọng.

Shared Memory (SMEM) là một vùng nhớ on-chip (nằm ngay bên trong SM), có dung lượng nhỏ (thường từ 48KB đến 128KB mỗi SM) nhưng tốc độ truy cập và băng thông cực cao. Độ trễ của SMEM thấp hơn L2 cache và chỉ cao hơn một chút so với register, khiến nó trở thành một công cụ hiệu quả để tái sử dụng dữ liệu (data reuse) trong một block. ¹

Chiến lược chung là:

• Load một khối dữ liệu từ Global Memory vào Shared Memory (chỉ một lần, và cố gắng load một cách coalesced - xem mục 3).
• Để các thread trong block thực hiện nhiều phép tính toán bằng cách truy cập dữ liệu trong Shared Memory.
• Ghi kết quả cuối cùng trở lại Global Memory.

Ví dụ 1: Nhân ma trận (Matrix Multiplication)

Hãy xem xét bài toán nhân hai ma trận $C = A \times B$ (kích thước $M \times N$ và $N \times K$).

Naive Kernel (Không dùng SMEM)

Kernel này đơn giản, mỗi thread tính một phần tử của ma trận C. Tuy nhiên, mỗi thread phải thực hiện N lần đọc từ ma trận A và N lần đọc từ ma trận B—tất cả đều từ Global Memory.

__global__ void matrix_multiplication_naive(
    const float *A, const float *B, float *C, 
    int M, int N, int K)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
 
    if (row < M && col < K)
    {
        float sum = 0.0f;
        // Lặp N lần, mỗi lần đều truy cập Global Memory
        for (int k = 0; k < N; ++k)
        {
            sum += A[row * N + k] * B[k * K + col];
        }
        C[row * K + col] = sum;
    }
}

Kernel Tối ưu (Tiling & SMEM)

Chúng ta sử dụng kỹ thuật gọi là "tiling" (chia ô). Mỗi block sẽ chịu trách nhiệm tính một "tile" (ô), ví dụ $16 \times 16$, của ma trận C. Để làm điều này, nó sẽ lặp và tải các tile $16 \times 16$ tương ứng từ A và B vào Shared Memory.

Ý tưởng thực thi của mỗi thread trong block:

• Vòng lặp (theo số lượng tile)
• Load: Load một phần tử của tile A và một phần tử của tile B (từ Global Memory) vào SMEM.
• Đồng bộ (sync): Đợi tất cả các thread trong block load xong (__syncthreads();).
• Tính toán: Tính toán tổng tích của các tile con (hiện đang nằm trong SMEM).
• Đồng bộ (sync): Đợi tất cả các thread tính toán xong trước khi load tile tiếp theo (__syncthreads();).
• Ghi: Ghi kết quả tổng cuối cùng vào ma trận C ở Global Memory (chỉ một lần khi kết thúc).

#define TILE_SIZE 16
 
__global__ void matrix_multiplication_smem(
    const float *A, const float *B, float *C, 
    int M, int N, int K)
{
    // Khai báo Shared Memory cho các tile A và B
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
 
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
 
    int tx = threadIdx.x; // Tọa độ x bên trong tile
    int ty = threadIdx.y; // Tọa độ y bên trong tile
 
    float sum = 0.0f;
 
    // Lặp qua tất cả các tile cần thiết
    for (int t = 0; t < (N + TILE_SIZE - 1) / TILE_SIZE; ++t)
    {
        // 1. Load tile từ Global Memory vào Shared Memory
        // Tính toán chỉ số cột 'A' và hàng 'B' trong global memory
        int A_col = t * TILE_SIZE + tx;
        int B_row = t * TILE_SIZE + ty;
 
        // Mỗi thread load 1 phần tử của A vào As
        if (row < M && A_col < N) {
            As[ty][tx] = A[row * N + A_col]; // ty, tx là tọa độ trong SMEM
        } else {
            As[ty][tx] = 0.0f; // Padding nếu ra ngoài biên
        }
 
        // Mỗi thread load 1 phần tử của B vào Bs
        if (B_row < N && col < K) {
            Bs[ty][tx] = B[B_row * K + col];
        } else {
            Bs[ty][tx] = 0.0f; // Padding nếu ra ngoài biên
        }
 
        // 2. Đồng bộ để đảm bảo tất cả thread đã load xong
        __syncthreads();
 
        // 3. Tính toán từ Shared Memory
        #pragma unroll // Gợi ý cho compiler unroll vòng lặp này
        for (int i = 0; i < TILE_SIZE; ++i)
        {
            sum += As[ty][i] * Bs[i][tx];
        }
 
        // 4. Đồng bộ trước khi sang tile tiếp theo
        __syncthreads();
    }
 
    // 5. Ghi kết quả cuối cùng ra Global Memory
    if (row < M && col < K)
    {
        C[row * K + col] = sum;
    }
}

Kết quả (Ví dụ 1: Nhân ma trận)

Phân tích: Dữ liệu đo lường này cho chúng ta một bài học cực kỳ quan trọng: Tối ưu hóa phụ thuộc mạnh vào kiến trúc (architecture-dependent).

• Trên MX330 (Pascal): Đúng như dự đoán, kỹ thuật Tiling & SMEM đã mang lại hiệu quả rõ rệt, tăng tốc $\approx$2.35 lần (37.240 / 15.846) so với naive kernel. Trên phần cứng cũ với L2 cache nhỏ và băng thông bộ nhớ thấp, việc giảm tải truy cập global memory bằng SMEM mang lại một cải thiện đáng kể về hiệu năng.

• Trên RTX A4000 (Ampere): Một kết quả đáng ngạc nhiên xuất hiện. Kernel Tiling & SMEM của chúng ta (2.077 ms) thực sự chậm hơn $\approx$1.1 lần so với naive kernel (1.861 ms).

  • Điều này không có nghĩa là SMEM vô dụng. Nó cho thấy naive kernel, mặc dù đơn giản, nhưng đang được hưởng lợi rất nhiều từ các cơ chế phần cứng của Ampere. Với L2 cache lớn (4MB) và băng thông bộ nhớ cực cao (448 GB/s), phần cứng A4000 có thể đã tự động che giấu (hide) phần lớn độ trễ truy cập global memory mà không cần sự can thiệp thủ công của chúng ta.
  • Trong khi đó, kernel Tiling & SMEM lại đưa vào các chi phí "overhead" mới: logic tính toán chỉ số phức tạp hơn và đặc biệt là các lệnh __syncthreads(). Trên một kiến trúc mạnh như Ampere, chi phí cho các rào cản đồng bộ (synchronization barrier) này có thể trở nên lớn hơn lợi ích thu được từ việc sử dụng SMEM, nhất là khi naive kernel vốn đã chạy rất nhanh.

• Thư viện cuBLAS: Trong cả hai trường hợp, cuBLAS đều vượt trội. Đặc biệt trên A4000, nó nhanh hơn naive kernel $\approx$9 lần và nhanh hơn kernel SMEM của chúng ta $\approx$10 lần. Điều này là do cuBLAS không chỉ được tối ưu ở mức assembly mà còn được thiết kế để khai thác triệt để Tensor Cores (có trên kiến trúc Ampere), một loại phần cứng chuyên dụng cho các phép toán nhân ma trận mà kernel float cơ bản của chúng ta không sử dụng tới.

Ví dụ 2: Tối ưu Atomic Operations

atomicAdd() là một hàm cần thiết để tránh race condition (cuộc đua tài nguyên) khi nhiều thread cùng cập nhật một biến trong global memory. ¹

Bài toán: Đếm phần tử mảng (Count Array Elements)

Cho một mảng đầu vào input (kích thước N), chúng ta cần đếm tổng số lượng phần tử trong mảng có giá trị bằng một hằng số K cho trước. Kết quả tổng này phải được tích lũy và lưu trữ vào một biến duy nhất output trong global memory.

Naive Kernel (Naive Atomic)

Giải pháp đơn giản nhất là mỗi thread tìm thấy một phần tử hợp lệ sẽ gọi atomicAdd() trực tiếp lên biến output trong global memory.

__global__ void count_equal_naive(
    const int *input, int *output, int N, int K)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N && input[idx] == K)
    {
        // Rất nhiều lệnh atomic, gây "thắt cổ chai"
        atomicAdd(output, 1);
    }
}

Vấn đề: Nếu input có rất nhiều phần tử bằng K, hàng ngàn thread sẽ cùng lúc "tranh chấp" (contention) để gọi atomicAdd vào cùng một địa chỉ. Các truy cập này sẽ bị tuần tự hóa (serialize), làm mất hoàn toàn tính song song.

Kernel Tối ưu (SMEM & Parallel Reduction)

Chiến lược tốt hơn là giảm thiểu số lần gọi atomicAdd vào global memory. Chúng ta chỉ gọi nó một lần cho mỗi block.

• Grid-Stride Loop: Mỗi thread sẽ xử lý nhiều phần tử (cách nhau một stride). Đây là một pattern mạnh mẽ để xử lý các mảng có kích thước bất kỳ.
• Mỗi thread đếm cục bộ số phần tử nó tìm thấy.
• Kết quả đếm cục bộ này được lưu vào một mảng trong Shared Memory (I[]).
• Parallel Reduction: Chúng ta thực hiện một phép tính tổng song song (reduction) ngay trong shared memory để cộng tất cả các biến đếm cục bộ của block lại. Phép 'reduction' này sử dụng SMEM và các luồng trong block, hoàn toàn không có tranh chấp như atomicAdd vào global memory. Nó cực kỳ hiệu quả.
• Cuối cùng, chỉ một thread (ví dụ threadIdx.x == 0) trong block gọi atomicAdd một lần duy nhất để cộng tổng của block vào output.

__global__ void count_equal_optimized(
    const int *input, int *output, int N, int K)
{
    // Dùng 'extern' để cấp phát SMEM động khi launch kernel
    extern __shared__ int I[];
    
    int g_idx = blockIdx.x * blockDim.x + threadIdx.x;
    int l_idx = threadIdx.x;
    // Stride là tổng số thread trong grid
    int stride = gridDim.x * blockDim.x;
 
    // 1. Mỗi thread đếm cục bộ (Grid-Stride Loop)
    int count = 0;
    while (g_idx < N)
    {
        if (input[g_idx] == K)
        {
            count++;
        }
        g_idx += stride; // Nhảy đến phần tử tiếp theo
    }
    
    // 2. Lưu kết quả cục bộ vào SMEM
    I[l_idx] = count;
    __syncthreads();
 
    // 3. Thực hiện Parallel Reduction trong SMEM
    for (int s = blockDim.x / 2; s > 0; s >>= 1)
    {
        if (l_idx < s)
        {
            I[l_idx] += I[l_idx + s];
        }
        __syncthreads();
    }
 
    // 4. Chỉ thread 0 ghi tổng của block ra Global Memory
    if (l_idx == 0)
    {
        atomicAdd(output, I[0]);
    }
}

Kết quả (Ví dụ 2: Tối ưu Atomic Operations)

Phân tích: Kết quả này tiếp tục củng cố nhận định về tầm quan trọng của kiến trúc phần cứng.

• Trên MX330 (Pascal): Kỹ thuật tối ưu (giảm tải atomic vào global memory) đã mang lại hiệu quả $\approx$1.34 lần (12.271 / 9.1230). Điều này hợp lý vì trên kiến trúc Pascal cũ, việc tranh chấp (contention) atomicAdd trên global memory là một "thắt cổ chai" (bottleneck) tốn kém. Việc chuyển sang reduction trên SMEM đã giảm đáng kể sự tranh chấp này.

• Trên RTX A4000 (Ampere): Hiệu năng tăng tốc gần như không tồn tại, chỉ cải thiện $\approx$1% (1.000 / 0.990).

  • Các kiến trúc GPU hiện đại như Ampere (Compute Capability 8.6) có hệ thống xử lý atomic hiệu quả hơn rất nhiều so với Pascal. Các phép atomicAdd vào global memory thường được xử lý và hợp nhất (coalesced) hiệu quả ngay tại L2 cache, làm giảm đáng kể chi phí tranh chấp.
  • Kernel "tối ưu" của chúng ta đã thay thế một bottleneck (global atomics) vốn đã được phần cứng Ampere giải quyết gần như triệt để, bằng một loạt các thao tác khác: ghi vào SMEM, nhiều lệnh __syncthreads(), và logic phức tạp.
  • Kết quả là, tổng thời gian overhead của phép parallel reduction trong SMEM gần như bằng với thời gian tiết kiệm được từ việc giảm tải global atomic. Nói cách khác, chúng ta đã "tối ưu hóa" một vấn đề không còn là vấn đề lớn trên phần cứng hiện đại.

Kết luận chung: Những kỹ thuật tối ưu hóa cổ điển (như SMEM tiling, giảm tải atomic) vẫn cực kỳ giá trị, nhưng hiệu quả của chúng không còn là tuyệt đối. Một kỹ thuật có thể là "cứu cánh" trên thế hệ phần cứng này, nhưng lại trở nên kém hiệu quả (thậm chí phản tác dụng) trên thế hệ kế tiếp do những cải tiến trong bộ đệm (cache), băng thông, và các đơn vị xử lý chuyên dụng (như L2 cache atomics, Tensor Cores).

2. Memory Coalescing#

Đây là một trong những khái niệm quan trọng nhất khi tối ưu truy cập Global Memory.

Khi 32 thread trong một warp (một nhóm 32 thread thực thi đồng bộ) cùng truy cập global memory, phần cứng GPU sẽ cố gắng "hợp nhất" (coalesce) 32 yêu cầu này thành càng ít transaction (giao dịch bộ nhớ) càng tốt.

• Trường hợp lý tưởng (Coalesced): 32 thread truy cập 32 vị trí liên tục trong bộ nhớ (ví dụ A[idx], A[idx+1], ..., A[idx+31]). Phần cứng GPU đọc bộ nhớ theo các 'segment' 32-byte hoặc 128-byte. Khi truy cập là coalesced, một warp (32 thread) đọc 32 giá trị float (32 * 4 = 128 bytes) chỉ cần một giao dịch 128-byte duy nhất.
• Trường hợp xấu (Uncoalesced): 32 thread truy cập 32 vị trí ngẫu nhiên hoặc cách xa nhau (strided access). Ví dụ, nếu 32 thread truy cập A[idx * 100], chúng có thể rơi vào 32 segment bộ nhớ khác nhau. GPU sẽ phải thực hiện 32 giao dịch 128-byte riêng biệt, lãng phí 31/32 băng thông!

Trong ví dụ nhân ma trận naive, việc truy cập A[row * N + k] (theo hàng) thường là coalesced (vì các thread liền kề trong warp có col liền kề). Ngược lại, truy cập B[k * K + col] (theo cột) thường là uncoalesced, gây lãng phí băng thông nghiêm trọng. ¹

3. Bank Conflicts#

Vấn đề này xảy ra trên Shared Memory. SMEM không phải là một khối đồng nhất; nó được chia thành 32 "bank" (ngân hàng) bộ nhớ. Các thread trong một warp có thể truy cập SMEM song song nếu chúng truy cập vào các bank khác nhau.

Bank Conflict: Xảy ra khi hai hay nhiều thread trong một warp cố gắng truy cập vào các địa chỉ nằm trên cùng một bank. Các truy cập này sẽ bị tuần tự hóa (serialize), làm mất tốc độ của SMEM.

Quy tắc: Cụ thể, với các SMEM có word-size là 4-byte (dành cho float hoặc int), bank $k$ sẽ chứa các word 4-byte ở địa chỉ $addr$ sao cho $(addr / 4) \pmod{32} = k$.

Ví dụ kinh điển:

__shared__ float A[32][32];

• Truy cập theo hàng: A[my_row][threadIdx.x]
  • 32 thread truy cập A[r][0], A[r][1], ..., A[r][31].
  • Các địa chỉ này nằm liên tiếp, do đó chúng rơi vào 32 bank khác nhau (Bank 0, 1, 2, ...).
  • Kết quả: Rất nhanh, không conflict.
• Truy cập theo cột: A[threadIdx.x][my_col]
  • Thread 0 truy cập A[0][c] (ví dụ Bank 1)
  • Thread 1 truy cập A[1][c] (cách A[0][c] 32 phần tử). Địa chỉ của nó cũng mod 32 và ra Bank 1.
  • Thread 2 truy cập A[2][c] (cách A[0][c] 64 phần tử). Địa chỉ của nó cũng mod 32 và ra Bank 1.
  • Kết quả: Thảm họa! 32-way bank conflict! 32 thread cùng truy cập Bank 1 và bị tuần tự hóa.

Hãy hình dung 32 bank như 32 quầy thu ngân:

• Truy cập theo hàng: 32 người (thread) đi đến 32 quầy khác nhau (Bank 0...31). Tất cả được phục vụ song song.
• Truy cập theo cột: Cả 32 người (thread) cùng xếp hàng tại một quầy duy nhất (ví dụ Bank 1). 31 người phải chờ.

Cách tránh: Padding (đệm). Chúng ta phá vỡ pattern truy cập bằng cách thay đổi kích thước mảng:

// Thêm 1 cột đệm
__shared__ float A[32][33]; 

Bây giờ, khi truy cập theo cột A[threadIdx.x][my_col]:

• Thread 0 truy cập A[0][c] (Bank c % 32)
• Thread 1 truy cập A[1][c] (cách 33 phần tử). Bank (c + 33) % 32 = Bank (c+1) % 32
• Thread 2 truy cập A[2][c] (cách 66 phần tử). Bank (c + 66) % 32 = Bank (c+2) % 32
• Kết quả: Không còn conflict! Chúng ta đã "hy sinh" một chút SMEM để đổi lấy tốc độ truy cập song song.

4. Occupancy#

Occupancy là tỷ lệ giữa số lượng warp đang hoạt động trên một SM so với số lượng warp tối đa mà SM đó có thể hỗ trợ (ví dụ: 32 active warps / 64 max warps = 50% occupancy).

Occupancy là yếu tố cốt lõi của Latency Hiding (Che giấu độ trễ). Khi một warp phải dừng lại (stall) – ví dụ như chờ dữ liệu từ global memory – bộ lập lịch (scheduler) của SM có thể ngay lập tức chuyển sang thực thi một warp khác đang ở trạng thái sẵn sàng.

• Occupancy thấp: SM không có đủ warp sẵn sàng để chuyển đổi. Khi warp duy nhất bị stall, toàn bộ SM sẽ bị "chết" (idle), lãng phí tài nguyên tính toán.
• Occupancy cao: SM có nhiều warp để lựa chọn, giúp che giấu độ trễ bộ nhớ và giữ cho các lõi tính toán luôn bận rộn.

Đây là sự khác biệt cơ bản: CPU là latency-oriented (cố gắng hoàn thành 1 tác vụ thật nhanh) trong khi GPU là throughput-oriented (cố gắng hoàn thành nhiều tác vụ nhất trong một đơn vị thời gian). Occupancy cao là chìa khóa cho mô hình throughput này.

Occupancy bị giới hạn bởi tài nguyên nào cạn kiệt trước tiên trên SM:

• Registers: Nếu mỗi thread dùng quá nhiều register, SM sẽ không đủ register để chứa nhiều thread $\rightarrow$ occupancy giảm.
• Shared Memory: Nếu mỗi block dùng quá nhiều SMEM, SM sẽ không đủ SMEM để chứa nhiều block $\rightarrow$ occupancy giảm.
• Số lượng thread/block: Nếu đặt số lượng thread mỗi block quá thấp (ví dụ 64), bạn sẽ không bao giờ đạt được occupancy cao vì một warp là 32 thread.
• Giới hạn phần cứng: Mỗi SM có một số lượng block tối đa và thread tối đa mà nó có thể quản lý. Ví dụ, A4000 (CC 8.6) có thể xử lý tối đa 1536 thread / SM (tương đương 48 warp) và 32 block / SM. ¹

5. Kernel Fusion#

Đây là kỹ thuật gộp nhiều kernel (chạy tuần tự) thành một kernel duy nhất.

Ví dụ 1: Thay vì chạy:

• kernel_add(A, B, C); // C = A + B
• kernel_scale(C, D); // D = C * alpha

Chúng ta sẽ gộp chúng lại thành:

• kernel_add_and_scale(A, B, D); // D = (A + B) * alpha

Ví dụ 2: Phép toán 'SAXPY' (Y = a*X + Y) hoặc Y = a*X + b.

• Chưa tối ưu:
  • kernel_scale(X, a, Temp); // Temp = a*X (Ghi Temp ra Global Mem)
  • kernel_add(Temp, b, Y); // Y = Temp + b (Đọc Temp từ Global Mem)
• Đã tối ưu (Fusion):
  • kernel_fused(X, a, b, Y); // Một kernel duy nhất tính Y[i] = a*X[i] + b[i]. Temp chỉ tồn tại trong register, loại bỏ hoàn toàn việc đọc/ghi Temp ra global memory.

Lợi ích:

• Giảm chi phí launch kernel: Mỗi lần gọi __global__ đều có một chi phí (overhead) nhỏ.
• Loại bỏ truy cập Global Memory: Đây là lợi ích lớn nhất. Trong các ví dụ trên, biến trung gian (C hoặc Temp) phải được ghi ra global memory, và kernel tiếp theo lại phải đọc nó từ global memory. Kernel hợp nhất có thể giữ giá trị trung gian này tạm thời trong register, loại bỏ hoàn toàn 2 lượt truy cập global memory đắt đỏ.

Đánh đổi:

• Tăng độ phức tạp: Kernel hợp nhất khó viết và debug hơn.
• Tăng áp lực tài nguyên: Kernel mới làm nhiều việc hơn, có thể cần nhiều register hơn (vì phải giữ nhiều biến tạm), điều này có thể làm giảm occupancy.

Kernel fusion là một sự đánh đổi: bạn giảm yêu cầu về băng thông bộ nhớ, nhưng có thể tăng yêu cầu về tài nguyên tính toán trên SM. Cần phải đo đạc để xem sự đánh đổi có xứng đáng hay không.

Conclusion#

Tối ưu hóa GPU là một quá trình đòi hỏi cả hiểu biết lý thuyết lẫn kinh nghiệm thực hành. Nó buộc chúng ta phải nhìn lại những bài toán tưởng như quen thuộc — như nhân ma trận hay tính tổng mảng — dưới một góc nhìn song song hoàn toàn mới.

Cốt lõi của quá trình này là xác định điểm nghẽn hiệu năng (bottleneck) trong kernel: liệu nó đang bị giới hạn bởi băng thông bộ nhớ (memory-bound), tính toán (compute-bound) hay độ trễ (latency-bound). Sử dụng các công cụ như NVIDIA Nsight là tối quan trọng để 'profile' kernel và tìm ra chính xác nơi cần tập trung nỗ lực.

Chúng ta đã đi qua các kỹ thuật cơ bản nhưng cực kỳ quan trọng. Trong các bài viết tương lai, chúng ta có thể khám phá các chủ đề nâng cao hơn như Dynamic Parallelism, sức mạnh của Tensor Cores (Tensor Cores khả dụng từ kiến trúc Volta (CC ≥ 7.0) trở lên) cho deep learning, và các kỹ thuật xử lý bất đồng bộ (Asynchronous Operations) với streams.