Bài 10: Bandwidth, Throughput và Latency trong Profiling

Trong bài viết này, chúng ta sẽ khám phá ba khái niệm quan trọng trong việc profiling hiệu suất của một đoạn mã hoặc chương trình: Bandwidth, Throughput và Latency. Ngoài việc tìm hiểu từng khái niệm cụ thể, chúng ta cũng sẽ thấy sự liên kết giữa chúng và cách đánh giá hiệu suất toàn diện.

Bandwidth, Throughput và Latency là gì?

Khi đánh giá một đoạn mã, ba yếu tố này thường được xem xét một cách đồng thời. Thế nhưng, rất dễ để nhầm lẫn giữa chúng, dẫn đến việc đưa ra kết luận không chính xác về hiệu suất. Mỗi máy tính có thể có thông số về Latency hoặc Bandwidth khác nhau, vì vậy cần phải cung cấp thông tin đầy đủ để phản ánh hiệu suất thực sự của đoạn mã.

1. Latency (Thời gian trễ)

Latency (s) là thời gian cần để hoàn thành một tác vụ. Một điều cần lưu ý là khi profiling mã (sử dụng cudaEvent_t start, stop), hiệu suất có thể bị ảnh hưởng. Do đó, khi chạy chương trình cuối cùng, hãy nhớ xóa các dòng mã profiling.

Thay vì sử dụng cudaEvent_t start, stop để kiểm tra Latency, bạn có thể sử dụng Nsight Systems bằng lệnh sau:

Copy

nsysp profile -o timeline --trace cuda,nvtx,osrt,openacc ./a.out

Tất cả các hoạt động trên GPU sẽ được đo lường chi tiết.

2. Bandwidth (Băng thông)

Bandwidth (GB/s) thể hiện tốc độ đọc/ghi dữ liệu. Khi nói về Bandwidth, có hai khái niệm cần chú ý:

Theoretical Peak Bandwidth: tốc độ lý tưởng chuyển giao dữ liệu. Để chuyển đổi từ bit sang byte, bạn cần chia cho 8 và để đổi sang GB/s, chia cho 10^9. Các hệ SDR và DDR cũng ảnh hưởng đến cách tính toán này.

Để xác định máy của bạn là DDR hay SDR, hãy sử dụng lệnh:

Copy

sudo lshw -c memory

Effective Bandwidth: tốc độ thực tế khi thực hiện chuyển giao dữ liệu. Công thức tính toán là:

R(B): số byte được đọc mỗi kernel,
W(B): số byte được ghi mỗi kernel,
t(s): latency.

Ví dụ về mã

Chúng ta sẽ thực hiện bài toán sau: y[i] = a*x[i] + y[i] với N = 20 * (1 << 20).

c Copy

__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];
}

Mã profiling có thể được thực hiện như sau:

c Copy

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);

// Thực hiện SAXPY với 1M phần tử
axpy<<<(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);

Mã nguồn trên sẽ cho phép bạn tính toán hiệu suất của đoạn mã. Sự khác biệt giữa các kết quả đo có thể được xác định bằng cách sử dụng Nsight Compute.

3. Throughput (Thông lượng tính toán)

Throughput (GFLOP/s): Là số lượng phép toán dấu chấm động FLOP mà kernel có thể thực hiện trong một giây. Khi xem xét hiệu suất, câu hỏi thường đặt ra là nên tăng Bandwidth hay Throughput? Đó là tình huống mà một giá trị tăng có thể dẫn đến sự giảm của giá trị khác.

Để xác định giá trị chuẩn, một kỹ thuật hữu ích là sử dụng roofline chart, có thể được đề cập trong các bài viết sau.

Kết luận, việc hiểu rõ Bandwidth, Throughput và Latency sẽ giúp bạn đánh giá và tối ưu hóa chương trình của mình một cách chính xác hơn.
source: viblo

[Hướng dẫn sử dụng công cụ NVIDIA] Bài 10: Hiểu biết về Bandwidth, Throughput và Latency trong Profiling