Hướng Dẫn Hiểu Về Occupancy Trong Lập Trình CUDA

Trong bài viết trước, mình đã đề cập đến cách chọn số lượng thread phù hợp khi lập trình CUDA. Trong bài này, mình sẽ chia sẻ một phương pháp được sử dụng phổ biến để xác định occupancy, cùng với lý do tại sao việc thử nghiệm nhiều trường hợp với các số lượng thread không phải là lựa chọn tối ưu trong mọi tình huống.

Occupancy Là Gì?

Trước khi đi sâu vào các khái niệm, hãy cùng mình hình dung occupancy. Giả sử bạn có 6 nhân công và 6 công việc. Cách đơn giản nhất là chia mỗi nhân công một công việc. Tuy nhiên, nếu mỗi nhân công có thể thực hiện đồng thời 3 công việc, bạn chỉ cần 2 nhân công để hoàn thành 6 công việc. Điều này không chỉ tiết kiệm nhân lực mà còn giảm chi phí.

Trong ví dụ này, nhân công đại diện cho thread trong CUDA, còn công việc đại diện cho dữ liệu. Câu hỏi đặt ra là làm thế nào để xác định mỗi thread có thể xử lý bao nhiêu dữ liệu? NVIDIA đã phát triển khái niệm occupancy để giải quyết vấn đề này.

Occupancy được sử dụng để xác định số lượng thread tối ưu cho kernel nhằm đạt được hiệu suất cao nhất.

Ví Dụ Thực Tế

Giả sử ta có N = 1.000.000 và một kernel đơn giản như sau:

__global__ void square(int *array, int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    if (idx < N)
    {
        array[idx] *= array[idx];
    }
}

Cách thông thường để cấu hình số lượng block và thread là:

    blockSize = 1024;
    gridSize = (N + blockSize - 1) / blockSize;

Tuy nhiên, liệu con số 1024 đã tối ưu hay chưa? Chúng ta có thể sử dụng công cụ profiling để kiểm tra occupancy:

ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./a.out

Kết quả có thể cho thấy rằng việc sử dụng 1024 threads lãng phí tài nguyên (vì occupancy chỉ đạt 53.73%), chứng tỏ mỗi thread vẫn có thể thực hiện nhiều hơn một công việc.

Công Cụ Tính Toán Occupancy

NVIDIA cung cấp một hàm để xác định occupancy tối ưu:

template<class T>
__inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSize(
    int    *minGridSize,
    int    *blockSize,
    T       func,
    size_t  dynamicSMemSize = 0,
    int     blockSizeLimit = 0)
{
    return cudaOccupancyMaxPotentialBlockSizeVariableSMem(minGridSize, blockSize, func, __cudaOccupancyB2DHelper(dynamicSMemSize), blockSizeLimit);
}

Ở đây, minGridSize là kích thước grid tối thiểu gợi ý để đạt được mức sử dụng tối đa của máy, và blockSize là kích thước block tối ưu để đạt được occupancy cao nhất. Khi sử dụng hàm này, bạn có thể nhận được kết quả rằng số lượng thread cần thiết để đạt 100% occupancy sẽ khác nhau tùy theo cấu hình máy tính.

Sự Khác Biệt Giữa Theoretical Occupancy và Achieved Occupancy

Một điều thú vị là sao 2 giá trị occupancy lại khác nhau? Chúng ta có khái niệm mới là Theoretical Occupancy (occupancy lý thuyết) và Achieved Occupancy (occupancy đạt được). Lý do cho sự khác biệt này là khi chạy mã, nhiều yếu tố khác nhau có thể ảnh hưởng đến số lượng thread mà bạn có thể sử dụng một cách hiệu quả.

Lưu ý: Chúng ta chỉ nên chú trọng vào Achieved Occupancy và không quá quan tâm tới Theoretical Occupancy. Ví dụ:

1. Theoretical 100% và Achieved 50%
2. Theoretical 80% và Achieved 70%

Trong trường hợp trên, ta nên chọn trường hợp thứ hai vì đạt được hiệu suất cao hơn.

Kết Luận

Trong các bài viết tiếp theo, mình sẽ hướng dẫn cụ thể hơn về các phương pháp giúp tăng Achieved Occupancy nhằm tối ưu hóa hiệu suất cho các kernel trong CUDA. Hãy theo dõi để biết thêm chi tiết!
source: viblo