2.1 Warp và mô hình thực thi

Hiểu sâu đơn vị thực thi cốt lõi của GPU: Warp. Bài viết bắt đầu từ mô hình SIMT, phân tích chi phí của Warp Divergence và cách né, rồi Warp Shuffle để trao đổi dữ liệu giữa các thread với độ trễ gần bằng không, tạo nền tảng về lớp thực thi để viết CUDA Kernel hiệu quả.

Mục lục

• 1. Nhìn Warp từ góc phần cứng
• 2. Mô hình thực thi SIMT
• 3. Warp Divergence: cái giá của nhánh
• 4. Warp Shuffle: trao đổi dữ liệu giữa các thread
• 5. Primitive cấp Warp và mẹo nâng cao
• 6. Case study: reduction cấp Warp
• Tóm tắt
• Danh sách tự kiểm
• Tài liệu tham khảo

---

1. Nhìn Warp từ góc phần cứng

Hình dung một lớp học có 32 học sinh (thread). Giáo viên (scheduler) mỗi lần đưa ra một lệnh, cả lớp cùng làm một động tác. Đó là ý tưởng cốt lõi của Warp: đơn vị nhỏ nhất mà GPU lập lịch và thực thi, luôn gồm 32 thread liên tiếp.

1.1 Quy tắc hình thành Warp

Khi một Thread Block được gán lên SM, phần cứng tự nhóm các thread theo thứ tự threadIdx thành Warp:

Điểm then chốt: Việc chia Warp là hành vi phần cứng; lập trình viên không can thiệp được. Một Block 256 thread sẽ thành 8 Warp. Nếu kích thước Block không phải bội của 32 (ví dụ 48 thread), các vị trí thừa ở Warp cuối được “đệm” nhưng không làm việc hữu ích, lãng phí tài nguyên phần cứng.

1.2 Warp Scheduler

Mỗi SM có vài Warp Scheduler (kiến trúc hiện đại thường 4 cái). Mỗi chu kỳ đồng hồ, một scheduler có thể chọn một Warp sẵn sàng và phát lệnh. Từ đó có cơ chế ẩn độ trễ quan trọng:

Warp A phát lệnh đọc bộ nhớ → chờ dữ liệu về (~400 cycles)
Warp Scheduler chuyển sang Warp B → chạy lệnh tính toán
Chuyển sang Warp C → chạy lệnh tính toán
...
Dữ liệu của Warp A sẵn sàng → được lập lịch lại

Gợi ý: Chuyển giữa các Warp là zero-overhead (zero-overhead context switch), vì trạng thái thanh ghi của mỗi Warp luôn nằm trên SM, không cần save/restore. Đây khác bản chất so với chuyển thread trên CPU.

---

2. Mô hình thực thi SIMT

2.1 SIMT vs SIMD

NVIDIA gọi mô hình thực thi của mình là SIMT (Single Instruction, Multiple Threads), chứ không phải SIMD (Single Instruction, Multiple Data) truyền thống. Khác biệt tinh tế nhưng quan trọng:

Ý tưởng cốt lõi của SIMT: bạn viết mã scalar cho một thread; phần cứng broadcast cùng lệnh lên 32 thread. Ngưỡng vào song song hóa thấp hơn nhiều, nhưng cũng tạo ra một bẫy hiệu năng: phân kỳ nhánh (branch divergence).

2.2 Pipeline thực thi

Luồng thực thi lệnh của một Warp:

graph LR
    A["Fetch"] --> B["Decode"]
    B --> C["Issue"]
    C --> D["Execute"]
    D --> E["Write back"]

Đặc điểm chính:

• Thực thi đồng bộ: 32 thread trong một Warp thực thi cùng một lệnh trong cùng chu kỳ đồng hồ
• Trạng thái độc lập: mỗi thread có file thanh ghi và program counter riêng (về mặt logic)
• Lockstep: khi không có nhánh, mọi thread đi cùng nhịp

---

3. Warp Divergence: cái giá của nhánh

3.1 Warp Divergence là gì

Khi các thread trong cùng một Warp gặp nhánh điều kiện (if/else, switch, số vòng lặp khác nhau), chúng cần đi các đường khác nhau. Phần cứng tại một thời điểm chỉ phát được một lệnh, nên phải thực thi tuần tự mọi nhánh; thread không đi nhánh đó bị tạm tắt (masked off).

