Ở bài viết này mình sẽ giới thiệu 3 khái niệm rất quan trọng trong việc profiling là Bandwidth - Throughput - Latency
Bandwidth - Throughput - Latency
Để đánh giá 1 đoạn code, 1 chương trình thì ba khái niệm quan trọng cần được xem xét là bandwidth, throughput và latency . Tuy nhiên, chúng ta thường dễ nhầm lẫn khi chỉ đưa ra một trong ba thông tin này mà không kèm theo các thông tin khác, dẫn đến đánh giá không chính xác về hiệu năng - hiệu suất. Vì mỗi máy tính khác nhau có thể có dẫn đến thông số về latency hoặc bandwidth khác nhau, việc chỉ cung cấp một thông tin sẽ không phản ánh được hiệu năng thực sự của đoạn code.
Latency
Latency(s) : là thời gian để hoàn thành 1 task. 1 lưu ý vô cùng quan trọng là khi profile ( ví dụ như cudaEvent_t start, stop ) sẽ khiến performance bị ảnh hưởng nên vì vậy khi chạy chương trình 1 lần cuối nên xóa các dòng code profile.
Và thay vì dùng cudaEvent_t start, stop để kiểm tra latency thì chúng ta có thể dùng nsight system cho việc đó bằng command
nsys profile -o timeline --trace cuda,nvtx,osrt,openacc ./a.out
Thì lúc này tất cả những thứ chạy trên GPU sẽ được measure 1 cách chi tiết
Và như bài Giới thiệu Nsight Systems - Nsight Compute thì qua các thông số phía trên chúng ta có thể xác định là đoạn code của chúng ta cần cải thiện ở phần cấp phát data cho GPU (cudaMalloc)
Bandwidth
Bandwidth(GB/s) : thể hiện tốc độ read/write data
Khi nhắc tới Bandwidth chúng ta sẽ đề cập tới 2 khái niệm:
- Theoretical Peak Bandwidth: tốc độ lý tưởng trong việc transfer data
Chia cho 8 vì đổi từ bit sang byte
Chia cho 10^9 là dùng để đổi qua GB/s
DDR: double data rate thì là * 2
SDR: single data rate thì * 1
Cách xác định DDR hoặc SDR bằng command:
sudo lshw -c memory
Và vì máy mình là DDR nên
- Effective bandwidth : là tốc độ thực sự trong việc transfer data của kernel
R(B): số lượng byte được read mỗi kernel
W(B): số lượng byte được write mỗi kernel
t(s): latency
Code
Chúng ta sẽ thực hiện bài toán y[i] = a*x[i] + y[i] với N = 20 * (1 << 20)
__global__
void saxpy(int n, float a, float *x, float *y)
{ int i = blockIdx.x*blockDim.x + threadIdx.x; if (i < n) y[i] = a*x[i] + y[i];
}
cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start); // Perform SAXPY on 1M elements saxpy<<<(N+511)/512, 512>>>(N, 2.0f, d_x, d_y); cudaEventRecord(stop); cudaEventSynchronize(stop); float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); printf("time: %f\n", milliseconds); printf("Effective Bandwidth (GB/s): %f", N*4*3/milliseconds/1e6);
} Áp dụng công thức R + W thì N * 3 là vì read a + read y + write y và N * 4 là đổi qua byte ( float = 4 byte)
Và đây là output
Nhưng nếu các bạn profile bằng Nsight compute bằng command:
ncu -o profile --set full ./a.out
Thì sẽ thấy là
Từ đây chúng ta có thể tính toán
Như mình đã đề cập phía trên, nếu chúng ta profile ( cudaEvent_t start, stop ) sẽ khiến performance của code bị ảnh hưởng từ 91GB/s --> 87.8GB/s
Và chúng ta cũng có thể tính toán ngược lại để xác định
Việc bị lệch 1 xíu so với công thức phía trên ( 96 GB/s --> 95,8 GB/s là do các yếu tố như memory và kernel ảnh hưởng )
Từ đây chúng ta có thể kết luận code của chúng ta về mặt transfer data rất ổn - không có sự chênh lệch lớn giữa Theoretical - Effective
Các bạn có thể xác định effective bandwidth nhanh hơn bằng command:
Load: ncu --metrics l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum.per_second ./a.out
Store: ncu --metrics l1tex__t_bytes_pipe_lsu_mem_global_op_st.sum.per_second ./a.out
Ta có thể thấy 60.68 + 30.40 = 91.08 GB/s gần với 96GB/s ==> code ổn
Computational Throughput
Throughput(GFLOP/s): Là số lượng FLOP mà kernel có thể thực hiện trong một giây
FLOP (Floating Point Operations): Là một phép toán dấu chấm động. Một phép toán này có thể là phép cộng, trừ, nhân, chia, hoặc các phép toán phức tạp hơn như căn bậc hai, sin, cos,....
Câu hỏi đặt ra là nên cải thiện bandwidth hay là throughput? Vì ông này tăng là ông khác giảm vậy làm sao biết được giá trị đạt chuẩn?
Tăng bandwidth ==> nhiều data được read/write ==> số lượng phép toán tăng ==> throughput giảm và ngược lại
Để trả lời cho câu hỏi đó chúng ta có 1 kĩ thuật gọi là roofline chart và mình sẽ đề cập ở các bài sau
code mình sẽ để ở đây