Thuật toán sắp xếp sử dụng CUDA

(ashwanirathee.com)

3 điểm bởi GN⁺ 2025-03-13 | 2 bình luận | Chia sẻ qua WhatsApp

Cách cải thiện hiệu năng của thuật toán sắp xếp thông qua tính toán song song bằng CUDA
Khám phá khả năng cải thiện hiệu năng bằng cách triển khai merge sort cơ bản với CUDA

Merge sort đệ quy cơ bản (triển khai trên CPU)

Thuật toán sắp xếp chia mảng thành hai mảng con, sắp xếp từng mảng rồi hợp nhất lại
Chia mảng theo cách đệ quy, và khi kích thước còn 1 thì thực hiện thao tác hợp nhất
Các điểm chính liên quan đến triển khai
- Sử dụng uint8_t → giảm thiểu sử dụng bộ nhớ với các giá trị nhỏ (0~255)
- Sử dụng long long → có thể xử lý mảng lớn (tối đa 10¹⁸)
- Xác minh kết quả bằng std::sort để so sánh hiệu năng
- Độ phức tạp thời gian: trung bình O(n log n)
- Độ phức tạp không gian: O(n)

Merge sort đệ quy cơ bản trên CUDA

Tuân theo cùng một mẫu như triển khai trên CPU
Triển khai để thao tác hợp nhất được chạy song song trên CUDA
Các điểm chính liên quan đến triển khai
- Sử dụng cudaMalloc, cudaMemcpy, cudaFree → cấp phát bộ nhớ GPU và truyền dữ liệu
- merge<<<1, 1>>>(...) → chạy thao tác hợp nhất theo kiểu song song
- cudaDeviceSynchronize() → đồng bộ cho đến khi hoàn tất hợp nhất
- Vấn đề hiệu năng → CUDA không hiệu quả với xử lý đệ quy nên cần cách tiếp cận lặp

So sánh triển khai CPU và GPU

Hiệu năng suy giảm vì lời gọi đệ quy được thực thi trên CPU
Lời gọi đệ quy trong CUDA gây ra vấn đề kích thước stack và overhead khởi chạy kernel
Cách cải thiện hiệu năng: cần chuyển sang cách tiếp cận lặp (bottom-up)

Merge sort lặp bottom-up (triển khai trên CPU)

Hợp nhất dần từ các mảng con nhỏ → hiệu quả hơn trên CUDA
Tăng kích thước mảng hợp nhất theo 1, 2, 4, 8, … trong quá trình hợp nhất

Cấu trúc mã chính

MERGE_SORT(arr, temp, start, end)  
  FOR sub_size ← 1 TO end STEP 2 × sub_size DO  
      FOR left ← 0 TO end STEP 2 × sub_size DO  
          mid ← MIN(left + sub_size - 1, end)  
          right ← MIN(left + 2 * sub_size - 1, end)  
          MERGE(arr, temp, left, mid, right)  
      ENDFOR  
  ENDFOR  
END MERGE_SORT

Các điểm chính liên quan đến triển khai
- Nếu kích thước mảng không phải bội số của 2 thì giải quyết bằng cách clamp chỉ số
- Thực hiện thao tác hợp nhất thông qua vòng lặp
- Tiềm năng cải thiện hiệu năng lớn

Merge sort lặp bottom-up (triển khai trên CUDA)

Cải thiện hiệu năng bằng cách chạy merge sort lặp theo kiểu song song
Tính toán số lượng thread và block để chạy song song thao tác hợp nhất