Giống như bảo cả lớp “ai biết bơi thì chống đẩy, ai không biết thì gập bụng”, nhưng nội quy thể dục bắt mọi người phải làm cùng động tác. Kết quả: cả lớp chống đẩy trước (người không biết bơi đứng im), rồi cả lớp gập bụng (người biết bơi đứng im). Thời gian là tổng hai động tác, không phải thời gian song song của từng nhóm.

3.2 Ảnh hưởng hiệu năng của Divergence

Xét đoạn mã sau:

__global__ void divergent_kernel(float* data, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid % 2 == 0) {
        // Thread chẵn: đường A (10 lệnh)
        data[tid] = expf(data[tid]);
    } else {
        // Thread lẻ: đường B (10 lệnh)
        data[tid] = logf(data[tid]);
    }
}

Lưu ý: Trong mã trên, nửa trong 32 thread của mỗi Warp đi if, nửa đi else. Phần cứng chạy đường A trước (16 thread active, 16 idle), rồi đường B (16 active, 16 idle). Tổng thời gian khoảng gấp đôi trường hợp không có nhánh.

Nếu điều kiện còn mảnh hơn (32 thread đi 32 đường khác nhau), hiệu năng có thể tụt tới 32 lần.

3.3 Định lượng Divergence

Có thể dùng Branch Efficiency:

$$\text{Branch Efficiency} = \frac{\text{Non-Divergent Branches}}{\text{Total Branches}} \times 100\%$$

Nsight Compute cho thấy tỷ lệ thread active của mỗi lệnh nhánh. Lý tưởng thì mọi nhánh đều 100% active (32/32).

3.4 Chiến lược né

Chiến lược 1: Căn điều kiện nhánh theo Warp

// ❌ Không khuyến nghị: thread kề nhau đi nhánh khác nhau
if (threadIdx.x % 2 == 0) { ... }

// ✅ Khuyến nghị: thread trong cùng Warp đi cùng nhánh
if (threadIdx.x / 32 % 2 == 0) { ... }
// Hoặc nhánh theo warpId
int warpId = threadIdx.x / 32;
if (warpId < 4) { ... }

Chiến lược 2: Thay nhánh bằng phép số học

// ❌ Không khuyến nghị: nhánh điều kiện
if (x > 0) {
    y = x;
} else {
    y = 0;
}

// ✅ Khuyến nghị: ReLU không nhánh
y = x * (x > 0);  // tận dụng chuyển ngầm bool → int
// Hoặc dùng hàm dựng sẵn
y = fmaxf(x, 0.0f);

Chiến lược 3: Sắp xếp lại dữ liệu

Gom trước các phần tử sẽ đi cùng đường vào vị trí liên tiếp, để thread trong cùng Warp chạy cùng đường:

// Tiền xử lý: gom số dương và số âm riêng
// Dương ở data[0..pos_count-1]
// Âm ở data[pos_count..N-1]
// Warp xử lý số dương sẽ không còn phân kỳ nhánh

Chiến lược 4: Thống nhất điều kiện thoát vòng lặp Nói ngắn: biến sự không chắc “thread nào thoát vòng lúc nào” thành predication “trong thân vòng có chạy hay không”. Mọi thread đi cùng cấu trúc vòng, thân thiện hơn với nhánh, instruction cache, và mô hình SIMT của GPU.

// ❌ Không khuyến nghị: mỗi thread số vòng khác nhau
for (int i = 0; i < array_length[tid]; i++) {
    // logic xử lý
}

// ✅ Khuyến nghị: số vòng thống nhất, che bằng điều kiện bên trong
int max_len = block_max_length;  // giá trị max trong block
for (int i = 0; i < max_len; i++) {
    if (i < array_length[tid]) {
        // logic xử lý
    }
}

3.5 Independent Thread Scheduling (Volta+)

Từ kiến trúc Volta, NVIDIA đưa vào independent thread scheduling (lập lịch thread độc lập). Mỗi thread có program counter và call stack thật sự độc lập, nghĩa là:

• Khi phân kỳ nhánh, phần cứng có thể xen kẽ các đường linh hoạt hơn
• Hỗ trợ primitive đồng bộ cấp thread (như __syncwarp())
• Các thread trong Warp có thể đứng ở vị trí mã khác nhau

Lưu ý: Lập lịch độc lập không xóa chi phí hiệu năng của Divergence; nó chỉ làm lịch trình linh hoạt hơn. Muốn hiệu năng tốt vẫn nên tránh phân kỳ nhánh trong cùng Warp.

