Tìm hiểu về Shared Memory trong CUDA

Trong bài viết này, chúng ta sẽ cùng khám phá cách sử dụng shared memory trên GPU thông qua CUDA. Trước khi bắt đầu, bạn nên tham khảo bài viết trước về Các bộ nhớ trong GPU để có cái nhìn tổng quan hơn.

Shared Memory là gì?

Shared memory là loại bộ nhớ nhanh nhất trong GPU (sau register file) và phạm vi truy cập của nó giới hạn trong các thread cùng nằm trong một block. Đây là một đặc điểm quan trọng giúp tăng tốc độ xử lý trong các ứng dụng song song.

Khi sao chép dữ liệu từ global memory sang shared memory, chúng ta cần sử dụng hàm syncthreads để đồng bộ các thread trong cùng một block. Điều này giúp tránh tình trạng race condition, bởi vì mặc dù các thread trong một block chạy song song về mặt logic, không phải tất cả các thread đều có thể thực thi cùng một thời điểm trên phần cứng. Một số thread có thể hoàn thành trước, và do đó, việc đồng bộ là cần thiết.

Bạn có thể tham khảo thêm 2 bài viết sau để hiểu rõ hơn về: Data Hazard và Synchronization - Asynchronization.

Quy trình chuyển dữ liệu từ Global Memory sang Shared Memory

Trong bài viết này, chúng ta sẽ tập trung vào khái niệm cũng như cách sử dụng shared memory. Những bài viết sau sẽ hướng dẫn các kỹ thuật nhằm cải thiện và tối ưu hóa việc sử dụng shared memory.

Mã nguồn sử dụng Shared Memory

Chúng ta sẽ áp dụng hai khái niệm quen thuộc là static và dynamic memory trong bối cảnh shared memory:

Mã nguồn cho Static Memory

__global__ void staticReverse(int *data, int n) {
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n - t - 1;
  s[t] = data[t];
  __syncthreads();
  data[t] = s[tr];
}

Mã nguồn cho Dynamic Memory

__global__ void dynamicReverse(int *data, int n) {
  extern __shared__ int s[];
  int t = threadIdx.x;
  int tr = n - t - 1;
  s[t] = data[t];
  __syncthreads();
  data[t] = s[tr];
}

Giải thích quy trình

Trong bài toán của chúng ta, có hai bước chính:

1. Sao chép dữ liệu từ global memory sang shared memory theo thứ tự từ bé đến lớn.
2. Sao chép dữ liệu từ shared memory trở lại global memory theo thứ tự từ lớn đến bé.

Gọi kernel trên GPU

staticReverse<<<1, n>>>(d_data, n);
dynamicReverse<<<1, n, n * sizeof(int)>>>(d_data, n);

Trong đó, cú pháp <<<a, b, c>>> có nghĩa là:

• a: Số block
• b: Số thread trong một block
• c: Kích thước của shared memory

Hy vọng bài viết này sẽ giúp bạn hiểu rõ hơn về khái niệm và cách sử dụng shared memory trên GPU. Hãy theo dõi các bài viết tiếp theo để khám phá các kỹ thuật tối ưu hóa hiệu năng trong lập trình song song.

source: viblo