Cấu trúc mã chính

  void mergeSort(uint8_t* arr, uint8_t* temp, long long n) {  
      bool flipflop = true;  
      long long size;  
      for (size = 1; size < n; size *= 2) {  
          numThreads = max(n / (2 * size), (long long)1);  
          gridSize = (numThreads + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;  
          mergeKernel<<<gridSize, THREADS_PER_BLOCK>>>(flipflop ? arr : temp, flipflop ? temp : arr, size, n);  
          CUDA_CHECK(cudaGetLastError());  
          CUDA_CHECK(cudaDeviceSynchronize());  
          flipflop = !flipflop;  
      }  
      if (!flipflop) CUDA_CHECK(cudaMemcpy(arr, temp, n * sizeof(uint8_t), cudaMemcpyDeviceToDevice));  
  }

Các điểm chính
- flipflop → chuyển đổi vị trí lưu kết quả hợp nhất
- gridSize, THREADS_PER_BLOCK → thực hiện song song hóa thao tác hợp nhất
- mergeKernel → gán tác vụ hợp nhất riêng cho từng thread
- Quản lý chỉ số thông qua tính toán chỉ số thread và block

Kết quả hiệu năng

Merge sort bottom-up (CUDA) → cải thiện hiệu năng rõ rệt
- Mảng nhỏ → CPU nhanh hơn
- Mảng lớn → CUDA vượt trội về hiệu năng
thrust::sort → hiệu năng GPU tốt trên mảng lớn
Mức cải thiện hiệu năng của CUDA bị giới hạn bởi overhead truyền dữ liệu

Kết luận và công việc tiếp theo

Đã thành công trong việc cải thiện hiệu năng merge sort dựa trên CUDA
Các điểm chính đã học:
- Học được khái niệm xử lý song song của CUDA và chiến lược tuning hiệu năng
- Merge sort lặp → hiệu quả hơn trên CUDA so với cách tiếp cận đệ quy
- Phát hiện các nút thắt hiệu năng đặc thù của CUDA như đồng bộ thread, truyền bộ nhớ
Các hướng cải thiện tiếp theo:
- Tách và tối ưu công việc giữa CPU-GPU
- Thử nghiệm hiệu năng với các mảng lớn hơn
- Kết hợp thrust::sort với mã tự triển khai
- Tối ưu hiệu năng thông qua sử dụng shared memory

2 bình luận

kandk 2025-03-14

Có vẻ nó được triển khai bằng radix sort trên CUDA, và tôi cũng từng có kinh nghiệm triển khai y hệt như trong phần tham khảo.

GN⁺ 2025-03-13

Ý kiến Hacker News

Đây không phải là cách sắp xếp nhanh trên GPU. Thuật toán nhanh nhất trong CUDA là Onesweep, sử dụng các kỹ thuật phức tạp để tận dụng tính song song của GPU và vượt qua các giới hạn
- Linebender đang triển khai những ý tưởng này theo cách có tính di động hơn trên GPU
- Có thể xem tài liệu liên quan trên trang wiki của Linebender
Cũng như các ý kiến khác, thuật toán này không phù hợp. Những thuật toán như Onesweep rất ấn tượng nhưng khó hiểu
- Nếu nhìn vào radix sort, vốn là thuật toán cốt lõi, sẽ dễ hiểu hơn
- Radix sort có thể được triển khai để song song hóa rất đơn giản, và là một cách tiếp cận đẹp, thanh lịch
Đây là một bài toán đồ chơi thú vị để thử nghiệm. Có thể cải thiện hiệu năng nếu tận dụng các tùy chọn điều phối luồng
- Dùng Nsight để xác định các yếu tố ảnh hưởng đến hiệu năng cũng khá thú vị
- Cũng cần cân nhắc các thuật toán sắp xếp khác. Các mạng sắp xếp như bitonic sort cần nhiều công sức hơn, nhưng có thể chạy nhanh hơn trên phần cứng song song
- Tôi đã làm một bản triển khai đơn giản trên H100, sắp xếp 10M phần tử trong khoảng 10ms. Có thể làm nhanh hơn nữa nếu đầu tư thêm công sức
Ngôn ngữ Futhark giúp sử dụng những thuật toán này trên GPU thuận tiện hơn
- Đây là một ngôn ngữ mức rất cao, biên dịch thành lệnh GPU và có thể truy cập qua thư viện Python
- Trên website có ví dụ triển khai merge sort
Điều này làm tôi nhớ đến một dự án nhỏ thời đại học khi triển khai bitonic sort bằng CUDA
- Có thể xem phần triển khai bitonic sort trên GitHub
Ghi chú giải thích khái niệm đánh chỉ số luồng GPU khá hay
- Có giới thiệu một bài blog cá nhân về lợi ích hiệu năng của sắp xếp vector hóa
Tôi thích bản triển khai radix sort nhanh này
- Nó hoạt động dễ dàng với Cuda driver API và, khác với CUB, không bị giới hạn vào runtime API
- Cũng có bao gồm Onesweep, nhưng tôi chưa làm cho nó chạy được
Tôi từng định dùng nó với Unity, nhưng không vượt qua được nút thắt cổ chai truyền dữ liệu
- Khi dùng compute shader cũng có overhead, nhưng không quá lớn
Muốn đáng để sắp xếp trên GPU thì cần mảng đủ lớn
- Việc truyền dữ liệu giữa RAM và GPU tốn thời gian
Tóm tắt để tiết kiệm thời gian: có người viết một thuật toán sắp xếp trên GPU nhưng nó chậm
- Đây không phải kỹ thuật tối tân, và cũng không rõ tác giả có biết cách sử dụng GPU hiệu quả hay không
- Chỉ là một thử nghiệm lập trình GPU mang tính cá nhân