---

4. Warp Shuffle: trao đổi dữ liệu giữa các thread

4.1 Vì sao cần Warp Shuffle

Cách truyền thống trao đổi dữ liệu giữa thread phải qua shared memory:

Thread A ghi shared memory → __syncthreads() → Thread B đọc shared memory

Hai vấn đề: (1) tốn shared memory quý; (2) barrier đồng bộ thêm độ trễ.

Warp Shuffle cho phép thread trong cùng Warp đọc trực tiếp giá trị thanh ghi của nhau, không qua shared memory, cũng không cần đồng bộ tường minh, giống truyền giấy nháp cho nhau thay vì đặt lên bục rồi lấy.

4.2 Bốn lệnh Shuffle

Chữ ký của mọi lệnh Shuffle:

T __shfl_sync(unsigned mask, T var, int srcLane, int width=32);
T __shfl_up_sync(unsigned mask, T var, unsigned delta, int width=32);
T __shfl_down_sync(unsigned mask, T var, unsigned delta, int width=32);
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=32);

Tham số:

• mask: mặt nạ thread tham gia Shuffle; thường 0xFFFFFFFF cho đủ 32 thread
• var: giá trị thread hiện tại chia sẻ
• width: độ rộng sub-Warp logic (2, 4, 8, 16, 32)

4.3 Minh họa Shuffle

Ví dụ với 8 thread (thực tế là 32; rút gọn để minh họa):

Trạng thái ban đầu: Lane[0]=A  Lane[1]=B  Lane[2]=C  Lane[3]=D  Lane[4]=E  Lane[5]=F  Lane[6]=G  Lane[7]=H

__shfl_down_sync(mask, var, 2):
  Lane[0] nhận giá trị C của Lane[2]
  Lane[1] nhận giá trị D của Lane[3]
  Lane[2] nhận giá trị E của Lane[4]
  ...

__shfl_xor_sync(mask, var, 1):  (đổi cặp kề nhau)
  Lane[0] ↔ Lane[1]
  Lane[2] ↔ Lane[3]
  Lane[4] ↔ Lane[5]
  Lane[6] ↔ Lane[7]

4.4 Ví dụ thực dụng: broadcast trong Warp

__device__ float warp_broadcast(float val, int srcLane) {
    // Mọi thread lấy val của thread srcLane
    return __shfl_sync(0xFFFFFFFF, val, srcLane);
}

// Dùng: để mọi thread trong Warp biết kết quả của lane 0
float result = compute_something();
float broadcast_val = warp_broadcast(result, 0);

4.5 Ví dụ thực dụng: tổng trong Warp

__device__ float warp_reduce_sum(float val) {
    // Butterfly reduce: log2(32) = 5 bước
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    return val;  // chỉ lane 0 giữ kết quả đúng
}

Hình dung quá trình reduce (8 thread):

Ban đầu: [1] [2] [3] [4] [5] [6] [7] [8]

Bước 1 (offset=4): mỗi thread cộng giá trị cách 4 lane
       [1+5] [2+6] [3+7] [4+8] [5] [6] [7] [8]
       = [6]  [8]  [10]  [12]

Bước 2 (offset=2): mỗi thread cộng giá trị cách 2 lane
       [6+10] [8+12] [10] [12] ...
       = [16]  [20]

Bước 3 (offset=1): mỗi thread cộng giá trị cách 1 lane
       [16+20] [20] ...
       = [36]    ← lane 0 giữ tổng mọi phần tử

---

5. Primitive cấp Warp và mẹo nâng cao

5.1 Hàm Warp Vote

Warp Vote để cả Warp “bỏ phiếu” trên một điều kiện boolean:

// true nếu predicate của mọi thread đều true
int __all_sync(unsigned mask, int predicate);

// true nếu predicate của bất kỳ thread nào true
int __any_sync(unsigned mask, int predicate);

// Bitmap 32-bit; mỗi bit là predicate của một lane
unsigned __ballot_sync(unsigned mask, int predicate);

Tình huống thực dụng: tối ưu thoát sớm

__global__ void search_kernel(int* data, int target, int* found) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int pred = (data[tid] == target);

    // Nếu không thread nào trong Warp tìm thấy, cả Warp có thể bỏ qua xử lý sau
    if (__any_sync(0xFFFFFFFF, pred)) {
        // Ít nhất một thread tìm thấy mục tiêu
        if (pred) {
            atomicExch(found, tid);
        }
    }
}

