CUDA Execution Model

Kiến trúc của GPU được xây dựng xoay quanh một mảng có thể nhân rộng của Streaming Multiprocessors (SM). SM là trái tim của GPU

Mỗi GPU thiết kế để hỗ trợ việc chạy hàng ngàn thread trong cùng một thời điểm.

CUDA sử dụng kiến trúc SIMT (Single Instruction Multiple Thread) để quản lý và chạy 32 thread trong một nhóm gọi là Wraps. Tất cả thread trong cùng một wrap chạy cùng một lệnh, mỗi thread có địa chỉ lệnh riêng (instruction address counter) và trangjg thái register (register state) riêng và thực hiện lệnh trên dữ liệu riêng của thread.

• Shared Memory được chia đều giữa các block trên một SM

• Registers được chia đều giữa các thread. Thread trong cùng một block có thể phối hợp và trao đổi thông qua những tài nguyên này.

• Chia sẻ dữ liệu giữa những thread chạy song song có thể gây ra hiện tượng RACE CONDITION: đa luồng tương tác với cùng một dữ liệu với thứ tự không định trước, dẫn đến các hiện tượng bất thường

Số 32 là số đặc biệt trong CUDA đến từ phần cứng cho nên bạn có thể sử dụng con số này để tối ưu.

Tạo ra số lượng là bội của 32 mang đến hiệu quả tính toán.

Một block sẽ được cài đặt chạy trên một SM. Một SM có thể giữ hơn nhiều block ở một thời điểm.

Warps trong một block có thể hẹn lịch thực thi theo bất kỳ thứ tự nào, số lượng warps active tại một thời điểm bị giới hạn bởi SM.

Warp là đơn vị tính toán cơ bản của một SM.

Khi bạn chạy một grid các block các thread, thread trong block sẽ được chia thành các nhóm 32 thread gọi là wrap. 32 thread này sẽ cùng chạy một lệnh.

Block có thể cài đặt dưới dạng 1D, 2D hay 3D các thread tuy nhiên dưới góc nhìn của wrap, mọi thứ được tổ chức dưới dạng mảng 1D.

• Wrap 0 các thread với ID từ 1 đến 31

• Wrap 1 các thread với ID từ 32 đến 63

• Wrap 2 các thread với ID từ 64 đến 95

• Wrap 3 các thread với ID từ 96 đến 127

Ta có thể hoàn toàn chuyển chiều block các thread từ 2D hay 3D về 1D như những bài trước:

• Với 2D Block chuyển về 1D chỉ mục thread sẽ là threadIdx.y * blockDim.x + threadIdx.x

• Với 3D Block chuyển về 1D chỉ mục thread sẽ là threaIdx.z * blockDim.x * blockDimy + threadIdx.y * blockDim.x + threadIdx.x

Số lượng wrap sẽ bằng số lượng threads mỗi block chia cho chiều wrap (32).

Nếu số lượng thread trong block không phải bội 32 một số thread trong wrap cuối cùng sẽ để trạng thái không hoạt động.

Với Block có 74 thread, phần cứng sẽ phân bổ thành 3 Wrap.

Wrap cuối cùng sẽ không hoạt động, mặc dù 10 thread trong wrap này vẫn đang sử dụng nguồn lực của SM, ví dụ registers.

Cấu trúc điều khiển (Control Flow) là nền tảng của các ngôn ngữ lập trình.

GPUs hỗ trợ nền tảng cấu trúc điều khiển với cú pháp C.

CPUs có kiến trúc phức tạp để xử lý thực hiện dự đoán nhánh. Branch prediction (dự đoán nhánh) là một kỹ thuật trong vi kiến trúc của bộ xử lý nhằm cải thiện hiệu suất thực thi của các lệnh nhánh như if, else, và switch. Khi bộ xử lý gặp một lệnh nhánh, nó phải quyết định trước hướng đi nào sẽ được thực hiện (nhánh nào sẽ được theo).

Thay vì dừng và chờ cho đến khi điều kiện của nhánh được đánh giá, bộ xử lý sẽ "dự đoán" nhánh nào sẽ được thực hiện và bắt đầu thực hiện nó trước khi biết kết quả thực sự của điều kiện. Nếu dự đoán đúng, quá trình sẽ tiếp tục mà không bị gián đoạn. Nếu dự đoán sai, bộ xử lý phải hủy bỏ các thao tác đã thực hiện và thực hiện lại theo nhánh đúng, gây ra sự lãng phí thời gian (gọi là pipeline stall).

GPUs đơn giản hơn khi thiếu vắng cơ chế dự đoán nhánh bởi vì tất cả các thread trong một wrap phải thực thi cùng một lệnh trong cùng một cycle. Đây sẽ trở thành vấn đề nếu các luồng trong cùng một wrap thực thi các nhánh khác nhau.

Ví dụ đoạn code

if (cond) {
    ...
} else {
    ...
}

Giả sử có 16 thread trong một wrap thực thi đoạn code trên khi cond là true và 16 thread lại thực hiện code khi cond là false.

Một nửa sẽ thực thi code trong if block và nửa còn lại thực thi trong else block.

Thread trong wrap thực thi các lệnh khác nhau gọi là phân kỳ warp - warp divergence. Phân kỳ warp gây ra sự giảm hiệu năng đánh kể.

Để đạt được hiệu năng tốt, bạn cần tránh việc thực thi các lệnh khác nhau trong cùng một wrap. Chúng ta có thể chia data theo cách mà đảm bảo tất cả các thread trong cùng một wrap thực thi chung một lệnh.

Ví dụ bạn có hai nhánh, bạn có thể chia thread chẵn thực thi code trong thân hàm if còn thread lẻ thực hiện code trong thân hàm else.

__global__ void mathKernel1(float *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0.0f;

    if (tid % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }

    c[tid] = a + b;
}

Bạn có thể thay thế code này như sau:

Nếu bạn lặp data sử dụng phương pháp warp, bạn có thể tránh hiện tượng phân kỳ warp 100%. Hàm

(tid / warpSize) % 2 == 0

ép buộc độ chi tiết của nhánh phải là bội số của warp size; các even warps sẽ thực hiện điều kiện if, còn các odd warps sẽ thực hiện điều kiện else. Kernel này tạo ra cùng kết quả đầu ra, nhưng theo một thứ tự khác.”

__global__ void mathKernel2(void) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float a, b;
    a = b = 0.0f;

    // tid / warpSize gives the warp number the thread belongs to.
    // The % 2 == 0 checks if that warp number is even.

    if ((tid / warpSize) % 2 == 0) {
        a = 100.0f;
    } else {
        b = 200.0f;
    }

    c[tid] = a + b;
}

Đo lường hiệu năng của code