5.2 __activemask() và __syncwarp()

// Lấy mặt nạ các thread đang active (Volta+ mới có ý nghĩa rõ)
unsigned mask = __activemask();

// Barrier đồng bộ trong Warp (đảm bảo các thread chỉ định xong thao tác trước)
__syncwarp(0xFFFFFFFF);

Gợi ý: Trên kiến trúc trước Volta, __syncwarp() thường là no-op (Warp vốn đã lockstep). Từ Volta trở đi, vì lập lịch thread độc lập, đồng bộ Warp tường minh trở nên cần thiết.

5.3 Hàm Warp Match (Volta+)

// Trả về mặt nạ các thread có cùng giá trị val
unsigned __match_any_sync(unsigned mask, T val);

// Nếu mọi thread tham gia có cùng val, trả về mặt nạ đầy đủ; không thì 0
unsigned __match_all_sync(unsigned mask, T val, int* pred);

Ứng dụng: tự nhóm các thread cùng key, dùng cho group-by cấp Warp.

5.4 Dùng tham số width

Tham số width của Shuffle có thể chia một Warp thành nhiều sub-Warp logic:

// Coi Warp 32 thread như 4 nhóm độc lập 8 thread
// Mỗi nhóm tự reduce
float sub_warp_sum(float val) {
    for (int offset = 4; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset, 8);
    }
    return val;
}
// Lane 0, 8, 16, 24 lần lượt giữ tổng của từng nhóm

---

6. Case study: reduction cấp Warp

6.1 Cài đặt reduce cấp Block đầy đủ

Kết hợp Warp Shuffle với shared memory để tổng cấp Block hiệu quả:

__device__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    return val;
}

__global__ void block_reduce_kernel(float* input, float* output, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    // Mỗi thread nạp một phần tử
    float val = (tid < N) ? input[tid] : 0.0f;

    // Bước 1: reduce trong Warp
    val = warp_reduce_sum(val);

    // Bước 2: lane 0 của mỗi Warp ghi kết quả vào shared memory
    __shared__ float warp_sums[32];  // tối đa 32 Warp (1024 thread / 32)
    int laneId = threadIdx.x % 32;
    int warpId = threadIdx.x / 32;

    if (laneId == 0) {
        warp_sums[warpId] = val;
    }
    __syncthreads();

    // Bước 3: Warp đầu tiên reduce warp_sums
    int numWarps = blockDim.x / 32;
    val = (threadIdx.x < numWarps) ? warp_sums[threadIdx.x] : 0.0f;

    if (warpId == 0) {
        val = warp_reduce_sum(val);
    }

    // Lane 0 ghi kết quả reduce của Block này
    if (threadIdx.x == 0) {
        output[blockIdx.x] = val;
    }
}

6.2 So sánh hiệu năng

Ưu điểm của bản Warp Shuffle:

• Giảm mạnh shared memory (nhường cho việc khác)
• Ít barrier đồng bộ hơn
• Giao tiếp register-to-register, độ trễ rất thấp

---

Tóm tắt

---

Danh sách tự kiểm

• Giải thích được vì sao Warp có kích thước 32 và quy tắc hình thành Warp
• Nêu được khác biệt cốt lõi giữa SIMT và SIMD
• Nhận diện Warp Divergence trong mã và ước lượng ảnh hưởng hiệu năng
• Viết lại Kernel có if/else thành bản không nhánh hoặc căn theo Warp
• Dùng đúng bốn lệnh Warp Shuffle (shfl / shfl_up / shfl_down / shfl_xor)
• Dùng Warp Shuffle để cài reduce cấp Warp (sum/max/min)
• Giải thích vai trò của width và cài reduce sub-Warp
• Kết hợp Shuffle và shared memory để hoàn thành reduce cấp Block
• Dùng __ballot_sync để lọc điều kiện cấp Warp
• Nêu ảnh hưởng của independent thread scheduling trên Volta tới lập trình

---

Tài liệu tham khảo

• NVIDIA CUDA C++ Programming Guide - Warp Shuffle Functions
• NVIDIA CUDA C++ Programming Guide - SIMT Architecture
• NVIDIA Volta Architecture Whitepaper
• Using CUDA Warp-Level Primitives - NVIDIA Developer Blog
• Faster Parallel Reductions on Kepler - NVIDIA Developer Blog